Multi-Queue Support

blade-engine/src/lib.rs

@@ -458,7 +458,7 @@ impl Engine {
         if !self.track_hot_reloads {
             return;
         }
-        let sync_point = self.pacer.last_sync_point().unwrap();
+        let sync_point = self.pacer.last_sync_point();
         match self.renderer {
             Renderer::RayTracer { ref mut inner, .. } => {
                 inner.hot_reload(&self.asset_hub, &self.gpu_context, sync_point);

blade-graphics/src/gles/command.rs

@@ -170,6 +170,11 @@ impl super::CommandEncoder {
     }
 
     pub fn compute(&mut self, label: &str) -> super::PassEncoder<'_, super::ComputePipeline> {
+        assert_ne!(
+            self.queue_type,
+            crate::QueueType::AsyncTransfer,
+            "compute passes are not supported on transfer queues"
+        );
         self.begin_pass(label);
         self.pass(super::PassKind::Compute)
     }
@@ -179,6 +184,11 @@ impl super::CommandEncoder {
         label: &str,
         targets: crate::RenderTargetSet,
     ) -> super::PassEncoder<'_, super::RenderPipeline> {
+        assert_eq!(
+            self.queue_type,
+            crate::QueueType::Main,
+            "render passes are only supported on the main queue"
+        );
         self.begin_pass(label);
 
         let mut target_size = [0u16; 2];

blade-graphics/src/gles/mod.rs

@@ -382,6 +382,7 @@ struct TimingData {
 
 pub struct CommandEncoder {
     name: String,
+    queue_type: crate::QueueType,
     commands: Vec<Command>,
     plain_data: Vec<u8>,
     string_data: Vec<u8>,
@@ -437,9 +438,9 @@ pub struct PipelineContext<'a> {
     limits: &'a Limits,
 }
 
-#[derive(Clone, Debug)]
+#[derive(Clone, Debug, Default)]
 pub struct SyncPoint {
-    fence: glow::Fence,
+    fence: Option<glow::Fence>,
 }
 //TODO: destructor
 
@@ -458,6 +459,7 @@ impl Context {
             dual_source_blending: false,
             shader_float16: false,
             cooperative_matrix: crate::CooperativeMatrix::default(),
+            queues: vec![crate::QueueType::Main],
         }
     }
 
@@ -513,6 +515,7 @@ impl crate::traits::CommandDevice for Context {
         };
         CommandEncoder {
             name: desc.name.to_string(),
+            queue_type: desc.queue,
             commands: Vec::new(),
             plain_data: Vec::new(),
             string_data: Vec::new(),
@@ -537,7 +540,7 @@ impl crate::traits::CommandDevice for Context {
         }
     }
 
-    fn submit(&self, encoder: &mut CommandEncoder) -> SyncPoint {
+    fn submit(&self, encoder: &mut CommandEncoder, _after: &[SyncPoint]) -> SyncPoint {
         use glow::HasContext as _;
 
         let fence = {
@@ -583,12 +586,16 @@ impl crate::traits::CommandDevice for Context {
         for frame in encoder.present_frames.drain(..) {
             self.platform.present(frame);
         }
-        SyncPoint { fence }
+        SyncPoint { fence: Some(fence) }
     }
 
     fn wait_for(&self, sp: &SyncPoint, timeout_ms: u32) -> Result<bool, crate::DeviceError> {
         use glow::HasContext as _;
 
+        let fence = match sp.fence {
+            Some(fence) => fence,
+            None => return Ok(true), // default SyncPoint is already complete
+        };
         let gl = self.lock();
         let timeout_ns = if timeout_ms == !0 {
             !0
@@ -599,7 +606,7 @@ impl crate::traits::CommandDevice for Context {
         let timeout_ns_i32 = timeout_ns.min(MAX_TIMEOUT) as i32;
 
         let status =
-            unsafe { gl.client_wait_sync(sp.fence, glow::SYNC_FLUSH_COMMANDS_BIT, timeout_ns_i32) };
+            unsafe { gl.client_wait_sync(fence, glow::SYNC_FLUSH_COMMANDS_BIT, timeout_ns_i32) };
         match status {
             glow::ALREADY_SIGNALED | glow::CONDITION_SATISFIED => Ok(true),
             glow::TIMEOUT_EXPIRED => Ok(false),

blade-graphics/src/lib.rs

@@ -162,6 +162,10 @@ pub struct ContextDesc {
     pub overlay: bool,
     /// Force selection of a specific Device ID.
     pub device_id: Option<u32>,
+    /// Enable multi-queue support (async compute and transfer).
+    /// When enabled, every `submit` call must provide explicit
+    /// synchronization via a non-empty list of sync points.
+    pub multi_queue: bool,
 }
 
 #[derive(Debug)]
@@ -253,6 +257,8 @@ pub struct Capabilities {
     pub shader_float16: bool,
     /// Cooperative matrix support.
     pub cooperative_matrix: CooperativeMatrix,
+    /// Available GPU queues. Always contains [`QueueType::Main`].
+    pub queues: Vec<QueueType>,
 }
 
 #[derive(Clone, Debug)]
@@ -870,12 +876,16 @@ pub struct ShaderDesc<'a> {
     pub naga_module: Option<naga::Module>,
 }
 
-#[derive(Clone, Debug, Default, PartialEq)]
-pub enum CommandType {
-    Transfer,
-    Compute,
+/// Type of GPU queue to submit work to.
+#[derive(Clone, Copy, Debug, Default, PartialEq)]
+pub enum QueueType {
+    /// Main graphics+compute+transfer queue.
     #[default]
-    General,
+    Main,
+    /// Dedicated async compute queue.
+    AsyncCompute,
+    /// Dedicated async transfer queue.
+    AsyncTransfer,
 }
 
 pub struct CommandEncoderDesc<'a> {
@@ -884,6 +894,8 @@ pub struct CommandEncoderDesc<'a> {
     /// For example, one buffer is being run on GPU while the
     /// other is being actively encoded, which makes 2.
     pub buffer_count: u32,
+    /// Queue to submit commands to.
+    pub queue: QueueType,
 }
 
 pub struct ComputePipelineDesc<'a> {

blade-graphics/src/metal/command.rs

@@ -229,6 +229,11 @@ impl super::CommandEncoder {
         &mut self,
         label: &str,
     ) -> super::AccelerationStructureCommandEncoder<'_> {
+        assert_ne!(
+            self.queue_type,
+            crate::QueueType::AsyncTransfer,
+            "acceleration structure builds are not supported on transfer queues"
+        );
         let raw = objc2::rc::autoreleasepool(|_| unsafe {
             let descriptor = metal::MTLAccelerationStructurePassDescriptor::new();
 
@@ -255,6 +260,11 @@ impl super::CommandEncoder {
     }
 
     pub fn compute(&mut self, label: &str) -> super::ComputeCommandEncoder<'_> {
+        assert_ne!(
+            self.queue_type,
+            crate::QueueType::AsyncTransfer,
+            "compute passes are not supported on transfer queues"
+        );
         let raw = objc2::rc::autoreleasepool(|_| unsafe {
             let descriptor = metal::MTLComputePassDescriptor::new();
             if self.enable_dispatch_type {
@@ -290,6 +300,11 @@ impl super::CommandEncoder {
         label: &str,
         targets: crate::RenderTargetSet,
     ) -> super::RenderCommandEncoder<'_> {
+        assert_eq!(
+            self.queue_type,
+            crate::QueueType::Main,
+            "render passes are only supported on the main queue"
+        );
         let raw = objc2::rc::autoreleasepool(|_| {
             let descriptor = unsafe { metal::MTLRenderPassDescriptor::new() };
 

blade-graphics/src/metal/mod.rs

@@ -201,9 +201,9 @@ impl AccelerationStructure {
 }
 
 //TODO: make this copyable?
-#[derive(Clone, Debug)]
+#[derive(Clone, Debug, Default)]
 pub struct SyncPoint {
-    cmd_buf: Retained<ProtocolObject<dyn metal::MTLCommandBuffer>>,
+    cmd_buf: Option<Retained<ProtocolObject<dyn metal::MTLCommandBuffer>>>,
 }
 // Safe because all mutability is externalized
 unsafe impl Send for SyncPoint {}
@@ -219,6 +219,7 @@ pub struct CommandEncoder {
     raw: Option<RawCommandBuffer>,
     name: String,
     queue: Arc<Mutex<Retained<ProtocolObject<dyn metal::MTLCommandQueue>>>>,
+    queue_type: crate::QueueType,
     enable_debug_groups: bool,
     enable_dispatch_type: bool,
     has_open_debug_group: bool,
@@ -563,6 +564,7 @@ impl Context {
             } else {
                 crate::CooperativeMatrix::default()
             },
+            queues: vec![crate::QueueType::Main],
         }
     }
 
@@ -662,6 +664,7 @@ impl crate::traits::CommandDevice for Context {
             raw: None,
             name: desc.name.to_string(),
             queue: Arc::clone(&self.queue),
+            queue_type: desc.queue,
             enable_debug_groups: self.info.enable_debug_groups,
             enable_dispatch_type: self.info.enable_dispatch_type,
             has_open_debug_group: false,
@@ -672,18 +675,24 @@ impl crate::traits::CommandDevice for Context {
 
     fn destroy_command_encoder(&self, _command_encoder: &mut CommandEncoder) {}
 
-    fn submit(&self, encoder: &mut CommandEncoder) -> SyncPoint {
+    fn submit(&self, encoder: &mut CommandEncoder, _after: &[SyncPoint]) -> SyncPoint {
         use metal::MTLCommandBuffer as _;
         let cmd_buf = encoder.finish();
         cmd_buf.commit();
-        SyncPoint { cmd_buf }
+        SyncPoint {
+            cmd_buf: Some(cmd_buf),
+        }
     }
 
     fn wait_for(&self, sp: &SyncPoint, timeout_ms: u32) -> Result<bool, crate::DeviceError> {
         use metal::MTLCommandBuffer as _;
+        let cmd_buf = match sp.cmd_buf {
+            Some(ref buf) => buf,
+            None => return Ok(true), // default SyncPoint is already complete
+        };
         let start = time::Instant::now();
         loop {
-            match sp.cmd_buf.status() {
+            match cmd_buf.status() {
                 metal::MTLCommandBufferStatus::Completed => return Ok(true),
                 metal::MTLCommandBufferStatus::Error => return Err(crate::DeviceError::DeviceLost),
                 _ => {}

blade-graphics/src/traits.rs

@@ -43,11 +43,15 @@ pub trait ShaderDevice {
 
 pub trait CommandDevice {
     type CommandEncoder;
-    type SyncPoint: Clone + Debug;
+    type SyncPoint: Clone + Debug + Default;
 
     fn create_command_encoder(&self, desc: super::CommandEncoderDesc) -> Self::CommandEncoder;
     fn destroy_command_encoder(&self, encoder: &mut Self::CommandEncoder);
-    fn submit(&self, encoder: &mut Self::CommandEncoder) -> Self::SyncPoint;
+    fn submit(
+        &self,
+        encoder: &mut Self::CommandEncoder,
+        after: &[Self::SyncPoint],
+    ) -> Self::SyncPoint;
     fn wait_for(&self, sp: &Self::SyncPoint, timeout_ms: u32) -> Result<bool, super::DeviceError>;
 }
 

blade-graphics/src/vulkan/command.rs

@@ -400,6 +400,11 @@ impl super::CommandEncoder {
         &mut self,
         label: &str,
     ) -> super::AccelerationStructureCommandEncoder<'_> {
+        assert_ne!(
+            self.queue_type,
+            crate::QueueType::AsyncTransfer,
+            "acceleration structure builds are not supported on transfer queues"
+        );
         self.begin_pass(label);
         super::AccelerationStructureCommandEncoder {
             raw: self.buffers[0].raw,
@@ -408,6 +413,11 @@ impl super::CommandEncoder {
     }
 
     pub fn compute(&mut self, label: &str) -> super::ComputeCommandEncoder<'_> {
+        assert_ne!(
+            self.queue_type,
+            crate::QueueType::AsyncTransfer,
+            "compute passes are not supported on transfer queues"
+        );
         self.begin_pass(label);
         super::ComputeCommandEncoder {
             cmd_buf: self.buffers.first_mut().unwrap(),
@@ -421,6 +431,11 @@ impl super::CommandEncoder {
         label: &str,
         targets: crate::RenderTargetSet,
     ) -> super::RenderCommandEncoder<'_> {
+        assert_eq!(
+            self.queue_type,
+            crate::QueueType::Main,
+            "render passes are only supported on the main queue"
+        );
         self.begin_pass(label);
 
         let mut target_size = [0u16; 2];
@@ -575,6 +590,8 @@ impl crate::traits::CommandEncoder for super::CommandEncoder {
         let barrier = vk::ImageMemoryBarrier {
             old_layout: vk::ImageLayout::UNDEFINED,
             new_layout: vk::ImageLayout::GENERAL,
+            src_queue_family_index: vk::QUEUE_FAMILY_IGNORED,
+            dst_queue_family_index: vk::QUEUE_FAMILY_IGNORED,
             image: texture.raw,
             subresource_range: vk::ImageSubresourceRange {
                 aspect_mask: super::map_aspects(texture.format.aspects()),
@@ -620,6 +637,8 @@ impl crate::traits::CommandEncoder for super::CommandEncoder {
             let barrier = vk::ImageMemoryBarrier {
                 old_layout: vk::ImageLayout::GENERAL,
                 new_layout: vk::ImageLayout::PRESENT_SRC_KHR,
+                src_queue_family_index: vk::QUEUE_FAMILY_IGNORED,
+                dst_queue_family_index: vk::QUEUE_FAMILY_IGNORED,
                 image: frame.internal.image,
                 subresource_range: vk::ImageSubresourceRange {
                     aspect_mask: vk::ImageAspectFlags::COLOR,