From 08221868e0ee5fd84abebd04af68e23c07741d63 Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Mon, 13 Apr 2026 23:52:55 -0700 Subject: [PATCH 1/4] Extend CommandEncoderDesc, submit(), and ContextDesc with new APIs for multi-queue --- blade-graphics/src/gles/command.rs | 10 +++++++ blade-graphics/src/gles/mod.rs | 4 ++- blade-graphics/src/lib.rs | 20 +++++++++---- blade-graphics/src/metal/command.rs | 15 ++++++++++ blade-graphics/src/metal/mod.rs | 4 ++- blade-graphics/src/traits.rs | 6 +++- blade-graphics/src/vulkan/command.rs | 15 ++++++++++ blade-graphics/src/vulkan/init.rs | 1 + blade-graphics/src/vulkan/mod.rs | 43 +++++++++++++++++++--------- blade-render/src/util/frame_pacer.rs | 3 +- docs/CHANGELOG.md | 4 +++ examples-android/xr/xr.rs | 3 +- examples/bunnymark/example.rs | 3 +- examples/bunnymark/main.rs | 3 +- examples/matmul/main.rs | 3 +- examples/particle/main.rs | 3 +- examples/ray-query/example.rs | 3 +- examples/ray-query/main.rs | 3 +- tests/gpu_examples.rs | 10 +++++-- tests/snapshot.rs | 2 +- 20 files changed, 126 insertions(+), 32 deletions(-) diff --git a/blade-graphics/src/gles/command.rs b/blade-graphics/src/gles/command.rs index acf965b0..0c38aeb6 100644 --- a/blade-graphics/src/gles/command.rs +++ b/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]; diff --git a/blade-graphics/src/gles/mod.rs b/blade-graphics/src/gles/mod.rs index b0b6a754..448f7365 100644 --- a/blade-graphics/src/gles/mod.rs +++ b/blade-graphics/src/gles/mod.rs @@ -382,6 +382,7 @@ struct TimingData { pub struct CommandEncoder { name: String, + queue_type: crate::QueueType, commands: Vec, plain_data: Vec, string_data: Vec, @@ -513,6 +514,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 +539,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 = { diff --git a/blade-graphics/src/lib.rs b/blade-graphics/src/lib.rs index 771d422c..65125b59 100644 --- a/blade-graphics/src/lib.rs +++ b/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, + /// 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)] @@ -870,12 +874,16 @@ pub struct ShaderDesc<'a> { pub naga_module: Option, } -#[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 +892,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> { diff --git a/blade-graphics/src/metal/command.rs b/blade-graphics/src/metal/command.rs index fb87cac4..77247ab7 100644 --- a/blade-graphics/src/metal/command.rs +++ b/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() }; diff --git a/blade-graphics/src/metal/mod.rs b/blade-graphics/src/metal/mod.rs index 3c5df8de..1b938e81 100644 --- a/blade-graphics/src/metal/mod.rs +++ b/blade-graphics/src/metal/mod.rs @@ -219,6 +219,7 @@ pub struct CommandEncoder { raw: Option, name: String, queue: Arc>>>, + queue_type: crate::QueueType, enable_debug_groups: bool, enable_dispatch_type: bool, has_open_debug_group: bool, @@ -662,6 +663,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,7 +674,7 @@ 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(); diff --git a/blade-graphics/src/traits.rs b/blade-graphics/src/traits.rs index 8c721de7..2a6a0240 100644 --- a/blade-graphics/src/traits.rs +++ b/blade-graphics/src/traits.rs @@ -47,7 +47,11 @@ pub trait CommandDevice { 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; } diff --git a/blade-graphics/src/vulkan/command.rs b/blade-graphics/src/vulkan/command.rs index a914089b..0a1aad82 100644 --- a/blade-graphics/src/vulkan/command.rs +++ b/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]; diff --git a/blade-graphics/src/vulkan/init.rs b/blade-graphics/src/vulkan/init.rs index f20666b9..17b24b33 100644 --- a/blade-graphics/src/vulkan/init.rs +++ b/blade-graphics/src/vulkan/init.rs @@ -1375,6 +1375,7 @@ impl super::Context { cooperative_matrix: capabilities.cooperative_matrix, binding_array: capabilities.binding_array, memory_budget: capabilities.memory_budget, + multi_queue: desc.multi_queue, inner, xr, }) diff --git a/blade-graphics/src/vulkan/mod.rs b/blade-graphics/src/vulkan/mod.rs index 4868a59d..e0b2b5f4 100644 --- a/blade-graphics/src/vulkan/mod.rs +++ b/blade-graphics/src/vulkan/mod.rs @@ -277,6 +277,7 @@ pub struct Context { cooperative_matrix: crate::CooperativeMatrix, binding_array: bool, memory_budget: bool, + multi_queue: bool, inner: VulkanInstance, xr: Option>, } @@ -436,6 +437,7 @@ pub struct CommandEncoder { pool: vk::CommandPool, buffers: Box<[CommandBuffer]>, device: Device, + queue_type: crate::QueueType, update_data: Vec, present: Option, crash_handler: Option, @@ -571,6 +573,7 @@ impl crate::traits::CommandDevice for Context { pool, buffers, device: self.device.clone(), + queue_type: desc.queue, update_data: Vec::new(), present: None, crash_handler, @@ -616,37 +619,51 @@ impl crate::traits::CommandDevice for Context { }; } - fn submit(&self, encoder: &mut CommandEncoder) -> SyncPoint { + fn submit(&self, encoder: &mut CommandEncoder, after: &[SyncPoint]) -> SyncPoint { + assert!( + !self.multi_queue || !after.is_empty(), + "multi-queue mode requires explicit sync points in every submit" + ); let raw_cmd_buf = encoder.finish(); let mut queue = self.queue.lock().unwrap(); queue.last_progress += 1; let progress = queue.last_progress; let command_buffers = [raw_cmd_buf]; - let wait_values_all = [0]; - let mut wait_semaphores_all = [vk::Semaphore::null()]; - let wait_stages = [vk::PipelineStageFlags::ALL_COMMANDS]; + + // Build wait semaphore arrays: dependencies first, then optional acquire semaphore + let mut wait_semaphores = Vec::with_capacity(after.len() + 1); + let mut wait_values = Vec::with_capacity(after.len() + 1); + let mut wait_stages = Vec::with_capacity(after.len() + 1); + for sp in after { + wait_semaphores.push(queue.timeline_semaphore); + wait_values.push(sp.progress); + wait_stages.push(vk::PipelineStageFlags::ALL_COMMANDS); + } + let mut signal_semaphores_all = [queue.timeline_semaphore, vk::Semaphore::null()]; let signal_values_all = [progress, 0]; - let (num_wait_semaphores, num_signal_sepahores) = match encoder.present { + let num_signal_semaphores = match encoder.present { Some(Presentation::Window { acquire_semaphore, present_semaphore, .. }) => { - wait_semaphores_all[0] = acquire_semaphore; + wait_semaphores.push(acquire_semaphore); + wait_values.push(0); + wait_stages.push(vk::PipelineStageFlags::ALL_COMMANDS); signal_semaphores_all[1] = present_semaphore; - (1, 2) + 2 } - Some(Presentation::Xr { .. }) | None => (0, 1), + Some(Presentation::Xr { .. }) | None => 1, }; let mut timeline_info = vk::TimelineSemaphoreSubmitInfo::default() - .wait_semaphore_values(&wait_values_all[..num_wait_semaphores]) - .signal_semaphore_values(&signal_values_all[..num_signal_sepahores]); + .wait_semaphore_values(&wait_values) + .signal_semaphore_values(&signal_values_all[..num_signal_semaphores]); let vk_info = vk::SubmitInfo::default() .command_buffers(&command_buffers) - .wait_semaphores(&wait_semaphores_all[..num_wait_semaphores]) - .wait_dst_stage_mask(&wait_stages[..num_wait_semaphores]) - .signal_semaphores(&signal_semaphores_all[..num_signal_sepahores]) + .wait_semaphores(&wait_semaphores) + .wait_dst_stage_mask(&wait_stages) + .signal_semaphores(&signal_semaphores_all[..num_signal_semaphores]) .push_next(&mut timeline_info); let ret = unsafe { self.device diff --git a/blade-render/src/util/frame_pacer.rs b/blade-render/src/util/frame_pacer.rs index a5b75ac4..d997b4b2 100644 --- a/blade-render/src/util/frame_pacer.rs +++ b/blade-render/src/util/frame_pacer.rs @@ -17,6 +17,7 @@ impl FramePacer { let encoder = context.create_command_encoder(blade_graphics::CommandEncoderDesc { name: "main", buffer_count: 2, + queue: blade_graphics::QueueType::Main, }); Self { frame_index: 0, @@ -55,7 +56,7 @@ impl FramePacer { } pub fn end_frame(&mut self, context: &blade_graphics::Context) -> &blade_graphics::SyncPoint { - let sync_point = context.submit(&mut self.command_encoder); + let sync_point = context.submit(&mut self.command_encoder, &[]); self.frame_index += 1; // Wait for the previous frame immediately - this ensures that we are // only processing one frame at a time, and yet not stalling. diff --git a/docs/CHANGELOG.md b/docs/CHANGELOG.md index 54506d7b..882004d4 100644 --- a/docs/CHANGELOG.md +++ b/docs/CHANGELOG.md @@ -1,5 +1,9 @@ Changelog for *Blade* project +## blade-graphics-0.9 (TBD) + +- multi-queue support + ## blade-graphics-0.8.2 (4 Apr 2026) - add `ComputeCommandEncoder::barrier()` for inline compute-to-compute synchronization within a pass diff --git a/examples-android/xr/xr.rs b/examples-android/xr/xr.rs index 4bac360b..e81116dc 100644 --- a/examples-android/xr/xr.rs +++ b/examples-android/xr/xr.rs @@ -116,6 +116,7 @@ impl Example { let command_encoder = context.create_command_encoder(gpu::CommandEncoderDesc { name: "xr", buffer_count: 1, + queue: gpu::QueueType::Main, }); let params_buf = context.create_buffer(gpu::BufferDesc { name: "xr-params", @@ -227,7 +228,7 @@ impl Example { } self.command_encoder.present(frame); - let _sync_point = context.submit(&mut self.command_encoder); + let _sync_point = context.submit(&mut self.command_encoder, &[]); self.rendered_frames += 1; if self.xr_debug && (self.rendered_frames <= 5 || self.rendered_frames % 120 == 0) { info!("XR frame submitted: {}", self.rendered_frames); diff --git a/examples/bunnymark/example.rs b/examples/bunnymark/example.rs index 2736c212..4f913b08 100644 --- a/examples/bunnymark/example.rs +++ b/examples/bunnymark/example.rs @@ -174,13 +174,14 @@ impl Example { let mut command_encoder = context.create_command_encoder(gpu::CommandEncoderDesc { name: "init", buffer_count: 1, + queue: gpu::QueueType::Main, }); command_encoder.start(); command_encoder.init_texture(texture); if let mut transfer = command_encoder.transfer("init texture") { transfer.copy_buffer_to_texture(upload_buffer.into(), 4, texture.into(), extent); } - let sync_point = context.submit(&mut command_encoder); + let sync_point = context.submit(&mut command_encoder, &[]); let _ = context.wait_for(&sync_point, !0); context.destroy_command_encoder(&mut command_encoder); diff --git a/examples/bunnymark/main.rs b/examples/bunnymark/main.rs index bbaddde0..ca900f87 100644 --- a/examples/bunnymark/main.rs +++ b/examples/bunnymark/main.rs @@ -85,6 +85,7 @@ impl winit::application::ApplicationHandler for App { let command_encoder = context.create_command_encoder(gpu::CommandEncoderDesc { name: "main", buffer_count: 2, + queue: gpu::QueueType::Main, }); self.example = Some(example); @@ -173,7 +174,7 @@ impl winit::application::ApplicationHandler for App { command_encoder.init_texture(frame.texture()); example.render(command_encoder, frame.texture_view()); command_encoder.present(frame); - let sync_point = context.submit(command_encoder); + let sync_point = context.submit(command_encoder, &[]); if let Some(sp) = self.prev_sync_point.take() { let _ = context.wait_for(&sp, !0); } diff --git a/examples/matmul/main.rs b/examples/matmul/main.rs index e6ccee46..0e5db835 100644 --- a/examples/matmul/main.rs +++ b/examples/matmul/main.rs @@ -142,6 +142,7 @@ fn main() { let mut encoder = context.create_command_encoder(gpu::CommandEncoderDesc { name: "matmul", buffer_count: 1, + queue: gpu::QueueType::Main, }); encoder.start(); { @@ -158,7 +159,7 @@ fn main() { ); pe.dispatch([M / tile, N / tile, 1]); } - let sp = context.submit(&mut encoder); + let sp = context.submit(&mut encoder, &[]); let _ = context.wait_for(&sp, !0); // Read back results diff --git a/examples/particle/main.rs b/examples/particle/main.rs index 5f528f8f..660f2005 100644 --- a/examples/particle/main.rs +++ b/examples/particle/main.rs @@ -172,6 +172,7 @@ impl Example { let command_encoder = context.create_command_encoder(gpu::CommandEncoderDesc { name: "main", buffer_count: 2, + queue: gpu::QueueType::Main, }); Self { @@ -318,7 +319,7 @@ impl Example { } self.command_encoder.present(frame); - let sync_point = self.context.submit(&mut self.command_encoder); + let sync_point = self.context.submit(&mut self.command_encoder, &[]); self.gui_painter.after_submit(&sync_point); if let Some(sp) = self.prev_sync_point.take() { diff --git a/examples/ray-query/example.rs b/examples/ray-query/example.rs index fed4cb57..ed645de7 100644 --- a/examples/ray-query/example.rs +++ b/examples/ray-query/example.rs @@ -183,6 +183,7 @@ impl Example { let mut command_encoder = context.create_command_encoder(gpu::CommandEncoderDesc { name: "init", buffer_count: 1, + queue: gpu::QueueType::Main, }); command_encoder.start(); command_encoder.init_texture(target); @@ -199,7 +200,7 @@ impl Example { scratch_buffer.at(tlas_scratch_offset), ); } - let sync_point = context.submit(&mut command_encoder); + let sync_point = context.submit(&mut command_encoder, &[]); let _ = context.wait_for(&sync_point, !0); context.destroy_command_encoder(&mut command_encoder); diff --git a/examples/ray-query/main.rs b/examples/ray-query/main.rs index adedcd80..0c6e9e60 100644 --- a/examples/ray-query/main.rs +++ b/examples/ray-query/main.rs @@ -58,6 +58,7 @@ impl winit::application::ApplicationHandler for App { let command_encoder = context.create_command_encoder(gpu::CommandEncoderDesc { name: "main", buffer_count: 2, + queue: gpu::QueueType::Main, }); self.example = Some(example); @@ -107,7 +108,7 @@ impl winit::application::ApplicationHandler for App { command_encoder.init_texture(frame.texture()); example.render(command_encoder, frame.texture_view(), rotation_angle); command_encoder.present(frame); - let sync_point = context.submit(command_encoder); + let sync_point = context.submit(command_encoder, &[]); if let Some(sp) = self.prev_sync_point.take() { let _ = context.wait_for(&sp, !0); diff --git a/tests/gpu_examples.rs b/tests/gpu_examples.rs index 4d7db5ca..f8996287 100644 --- a/tests/gpu_examples.rs +++ b/tests/gpu_examples.rs @@ -242,6 +242,7 @@ fn dispatch_gpu_test() { let mut command_encoder = context.create_command_encoder(gpu::CommandEncoderDesc { name: "dispatch-test", buffer_count: 1, + queue: gpu::QueueType::Main, }); command_encoder.start(); if let mut compute = command_encoder.compute("dispatch") @@ -257,7 +258,7 @@ fn dispatch_gpu_test() { pass.dispatch([1, 1, 1]); } - let sync_point = context.submit(&mut command_encoder); + let sync_point = context.submit(&mut command_encoder, &[]); assert!(context.wait_for(&sync_point, 2000).unwrap()); let actual = unsafe { slice::from_raw_parts(output.data() as *const u32, 4) }; @@ -287,6 +288,7 @@ fn env_map_gpu_test() { let mut command_encoder = context.create_command_encoder(gpu::CommandEncoderDesc { name: "env-map-test", buffer_count: 1, + queue: gpu::QueueType::Main, }); command_encoder.start(); @@ -320,7 +322,7 @@ fn env_map_gpu_test() { ); } - let sync_point = context.submit(&mut command_encoder); + let sync_point = context.submit(&mut command_encoder, &[]); assert!(context.wait_for(&sync_point, 2000).unwrap()); let actual = unsafe { slice::from_raw_parts(readback.data(), 8) }; @@ -359,6 +361,7 @@ fn snapshot_bunnymark() { let mut command_encoder = context.create_command_encoder(gpu::CommandEncoderDesc { name: "snapshot-bunnymark", buffer_count: 1, + queue: gpu::QueueType::Main, }); command_encoder.start(); command_encoder.init_texture(target.texture); @@ -417,6 +420,7 @@ fn snapshot_ray_query() { let mut command_encoder = context.create_command_encoder(gpu::CommandEncoderDesc { name: "snapshot-ray-query", buffer_count: 1, + queue: gpu::QueueType::Main, }); command_encoder.start(); command_encoder.init_texture(target.texture); @@ -472,6 +476,7 @@ fn snapshot_particle() { let mut command_encoder = context.create_command_encoder(gpu::CommandEncoderDesc { name: "snapshot-particle", buffer_count: 1, + queue: gpu::QueueType::Main, }); command_encoder.start(); // Run several update cycles to emit and move particles @@ -613,6 +618,7 @@ fn snapshot_space_sky() { let mut command_encoder = context.create_command_encoder(gpu::CommandEncoderDesc { name: "sky-test", buffer_count: 1, + queue: gpu::QueueType::Main, }); command_encoder.start(); command_encoder.init_texture(target.texture); diff --git a/tests/snapshot.rs b/tests/snapshot.rs index 51721d72..54801984 100644 --- a/tests/snapshot.rs +++ b/tests/snapshot.rs @@ -65,7 +65,7 @@ impl OffscreenTarget { self.size, ); } - let sync_point = context.submit(encoder); + let sync_point = context.submit(encoder, &[]); assert!( context.wait_for(&sync_point, 5000).unwrap(), "GPU timed out during snapshot readback" From ecbfe3a0cb0a691ddd72211bfd608dd960fb2f8d Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Tue, 14 Apr 2026 01:24:45 -0700 Subject: [PATCH 2/4] Make SyncPoint: Default, implement queues on Vulkan --- blade-engine/src/lib.rs | 2 +- blade-graphics/src/gles/mod.rs | 12 +- blade-graphics/src/metal/mod.rs | 14 ++- blade-graphics/src/traits.rs | 2 +- blade-graphics/src/vulkan/init.rs | 158 ++++++++++++++++++++------ blade-graphics/src/vulkan/mod.rs | 63 +++++++--- blade-graphics/src/vulkan/resource.rs | 42 ++++++- blade-render/src/util/frame_pacer.rs | 16 ++- docs/CHANGELOG.md | 1 + examples/bunnymark/main.rs | 14 +-- examples/particle/main.rs | 63 +++++----- examples/ray-query/main.rs | 14 +-- examples/scene/main.rs | 2 +- 13 files changed, 288 insertions(+), 115 deletions(-) diff --git a/blade-engine/src/lib.rs b/blade-engine/src/lib.rs index 0876c68d..f46c3330 100644 --- a/blade-engine/src/lib.rs +++ b/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); diff --git a/blade-graphics/src/gles/mod.rs b/blade-graphics/src/gles/mod.rs index 448f7365..d845e6e2 100644 --- a/blade-graphics/src/gles/mod.rs +++ b/blade-graphics/src/gles/mod.rs @@ -438,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, } //TODO: destructor @@ -585,12 +585,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 { 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 @@ -601,7 +605,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), diff --git a/blade-graphics/src/metal/mod.rs b/blade-graphics/src/metal/mod.rs index 1b938e81..ffcc701d 100644 --- a/blade-graphics/src/metal/mod.rs +++ b/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>, + cmd_buf: Option>>, } // Safe because all mutability is externalized unsafe impl Send for SyncPoint {} @@ -678,14 +678,20 @@ impl crate::traits::CommandDevice for Context { 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 { 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), _ => {} diff --git a/blade-graphics/src/traits.rs b/blade-graphics/src/traits.rs index 2a6a0240..9069567d 100644 --- a/blade-graphics/src/traits.rs +++ b/blade-graphics/src/traits.rs @@ -43,7 +43,7 @@ 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); diff --git a/blade-graphics/src/vulkan/init.rs b/blade-graphics/src/vulkan/init.rs index 17b24b33..e1db15cc 100644 --- a/blade-graphics/src/vulkan/init.rs +++ b/blade-graphics/src/vulkan/init.rs @@ -93,6 +93,8 @@ struct AdapterCapabilities { properties: vk::PhysicalDeviceProperties, device_information: crate::DeviceInformation, queue_family_index: u32, + async_compute_queue_family: Option, + async_transfer_queue_family: Option, layered: bool, binding_array: bool, ray_tracing: Option, @@ -287,7 +289,42 @@ fn inspect_adapter( intel_fix_descriptor_pool_leak: cfg!(windows) && properties.vendor_id == db::intel::VENDOR, }; - let queue_family_index = 0; //TODO + let queue_families = unsafe { + instance + .core + .get_physical_device_queue_family_properties(phd) + }; + // Find main (graphics) queue family + let queue_family_index = queue_families + .iter() + .position(|qf| qf.queue_flags.contains(vk::QueueFlags::GRAPHICS)) + .unwrap_or(0) as u32; + // Find dedicated async compute queue family (compute+transfer but NOT graphics) + let async_compute_queue_family = if desc.multi_queue { + queue_families + .iter() + .position(|qf| { + qf.queue_flags.contains(vk::QueueFlags::COMPUTE) + && !qf.queue_flags.contains(vk::QueueFlags::GRAPHICS) + }) + .map(|i| i as u32) + } else { + None + }; + // Find dedicated async transfer queue family (transfer only, not compute or graphics) + let async_transfer_queue_family = if desc.multi_queue { + queue_families + .iter() + .position(|qf| { + qf.queue_flags.contains(vk::QueueFlags::TRANSFER) + && !qf.queue_flags.contains(vk::QueueFlags::COMPUTE) + && !qf.queue_flags.contains(vk::QueueFlags::GRAPHICS) + }) + .map(|i| i as u32) + } else { + None + }; + if desc.presentation && is_presentation_broken(properties.vendor_id, gpu_vendors, display_server) { @@ -534,6 +571,8 @@ fn inspect_adapter( properties, device_information, queue_family_index, + async_compute_queue_family, + async_transfer_queue_family, layered: portability_subset_properties.min_vertex_input_binding_stride_alignment != 0, binding_array, ray_tracing, @@ -892,10 +931,25 @@ impl super::Context { } let device_core = { - let family_info = vk::DeviceQueueCreateInfo::default() - .queue_family_index(capabilities.queue_family_index) - .queue_priorities(&[1.0]); - let family_infos = [family_info]; + let mut family_infos = vec![ + vk::DeviceQueueCreateInfo::default() + .queue_family_index(capabilities.queue_family_index) + .queue_priorities(&[1.0]), + ]; + if let Some(qfi) = capabilities.async_compute_queue_family { + family_infos.push( + vk::DeviceQueueCreateInfo::default() + .queue_family_index(qfi) + .queue_priorities(&[1.0]), + ); + } + if let Some(qfi) = capabilities.async_transfer_queue_family { + family_infos.push( + vk::DeviceQueueCreateInfo::default() + .queue_family_index(qfi) + .queue_priorities(&[1.0]), + ); + } let mut device_extensions = REQUIRED_DEVICE_EXTENSIONS.to_vec(); if capabilities.max_inline_uniform_block_size > 0 { @@ -1262,26 +1316,41 @@ impl super::Context { } }; - let queue = unsafe { - device - .core - .get_device_queue(capabilities.queue_family_index, 0) - }; - let last_progress = 0; - let mut timeline_info = vk::SemaphoreTypeCreateInfo { - semaphore_type: vk::SemaphoreType::TIMELINE, - initial_value: last_progress, - ..Default::default() - }; - let timeline_semaphore_create_info = - vk::SemaphoreCreateInfo::default().push_next(&mut timeline_info); - let timeline_semaphore = unsafe { - device - .core - .create_semaphore(&timeline_semaphore_create_info, None) - .unwrap() + let create_queue = |queue_family_index: u32| -> super::Queue { + let raw = unsafe { device.core.get_device_queue(queue_family_index, 0) }; + let mut timeline_info = vk::SemaphoreTypeCreateInfo { + semaphore_type: vk::SemaphoreType::TIMELINE, + initial_value: 0, + ..Default::default() + }; + let semaphore_info = vk::SemaphoreCreateInfo::default().push_next(&mut timeline_info); + let timeline_semaphore = + unsafe { device.core.create_semaphore(&semaphore_info, None).unwrap() }; + super::Queue { + raw, + timeline_semaphore, + last_progress: 0, + family_index: queue_family_index, + } }; + let queue = create_queue(capabilities.queue_family_index); + let async_compute_queue = capabilities.async_compute_queue_family.map(&create_queue); + let async_transfer_queue = capabilities.async_transfer_queue_family.map(&create_queue); + + if async_compute_queue.is_some() { + log::info!( + "Async compute queue family: {}", + capabilities.async_compute_queue_family.unwrap() + ); + } + if async_transfer_queue.is_some() { + log::info!( + "Async transfer queue family: {}", + capabilities.async_transfer_queue_family.unwrap() + ); + } + let mut naga_flags = spv::WriterFlags::FORCE_POINT_SIZE; let shader_debug_path = if desc.validation || desc.capture { use std::{env, fs}; @@ -1345,15 +1414,23 @@ impl super::Context { None }; + // Collect all queue family indices for concurrent resource sharing + let mut queue_family_indices = vec![capabilities.queue_family_index]; + if let Some(qfi) = capabilities.async_compute_queue_family { + queue_family_indices.push(qfi); + } + if let Some(qfi) = capabilities.async_transfer_queue_family { + queue_family_indices.push(qfi); + } + Ok(super::Context { memory: Mutex::new(memory_manager), device, queue_family_index: capabilities.queue_family_index, - queue: Mutex::new(super::Queue { - raw: queue, - timeline_semaphore, - last_progress, - }), + queue: Mutex::new(queue), + async_compute_queue: async_compute_queue.map(Mutex::new), + async_transfer_queue: async_transfer_queue.map(Mutex::new), + queue_family_indices: queue_family_indices.into_boxed_slice(), physical_device, naga_flags, shader_debug_path, @@ -1472,14 +1549,25 @@ impl super::Context { impl Drop for super::Context { fn drop(&mut self) { - unsafe { - self.xr = None; - if let Ok(queue) = self.queue.lock() { - let _ = self.device.core.queue_wait_idle(queue.raw); - self.device - .core - .destroy_semaphore(queue.timeline_semaphore, None); + let destroy_queue = |device: &super::Device, queue: &Mutex| { + if let Ok(queue) = queue.lock() { + unsafe { + let _ = device.core.queue_wait_idle(queue.raw); + device + .core + .destroy_semaphore(queue.timeline_semaphore, None); + } } + }; + self.xr = None; + if let Some(ref q) = self.async_transfer_queue { + destroy_queue(&self.device, q); + } + if let Some(ref q) = self.async_compute_queue { + destroy_queue(&self.device, q); + } + destroy_queue(&self.device, &self.queue); + unsafe { self.device.core.destroy_device(None); // inner.drop() destroys the Vulkan instance } diff --git a/blade-graphics/src/vulkan/mod.rs b/blade-graphics/src/vulkan/mod.rs index e0b2b5f4..05cba0a1 100644 --- a/blade-graphics/src/vulkan/mod.rs +++ b/blade-graphics/src/vulkan/mod.rs @@ -84,6 +84,7 @@ struct Queue { raw: vk::Queue, timeline_semaphore: vk::Semaphore, last_progress: u64, + family_index: u32, } #[derive(Clone, Copy, Debug, Default, PartialEq)] @@ -266,6 +267,10 @@ pub struct Context { device: Device, queue_family_index: u32, queue: Mutex, + async_compute_queue: Option>, + async_transfer_queue: Option>, + /// All unique queue family indices, for CONCURRENT sharing mode. + queue_family_indices: Box<[u32]>, physical_device: vk::PhysicalDevice, naga_flags: naga::back::spv::WriterFlags, shader_debug_path: Option, @@ -474,9 +479,10 @@ pub struct PipelineEncoder<'a, 'p> { update_data: &'a mut Vec, } -#[derive(Clone, Debug)] +#[derive(Clone, Debug, Default)] pub struct SyncPoint { progress: u64, + timeline_semaphore: vk::Semaphore, } #[hidden_trait::expose] @@ -485,8 +491,22 @@ impl crate::traits::CommandDevice for Context { type SyncPoint = SyncPoint; fn create_command_encoder(&self, desc: super::CommandEncoderDesc) -> CommandEncoder { + let queue_family_index = match desc.queue { + crate::QueueType::Main => self.queue_family_index, + crate::QueueType::AsyncCompute => self + .async_compute_queue + .as_ref() + .map(|q| q.lock().unwrap().family_index) + .unwrap_or(self.queue_family_index), + crate::QueueType::AsyncTransfer => self + .async_transfer_queue + .as_ref() + .map(|q| q.lock().unwrap().family_index) + .unwrap_or(self.queue_family_index), + }; let pool_info = vk::CommandPoolCreateInfo { flags: vk::CommandPoolCreateFlags::RESET_COMMAND_BUFFER, + queue_family_index, ..Default::default() }; let pool = unsafe { @@ -620,22 +640,36 @@ impl crate::traits::CommandDevice for Context { } fn submit(&self, encoder: &mut CommandEncoder, after: &[SyncPoint]) -> SyncPoint { - assert!( - !self.multi_queue || !after.is_empty(), - "multi-queue mode requires explicit sync points in every submit" - ); + if self.multi_queue && after.is_empty() { + log::warn!("multi-queue mode: submit without explicit sync points"); + } let raw_cmd_buf = encoder.finish(); - let mut queue = self.queue.lock().unwrap(); + + // Lock the target queue based on encoder type, falling back to main + let queue_mutex = match encoder.queue_type { + crate::QueueType::AsyncCompute => { + self.async_compute_queue.as_ref().unwrap_or(&self.queue) + } + crate::QueueType::AsyncTransfer => { + self.async_transfer_queue.as_ref().unwrap_or(&self.queue) + } + crate::QueueType::Main => &self.queue, + }; + let mut queue = queue_mutex.lock().unwrap(); queue.last_progress += 1; let progress = queue.last_progress; let command_buffers = [raw_cmd_buf]; - // Build wait semaphore arrays: dependencies first, then optional acquire semaphore + // Build wait semaphore arrays: dependencies first, then optional acquire semaphore. + // Each SyncPoint carries its own timeline semaphore, so cross-queue waits work. let mut wait_semaphores = Vec::with_capacity(after.len() + 1); let mut wait_values = Vec::with_capacity(after.len() + 1); let mut wait_stages = Vec::with_capacity(after.len() + 1); for sp in after { - wait_semaphores.push(queue.timeline_semaphore); + if sp.timeline_semaphore == vk::Semaphore::null() { + continue; // skip default (no-op) sync points + } + wait_semaphores.push(sp.timeline_semaphore); wait_values.push(sp.progress); wait_stages.push(vk::PipelineStageFlags::ALL_COMMANDS); } @@ -777,14 +811,17 @@ impl crate::traits::CommandDevice for Context { } } - SyncPoint { progress } + SyncPoint { + progress, + timeline_semaphore: queue.timeline_semaphore, + } } fn wait_for(&self, sp: &SyncPoint, timeout_ms: u32) -> Result { - //Note: technically we could get away without locking the queue, - // but also this isn't time-sensitive, so it's fine. - let timeline_semaphore = self.queue.lock().unwrap().timeline_semaphore; - let semaphores = [timeline_semaphore]; + if sp.timeline_semaphore == vk::Semaphore::null() { + return Ok(true); // default SyncPoint is already complete + } + let semaphores = [sp.timeline_semaphore]; let semaphore_values = [sp.progress]; let wait_info = vk::SemaphoreWaitInfoKHR::default() .semaphores(&semaphores) diff --git a/blade-graphics/src/vulkan/resource.rs b/blade-graphics/src/vulkan/resource.rs index 1e093536..0f26d943 100644 --- a/blade-graphics/src/vulkan/resource.rs +++ b/blade-graphics/src/vulkan/resource.rs @@ -297,6 +297,16 @@ impl crate::traits::ResourceDevice for super::Context { handle_types: external_source_handle_type(e), ..Default::default() }); + let (sharing_mode, queue_family_index_count, p_queue_family_indices) = + if self.queue_family_indices.len() > 1 { + ( + vk::SharingMode::CONCURRENT, + self.queue_family_indices.len() as u32, + self.queue_family_indices.as_ptr(), + ) + } else { + (vk::SharingMode::EXCLUSIVE, 0, ptr::null()) + }; let mut vk_info = vk::BufferCreateInfo { size: desc.size, usage: Buf::TRANSFER_SRC @@ -305,7 +315,9 @@ impl crate::traits::ResourceDevice for super::Context { | Buf::INDEX_BUFFER | Buf::VERTEX_BUFFER | Buf::INDIRECT_BUFFER, - sharing_mode: vk::SharingMode::EXCLUSIVE, + sharing_mode, + queue_family_index_count, + p_queue_family_indices, ..Default::default() }; // Always include UNIFORM_BUFFER usage: even when inline uniform blocks @@ -386,6 +398,16 @@ impl crate::traits::ResourceDevice for super::Context { ..Default::default() }); + let (sharing_mode, queue_family_index_count, p_queue_family_indices) = + if self.queue_family_indices.len() > 1 { + ( + vk::SharingMode::CONCURRENT, + self.queue_family_indices.len() as u32, + self.queue_family_indices.as_ptr(), + ) + } else { + (vk::SharingMode::EXCLUSIVE, 0, ptr::null()) + }; let mut vk_info = vk::ImageCreateInfo { flags: create_flags, image_type: map_texture_dimension(desc.dimension), @@ -396,7 +418,9 @@ impl crate::traits::ResourceDevice for super::Context { samples: vk::SampleCountFlags::from_raw(desc.sample_count), tiling: vk::ImageTiling::OPTIMAL, usage: map_texture_usage(desc.usage, desc.format.aspects()), - sharing_mode: vk::SharingMode::EXCLUSIVE, + sharing_mode, + queue_family_index_count, + p_queue_family_indices, ..Default::default() }; @@ -528,11 +552,23 @@ impl crate::traits::ResourceDevice for super::Context { &self, desc: crate::AccelerationStructureDesc, ) -> super::AccelerationStructure { + let (sharing_mode, queue_family_index_count, p_queue_family_indices) = + if self.queue_family_indices.len() > 1 { + ( + vk::SharingMode::CONCURRENT, + self.queue_family_indices.len() as u32, + self.queue_family_indices.as_ptr(), + ) + } else { + (vk::SharingMode::EXCLUSIVE, 0, ptr::null()) + }; let buffer_info = vk::BufferCreateInfo { size: desc.size, usage: vk::BufferUsageFlags::ACCELERATION_STRUCTURE_STORAGE_KHR | vk::BufferUsageFlags::SHADER_DEVICE_ADDRESS, - sharing_mode: vk::SharingMode::EXCLUSIVE, + sharing_mode, + queue_family_index_count, + p_queue_family_indices, ..Default::default() }; diff --git a/blade-render/src/util/frame_pacer.rs b/blade-render/src/util/frame_pacer.rs index d997b4b2..d233fa35 100644 --- a/blade-render/src/util/frame_pacer.rs +++ b/blade-render/src/util/frame_pacer.rs @@ -7,7 +7,7 @@ use std::mem; pub struct FramePacer { frame_index: usize, prev_resources: FrameResources, - prev_sync_point: Option, + prev_sync_point: blade_graphics::SyncPoint, command_encoder: blade_graphics::CommandEncoder, next_resources: FrameResources, } @@ -22,7 +22,7 @@ impl FramePacer { Self { frame_index: 0, prev_resources: FrameResources::default(), - prev_sync_point: None, + prev_sync_point: blade_graphics::SyncPoint::default(), command_encoder: encoder, next_resources: FrameResources::default(), } @@ -30,9 +30,7 @@ impl FramePacer { #[profiling::function] pub fn wait_for_previous_frame(&mut self, context: &blade_graphics::Context) { - if let Some(sp) = self.prev_sync_point.take() { - let _ = context.wait_for(&sp, !0); - } + let _ = context.wait_for(&self.prev_sync_point, !0); for buffer in self.prev_resources.buffers.drain(..) { context.destroy_buffer(buffer); } @@ -41,8 +39,8 @@ impl FramePacer { } } - pub fn last_sync_point(&self) -> Option<&blade_graphics::SyncPoint> { - self.prev_sync_point.as_ref() + pub fn last_sync_point(&self) -> &blade_graphics::SyncPoint { + &self.prev_sync_point } pub fn destroy(&mut self, context: &blade_graphics::Context) { @@ -61,9 +59,9 @@ impl FramePacer { // Wait for the previous frame immediately - this ensures that we are // only processing one frame at a time, and yet not stalling. self.wait_for_previous_frame(context); - self.prev_sync_point = Some(sync_point); + self.prev_sync_point = sync_point; mem::swap(&mut self.prev_resources, &mut self.next_resources); - self.prev_sync_point.as_ref().unwrap() + &self.prev_sync_point } pub fn timings(&self) -> &blade_graphics::Timings { diff --git a/docs/CHANGELOG.md b/docs/CHANGELOG.md index 882004d4..bcd7a595 100644 --- a/docs/CHANGELOG.md +++ b/docs/CHANGELOG.md @@ -3,6 +3,7 @@ Changelog for *Blade* project ## blade-graphics-0.9 (TBD) - multi-queue support +- `SyncPoint: Default` ## blade-graphics-0.8.2 (4 Apr 2026) diff --git a/examples/bunnymark/main.rs b/examples/bunnymark/main.rs index ca900f87..9ec23766 100644 --- a/examples/bunnymark/main.rs +++ b/examples/bunnymark/main.rs @@ -21,7 +21,7 @@ fn make_surface_config(size: winit::dpi::PhysicalSize) -> gpu::SurfaceConfi struct App { example: Option, command_encoder: Option, - prev_sync_point: Option, + prev_sync_point: gpu::SyncPoint, surface: Option, context: Option, window: Option, @@ -175,10 +175,8 @@ impl winit::application::ApplicationHandler for App { example.render(command_encoder, frame.texture_view()); command_encoder.present(frame); let sync_point = context.submit(command_encoder, &[]); - if let Some(sp) = self.prev_sync_point.take() { - let _ = context.wait_for(&sp, !0); - } - self.prev_sync_point = Some(sync_point); + let _ = context.wait_for(&self.prev_sync_point, !0); + self.prev_sync_point = sync_point; } _ => {} } @@ -193,7 +191,7 @@ fn main() { let mut app = App { example: None, command_encoder: None, - prev_sync_point: None, + prev_sync_point: gpu::SyncPoint::default(), surface: None, context: None, window: None, @@ -204,9 +202,7 @@ fn main() { event_loop.run_app(&mut app).unwrap(); let context = app.context.as_ref().unwrap(); - if let Some(sp) = app.prev_sync_point.take() { - let _ = context.wait_for(&sp, !0); - } + let _ = context.wait_for(&app.prev_sync_point, !0); if let Some(mut example) = app.example.take() { example.deinit(context); } diff --git a/examples/particle/main.rs b/examples/particle/main.rs index 660f2005..9b7cb65b 100644 --- a/examples/particle/main.rs +++ b/examples/particle/main.rs @@ -4,7 +4,8 @@ use blade_graphics as gpu; struct Example { command_encoder: gpu::CommandEncoder, - prev_sync_point: Option, + compute_encoder: gpu::CommandEncoder, + prev_sync_point: gpu::SyncPoint, context: gpu::Context, surface: gpu::Surface, gui_painter: blade_egui::GuiPainter, @@ -27,9 +28,7 @@ impl Example { let surface_info = self.surface.info(); - if let Some(sp) = self.prev_sync_point.take() { - let _ = self.context.wait_for(&sp, !0); - } + let _ = self.context.wait_for(&self.prev_sync_point, !0); if let Some(msaa_view) = self.msaa_view.take() { self.context.destroy_texture_view(msaa_view); @@ -120,7 +119,7 @@ impl Example { validation: cfg!(debug_assertions), timing: true, capture: false, - + multi_queue: true, ..Default::default() }) .unwrap() @@ -174,10 +173,16 @@ impl Example { buffer_count: 2, queue: gpu::QueueType::Main, }); + let compute_encoder = context.create_command_encoder(gpu::CommandEncoderDesc { + name: "async-compute", + buffer_count: 2, + queue: gpu::QueueType::AsyncCompute, + }); Self { command_encoder, - prev_sync_point: None, + compute_encoder, + prev_sync_point: gpu::SyncPoint::default(), context, surface, gui_painter, @@ -192,11 +197,11 @@ impl Example { } fn destroy(&mut self) { - if let Some(sp) = self.prev_sync_point.take() { - let _ = self.context.wait_for(&sp, !0); - } + let _ = self.context.wait_for(&self.prev_sync_point, !0); self.context .destroy_command_encoder(&mut self.command_encoder); + self.context + .destroy_command_encoder(&mut self.compute_encoder); self.gui_painter.destroy(&self.context); self.particle_system.destroy(&self.context); self.particle_pipeline.destroy(&self.context); @@ -242,14 +247,6 @@ impl Example { let frame = self.surface.acquire_frame(); let frame_view = frame.texture_view(); - self.command_encoder.start(); - if let Some(msaa_texture) = self.msaa_texture { - self.command_encoder.init_texture(msaa_texture); - } - self.command_encoder.init_texture(frame.texture()); - - self.gui_painter - .update_textures(&mut self.command_encoder, gui_textures, &self.context); // Orbit the emitter and rotate its axis self.time += 0.01; @@ -261,8 +258,24 @@ impl Example { ]; let spray_angle = self.time * 3.0; self.particle_system.axis = [spray_angle.cos(), spray_angle.sin(), 0.0]; + + // Submit particle update on async compute queue. + // Wait for previous frame's render to finish before reusing particle buffers. + let prev_sp = self.prev_sync_point.clone(); + self.compute_encoder.start(); self.particle_system - .update(&self.particle_pipeline, &mut self.command_encoder, 0.01); + .update(&self.particle_pipeline, &mut self.compute_encoder, 0.01); + let sp_compute = self.context.submit(&mut self.compute_encoder, &[prev_sp]); + + // Record rendering on main queue, waiting for compute to finish. + self.command_encoder.start(); + if let Some(msaa_texture) = self.msaa_texture { + self.command_encoder.init_texture(msaa_texture); + } + self.command_encoder.init_texture(frame.texture()); + + self.gui_painter + .update_textures(&mut self.command_encoder, gui_textures, &self.context); let camera = self.make_camera(screen_desc.physical_size); @@ -319,13 +332,13 @@ impl Example { } self.command_encoder.present(frame); - let sync_point = self.context.submit(&mut self.command_encoder, &[]); + let sync_point = self + .context + .submit(&mut self.command_encoder, &[sp_compute]); self.gui_painter.after_submit(&sync_point); - if let Some(sp) = self.prev_sync_point.take() { - let _ = self.context.wait_for(&sp, !0); - } - self.prev_sync_point = Some(sync_point); + let _ = self.context.wait_for(&self.prev_sync_point, !0); + self.prev_sync_point = sync_point; } fn add_gui(&mut self, ui: &mut egui::Ui) { @@ -354,9 +367,7 @@ impl Example { .selectable_value(&mut self.sample_count, i, format!("x{i}")) .changed() { - if let Some(sp) = self.prev_sync_point.take() { - let _ = self.context.wait_for(&sp, !0); - } + let _ = self.context.wait_for(&self.prev_sync_point, !0); let old_effect = self.particle_system.effect.clone(); self.particle_system.destroy(&self.context); diff --git a/examples/ray-query/main.rs b/examples/ray-query/main.rs index 0c6e9e60..ec86371e 100644 --- a/examples/ray-query/main.rs +++ b/examples/ray-query/main.rs @@ -8,7 +8,7 @@ use std::time; struct App { example: Option, command_encoder: Option, - prev_sync_point: Option, + prev_sync_point: gpu::SyncPoint, surface: Option, context: Option, window: Option, @@ -110,10 +110,8 @@ impl winit::application::ApplicationHandler for App { command_encoder.present(frame); let sync_point = context.submit(command_encoder, &[]); - if let Some(sp) = self.prev_sync_point.take() { - let _ = context.wait_for(&sp, !0); - } - self.prev_sync_point = Some(sync_point); + let _ = context.wait_for(&self.prev_sync_point, !0); + self.prev_sync_point = sync_point; } winit::event::WindowEvent::CloseRequested => { event_loop.exit(); @@ -130,7 +128,7 @@ fn main() { let mut app = App { example: None, command_encoder: None, - prev_sync_point: None, + prev_sync_point: gpu::SyncPoint::default(), surface: None, context: None, window: None, @@ -139,9 +137,7 @@ fn main() { event_loop.run_app(&mut app).unwrap(); let context = app.context.as_ref().unwrap(); - if let Some(sp) = app.prev_sync_point.take() { - let _ = context.wait_for(&sp, !0); - } + let _ = context.wait_for(&app.prev_sync_point, !0); if let Some(example) = app.example.take() { example.deinit(context); } diff --git a/examples/scene/main.rs b/examples/scene/main.rs index a7f78467..5465d808 100644 --- a/examples/scene/main.rs +++ b/examples/scene/main.rs @@ -398,7 +398,7 @@ impl Example { self.need_accumulation_reset |= self.renderer.hot_reload( &self.asset_hub, &self.context, - self.pacer.last_sync_point().unwrap(), + self.pacer.last_sync_point(), ); } From 2967b65418eefdddbe87d8815b890cf3133e06ca Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Thu, 16 Apr 2026 00:10:45 -0700 Subject: [PATCH 3/4] Expose the queue types in capabilities, rewire the particle example to be parallel --- blade-graphics/src/gles/mod.rs | 1 + blade-graphics/src/lib.rs | 2 + blade-graphics/src/metal/mod.rs | 1 + blade-graphics/src/vulkan/command.rs | 4 + blade-graphics/src/vulkan/init.rs | 16 ++ examples/particle/main.rs | 314 +++++++++++++++++++-------- 6 files changed, 246 insertions(+), 92 deletions(-) diff --git a/blade-graphics/src/gles/mod.rs b/blade-graphics/src/gles/mod.rs index d845e6e2..c15d365e 100644 --- a/blade-graphics/src/gles/mod.rs +++ b/blade-graphics/src/gles/mod.rs @@ -459,6 +459,7 @@ impl Context { dual_source_blending: false, shader_float16: false, cooperative_matrix: crate::CooperativeMatrix::default(), + queues: vec![crate::QueueType::Main], } } diff --git a/blade-graphics/src/lib.rs b/blade-graphics/src/lib.rs index 65125b59..5f478c1f 100644 --- a/blade-graphics/src/lib.rs +++ b/blade-graphics/src/lib.rs @@ -257,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, } #[derive(Clone, Debug)] diff --git a/blade-graphics/src/metal/mod.rs b/blade-graphics/src/metal/mod.rs index ffcc701d..8adf75e5 100644 --- a/blade-graphics/src/metal/mod.rs +++ b/blade-graphics/src/metal/mod.rs @@ -564,6 +564,7 @@ impl Context { } else { crate::CooperativeMatrix::default() }, + queues: vec![crate::QueueType::Main], } } diff --git a/blade-graphics/src/vulkan/command.rs b/blade-graphics/src/vulkan/command.rs index 0a1aad82..86aefef7 100644 --- a/blade-graphics/src/vulkan/command.rs +++ b/blade-graphics/src/vulkan/command.rs @@ -590,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()), @@ -635,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, diff --git a/blade-graphics/src/vulkan/init.rs b/blade-graphics/src/vulkan/init.rs index e1db15cc..48e96259 100644 --- a/blade-graphics/src/vulkan/init.rs +++ b/blade-graphics/src/vulkan/init.rs @@ -115,6 +115,13 @@ struct AdapterCapabilities { impl AdapterCapabilities { fn to_capabilities(&self) -> crate::Capabilities { + let mut queues = vec![crate::QueueType::Main]; + if self.async_compute_queue_family.is_some() { + queues.push(crate::QueueType::AsyncCompute); + } + if self.async_transfer_queue_family.is_some() { + queues.push(crate::QueueType::AsyncTransfer); + } crate::Capabilities { binding_array: self.binding_array, ray_query: match self.ray_tracing { @@ -127,6 +134,7 @@ impl AdapterCapabilities { dual_source_blending: self.dual_source_blending, shader_float16: self.shader_float16, cooperative_matrix: self.cooperative_matrix, + queues, } } } @@ -1471,6 +1479,13 @@ impl super::Context { } pub fn capabilities(&self) -> crate::Capabilities { + let mut queues = vec![crate::QueueType::Main]; + if self.async_compute_queue.is_some() { + queues.push(crate::QueueType::AsyncCompute); + } + if self.async_transfer_queue.is_some() { + queues.push(crate::QueueType::AsyncTransfer); + } crate::Capabilities { binding_array: self.binding_array, ray_query: match self.device.ray_tracing { @@ -1481,6 +1496,7 @@ impl super::Context { dual_source_blending: self.dual_source_blending, shader_float16: self.shader_float16, cooperative_matrix: self.cooperative_matrix, + queues, } } diff --git a/examples/particle/main.rs b/examples/particle/main.rs index 9b7cb65b..006169e9 100644 --- a/examples/particle/main.rs +++ b/examples/particle/main.rs @@ -5,13 +5,20 @@ use blade_graphics as gpu; struct Example { command_encoder: gpu::CommandEncoder, compute_encoder: gpu::CommandEncoder, - prev_sync_point: gpu::SyncPoint, + prev_compute_sync: gpu::SyncPoint, + prev_render_sync: gpu::SyncPoint, context: gpu::Context, surface: gpu::Surface, gui_painter: blade_egui::GuiPainter, particle_pipeline: blade_particle::ParticlePipeline, - particle_system: blade_particle::ParticleSystem, + /// Double-buffered particle systems for async compute overlap. + /// `particle_systems[frame_index % 2]` is being computed, + /// while the other is being rendered (from the previous frame's compute). + particle_systems: [blade_particle::ParticleSystem; 2], + frame_index: usize, + async_compute: bool, + parallel_update: bool, time: f32, sample_count: u32, @@ -28,7 +35,7 @@ impl Example { let surface_info = self.surface.info(); - let _ = self.context.wait_for(&self.prev_sync_point, !0); + let _ = self.context.wait_for(&self.prev_render_sync, !0); if let Some(msaa_view) = self.msaa_view.take() { self.context.destroy_texture_view(msaa_view); @@ -111,6 +118,29 @@ impl Example { } } + fn make_effect() -> blade_particle::ParticleEffect { + blade_particle::ParticleEffect { + capacity: 1000_000, + emitter: blade_particle::Emitter { + rate: 6400.0, + burst_count: 0, + shape: blade_particle::EmitterShape::Point, + cone_angle: 0.5, + }, + particle: blade_particle::ParticleConfig { + life: [1.0, 5.0], + speed: [50.0, 250.0], + scale: [1.0, 15.0], + color: blade_particle::ColorConfig::Palette(vec![ + [255, 100, 50, 255], + [50, 200, 255, 255], + [255, 220, 60, 255], + [100, 255, 120, 255], + ]), + }, + } + } + fn new(window: &winit::window::Window) -> Self { let window_size = window.inner_size(); let context = unsafe { @@ -130,6 +160,9 @@ impl Example { let surface_info = surface.info(); let caps = context.capabilities(); + let async_compute = caps.queues.contains(&gpu::QueueType::AsyncCompute); + log::info!("Async compute: {async_compute}"); + let sample_count = [4, 2, 1] .into_iter() .find(|&n| (caps.sample_count_mask & n) != 0) @@ -146,48 +179,41 @@ impl Example { sample_count, }, ); - let effect = blade_particle::ParticleEffect { - capacity: 100_000, - emitter: blade_particle::Emitter { - rate: 6400.0, - burst_count: 0, - shape: blade_particle::EmitterShape::Point, - cone_angle: 0.5, - }, - particle: blade_particle::ParticleConfig { - life: [1.0, 5.0], - speed: [50.0, 250.0], - scale: [1.0, 15.0], - color: blade_particle::ColorConfig::Palette(vec![ - [255, 100, 50, 255], - [50, 200, 255, 255], - [255, 220, 60, 255], - [100, 255, 120, 255], - ]), - }, - }; - let particle_system = particle_pipeline.create_system(&context, "particle system", &effect); + let effect = Self::make_effect(); + let particle_systems = [ + particle_pipeline.create_system(&context, "particles-A", &effect), + particle_pipeline.create_system(&context, "particles-B", &effect), + ]; let command_encoder = context.create_command_encoder(gpu::CommandEncoderDesc { name: "main", buffer_count: 2, queue: gpu::QueueType::Main, }); + let compute_queue = if async_compute { + gpu::QueueType::AsyncCompute + } else { + gpu::QueueType::Main + }; let compute_encoder = context.create_command_encoder(gpu::CommandEncoderDesc { - name: "async-compute", + name: "compute", buffer_count: 2, - queue: gpu::QueueType::AsyncCompute, + queue: compute_queue, }); Self { command_encoder, compute_encoder, - prev_sync_point: gpu::SyncPoint::default(), + prev_compute_sync: gpu::SyncPoint::default(), + prev_render_sync: gpu::SyncPoint::default(), context, surface, gui_painter, particle_pipeline, - particle_system, + particle_systems, + frame_index: 0, + async_compute, + parallel_update: async_compute, time: 0.0, sample_count, msaa_texture: None, @@ -197,13 +223,16 @@ impl Example { } fn destroy(&mut self) { - let _ = self.context.wait_for(&self.prev_sync_point, !0); + let _ = self.context.wait_for(&self.prev_render_sync, !0); + let _ = self.context.wait_for(&self.prev_compute_sync, !0); self.context .destroy_command_encoder(&mut self.command_encoder); self.context .destroy_command_encoder(&mut self.compute_encoder); self.gui_painter.destroy(&self.context); - self.particle_system.destroy(&self.context); + for ps in &mut self.particle_systems { + ps.destroy(&self.context); + } self.particle_pipeline.destroy(&self.context); self.context.destroy_surface(&mut self.surface); @@ -234,49 +263,13 @@ impl Example { } } - fn render( + fn render_particles( &mut self, + frame_view: gpu::TextureView, + draw_idx: usize, gui_primitives: &[egui::ClippedPrimitive], - gui_textures: &egui::TexturesDelta, screen_desc: &blade_egui::ScreenDescriptor, ) { - self.recreate_msaa_texutres_if_needed( - screen_desc.physical_size, - self.surface.info().format, - ); - - let frame = self.surface.acquire_frame(); - let frame_view = frame.texture_view(); - - // Orbit the emitter and rotate its axis - self.time += 0.01; - let orbit_radius = 60.0; - self.particle_system.origin = [ - orbit_radius * self.time.cos(), - orbit_radius * self.time.sin(), - 0.0, - ]; - let spray_angle = self.time * 3.0; - self.particle_system.axis = [spray_angle.cos(), spray_angle.sin(), 0.0]; - - // Submit particle update on async compute queue. - // Wait for previous frame's render to finish before reusing particle buffers. - let prev_sp = self.prev_sync_point.clone(); - self.compute_encoder.start(); - self.particle_system - .update(&self.particle_pipeline, &mut self.compute_encoder, 0.01); - let sp_compute = self.context.submit(&mut self.compute_encoder, &[prev_sp]); - - // Record rendering on main queue, waiting for compute to finish. - self.command_encoder.start(); - if let Some(msaa_texture) = self.msaa_texture { - self.command_encoder.init_texture(msaa_texture); - } - self.command_encoder.init_texture(frame.texture()); - - self.gui_painter - .update_textures(&mut self.command_encoder, gui_textures, &self.context); - let camera = self.make_camera(screen_desc.physical_size); if self.sample_count <= 1 && !self.export_image { @@ -291,8 +284,7 @@ impl Example { depth_stencil: None, }, ) { - self.particle_system - .draw(&self.particle_pipeline, &mut pass, &camera); + self.particle_systems[draw_idx].draw(&self.particle_pipeline, &mut pass, &camera); self.gui_painter .paint(&mut pass, gui_primitives, screen_desc, &self.context); } @@ -312,8 +304,7 @@ impl Example { depth_stencil: None, }, ) { - self.particle_system - .draw(&self.particle_pipeline, &mut pass, &camera); + self.particle_systems[draw_idx].draw(&self.particle_pipeline, &mut pass, &camera); } if let mut pass = self.command_encoder.render( "draw ui", @@ -330,30 +321,151 @@ impl Example { .paint(&mut pass, gui_primitives, screen_desc, &self.context); } } + } + + fn render( + &mut self, + gui_primitives: &[egui::ClippedPrimitive], + gui_textures: &egui::TexturesDelta, + screen_desc: &blade_egui::ScreenDescriptor, + ) { + profiling::scope!("frame"); + self.recreate_msaa_texutres_if_needed( + screen_desc.physical_size, + self.surface.info().format, + ); + + let frame = self.surface.acquire_frame(); + let frame_view = frame.texture_view(); + + // Orbit the emitter and rotate its axis + self.time += 0.01; + let orbit_radius = 60.0; + let origin = [ + orbit_radius * self.time.cos(), + orbit_radius * self.time.sin(), + 0.0, + ]; + let spray_angle = self.time * 3.0; + let axis = [spray_angle.cos(), spray_angle.sin(), 0.0]; + + // Pick which particle system to compute and which to draw. + // In parallel mode, they are different (double-buffered). + // In serial mode, both are the same (index 0). + let compute_idx = if self.parallel_update { + self.frame_index % 2 + } else { + 0 + }; + let draw_idx = if self.parallel_update { + (self.frame_index + 1) % 2 + } else { + 0 + }; - self.command_encoder.present(frame); - let sync_point = self - .context - .submit(&mut self.command_encoder, &[sp_compute]); - self.gui_painter.after_submit(&sync_point); + // Update emitter state on the system we're about to compute + self.particle_systems[compute_idx].origin = origin; + self.particle_systems[compute_idx].axis = axis; + + if self.parallel_update { + // === PARALLEL PATH === + // Compute and render touch different buffers, so they overlap on GPU. + + // Submit compute on async queue. Only waits for previous compute. + let compute_sync = { + profiling::scope!("submit compute"); + self.compute_encoder.start(); + self.particle_systems[compute_idx].update( + &self.particle_pipeline, + &mut self.compute_encoder, + 0.01, + ); + self.context + .submit(&mut self.compute_encoder, &[self.prev_compute_sync.clone()]) + }; + + // Submit render on main queue. Only waits for previous render. + let render_sync = { + profiling::scope!("submit render"); + self.command_encoder.start(); + if let Some(msaa_texture) = self.msaa_texture { + self.command_encoder.init_texture(msaa_texture); + } + self.command_encoder.init_texture(frame.texture()); + self.gui_painter.update_textures( + &mut self.command_encoder, + gui_textures, + &self.context, + ); + self.render_particles(frame_view, draw_idx, gui_primitives, screen_desc); + self.command_encoder.present(frame); + self.context + .submit(&mut self.command_encoder, &[self.prev_render_sync.clone()]) + }; + + self.gui_painter.after_submit(&render_sync); + let _ = self.context.wait_for(&self.prev_compute_sync, !0); + let _ = self.context.wait_for(&self.prev_render_sync, !0); + self.prev_compute_sync = compute_sync; + self.prev_render_sync = render_sync; + } else { + // === SERIAL PATH === + // Everything on the main encoder, single particle system. + + self.command_encoder.start(); + if let Some(msaa_texture) = self.msaa_texture { + self.command_encoder.init_texture(msaa_texture); + } + self.command_encoder.init_texture(frame.texture()); + self.gui_painter.update_textures( + &mut self.command_encoder, + gui_textures, + &self.context, + ); + self.particle_systems[0].update( + &self.particle_pipeline, + &mut self.command_encoder, + 0.01, + ); + self.render_particles(frame_view, 0, gui_primitives, screen_desc); + self.command_encoder.present(frame); + let render_sync = self + .context + .submit(&mut self.command_encoder, &[self.prev_render_sync.clone()]); + + self.gui_painter.after_submit(&render_sync); + let _ = self.context.wait_for(&self.prev_render_sync, !0); + self.prev_render_sync = render_sync; + } - let _ = self.context.wait_for(&self.prev_sync_point, !0); - self.prev_sync_point = sync_point; + self.frame_index += 1; + profiling::finish_frame!(); } fn add_gui(&mut self, ui: &mut egui::Ui) { ui.heading("Particle System"); - let p = &mut self.particle_system.effect.particle; + let compute_idx = self.frame_index % 2; + let p = &mut self.particle_systems[compute_idx].effect.particle; ui.add(egui::Slider::new(&mut p.life[1], 0.1..=20.0).text("life")); ui.add(egui::Slider::new(&mut p.speed[1], 1.0..=500.0).text("speed")); ui.add(egui::Slider::new(&mut p.scale[1], 0.1..=70.0).text("scale")); ui.add( - egui::Slider::new(&mut self.particle_system.effect.emitter.rate, 0.0..=20000.0) - .text("rate"), + egui::Slider::new( + &mut self.particle_systems[compute_idx].effect.emitter.rate, + 0.0..=20000.0, + ) + .text("rate"), ); ui.add_space(5.0); ui.heading("Rendering Settings"); + ui.add_enabled( + self.async_compute, + egui::Checkbox::new(&mut self.parallel_update, "Parallel update"), + ); + if !self.async_compute { + ui.label("(no dedicated compute queue)"); + } egui::ComboBox::new("msaa dropdown", "MSAA samples") .selected_text(format!("x{}", self.sample_count)) .show_ui(ui, |ui| { @@ -367,10 +479,13 @@ impl Example { .selectable_value(&mut self.sample_count, i, format!("x{i}")) .changed() { - let _ = self.context.wait_for(&self.prev_sync_point, !0); + let _ = self.context.wait_for(&self.prev_render_sync, !0); + let _ = self.context.wait_for(&self.prev_compute_sync, !0); - let old_effect = self.particle_system.effect.clone(); - self.particle_system.destroy(&self.context); + let effect = Self::make_effect(); + for ps in &mut self.particle_systems { + ps.destroy(&self.context); + } self.particle_pipeline.destroy(&self.context); self.particle_pipeline = blade_particle::ParticlePipeline::new( &self.context, @@ -381,11 +496,18 @@ impl Example { sample_count: self.sample_count, }, ); - self.particle_system = self.particle_pipeline.create_system( - &self.context, - "particle system", - &old_effect, - ); + self.particle_systems = [ + self.particle_pipeline.create_system( + &self.context, + "particles-A", + &effect, + ), + self.particle_pipeline.create_system( + &self.context, + "particles-B", + &effect, + ), + ]; if let Some(msaa_view) = self.msaa_view.take() { self.context.destroy_texture_view(msaa_view); @@ -398,7 +520,15 @@ impl Example { }); ui.add_space(5.0); - ui.heading("Timings"); + ui.heading("Timings (compute)"); + for (name, time) in self.compute_encoder.timings() { + let millis = time.as_secs_f32() * 1000.0; + ui.horizontal(|ui| { + ui.label(name); + ui.colored_label(egui::Color32::WHITE, format!("{:.2} ms", millis)); + }); + } + ui.heading("Timings (render)"); for (name, time) in self.command_encoder.timings() { let millis = time.as_secs_f32() * 1000.0; ui.horizontal(|ui| { From a24f3b72f92ea2a8aff4f09818c9d6a20de66a84 Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Thu, 16 Apr 2026 00:18:18 -0700 Subject: [PATCH 4/4] Bar chart for the particle example performance --- examples/particle/main.rs | 131 ++++++++++++++++++++++++++++++++++---- 1 file changed, 117 insertions(+), 14 deletions(-) diff --git a/examples/particle/main.rs b/examples/particle/main.rs index 006169e9..c7de0657 100644 --- a/examples/particle/main.rs +++ b/examples/particle/main.rs @@ -1,6 +1,9 @@ #![allow(irrefutable_let_patterns)] use blade_graphics as gpu; +use std::collections::VecDeque; + +const TIMING_HISTORY_SIZE: usize = 120; struct Example { command_encoder: gpu::CommandEncoder, @@ -19,6 +22,7 @@ struct Example { frame_index: usize, async_compute: bool, parallel_update: bool, + timing_history: VecDeque<[f32; 2]>, time: f32, sample_count: u32, @@ -214,6 +218,7 @@ impl Example { frame_index: 0, async_compute, parallel_update: async_compute, + timing_history: VecDeque::with_capacity(TIMING_HISTORY_SIZE), time: 0.0, sample_count, msaa_texture: None, @@ -438,6 +443,24 @@ impl Example { self.prev_render_sync = render_sync; } + // Record timing history + let compute_ms: f32 = self + .compute_encoder + .timings() + .iter() + .map(|(_, d)| d.as_secs_f32() * 1000.0) + .sum(); + let render_ms: f32 = self + .command_encoder + .timings() + .iter() + .map(|(_, d)| d.as_secs_f32() * 1000.0) + .sum(); + if self.timing_history.len() >= TIMING_HISTORY_SIZE { + self.timing_history.pop_front(); + } + self.timing_history.push_back([compute_ms, render_ms]); + self.frame_index += 1; profiling::finish_frame!(); } @@ -520,21 +543,101 @@ impl Example { }); ui.add_space(5.0); - ui.heading("Timings (compute)"); - for (name, time) in self.compute_encoder.timings() { - let millis = time.as_secs_f32() * 1000.0; - ui.horizontal(|ui| { - ui.label(name); - ui.colored_label(egui::Color32::WHITE, format!("{:.2} ms", millis)); - }); + ui.heading("GPU Timings"); + // Current frame numbers + let compute_ms: f32 = self + .compute_encoder + .timings() + .iter() + .map(|(_, d)| d.as_secs_f32() * 1000.0) + .sum(); + let render_ms: f32 = self + .command_encoder + .timings() + .iter() + .map(|(_, d)| d.as_secs_f32() * 1000.0) + .sum(); + if self.parallel_update { + ui.label(format!( + "compute: {compute_ms:.2} ms render: {render_ms:.2} ms" + )); + } else { + ui.label(format!("frame: {:.2} ms", compute_ms + render_ms)); } - ui.heading("Timings (render)"); - for (name, time) in self.command_encoder.timings() { - let millis = time.as_secs_f32() * 1000.0; - ui.horizontal(|ui| { - ui.label(name); - ui.colored_label(egui::Color32::WHITE, format!("{:.2} ms", millis)); - }); + + // Stacked bar chart of timing history + let plot_height = 80.0; + let (rect, _response) = ui.allocate_exact_size( + egui::vec2(ui.available_width(), plot_height), + egui::Sense::hover(), + ); + let painter = ui.painter_at(rect); + painter.rect_filled(rect, 0.0, egui::Color32::from_gray(20)); + + if !self.timing_history.is_empty() { + let max_ms = self + .timing_history + .iter() + .map(|[c, r]| c + r) + .fold(0.1_f32, f32::max); + let bar_w = rect.width() / TIMING_HISTORY_SIZE as f32; + let offset = TIMING_HISTORY_SIZE - self.timing_history.len(); + for (i, &[c_ms, r_ms]) in self.timing_history.iter().enumerate() { + let x = rect.left() + (offset + i) as f32 * bar_w; + // Compute bar (bottom) + let c_h = (c_ms / max_ms) * rect.height(); + let c_rect = egui::Rect::from_min_size( + egui::pos2(x, rect.bottom() - c_h), + egui::vec2(bar_w - 1.0, c_h), + ); + painter.rect_filled(c_rect, 0.0, egui::Color32::from_rgb(100, 200, 255)); + // Render bar (stacked on top) + let r_h = (r_ms / max_ms) * rect.height(); + let r_rect = egui::Rect::from_min_size( + egui::pos2(x, rect.bottom() - c_h - r_h), + egui::vec2(bar_w - 1.0, r_h), + ); + painter.rect_filled(r_rect, 0.0, egui::Color32::from_rgb(255, 150, 50)); + } + // Legend + let legend_y = rect.top() + 2.0; + painter.text( + egui::pos2(rect.left() + 4.0, legend_y), + egui::Align2::LEFT_TOP, + format!("{max_ms:.1} ms"), + egui::FontId::proportional(10.0), + egui::Color32::GRAY, + ); + painter.rect_filled( + egui::Rect::from_min_size( + egui::pos2(rect.right() - 90.0, legend_y), + egui::vec2(8.0, 8.0), + ), + 0.0, + egui::Color32::from_rgb(255, 150, 50), + ); + painter.text( + egui::pos2(rect.right() - 78.0, legend_y), + egui::Align2::LEFT_TOP, + "render", + egui::FontId::proportional(10.0), + egui::Color32::GRAY, + ); + painter.rect_filled( + egui::Rect::from_min_size( + egui::pos2(rect.right() - 44.0, legend_y), + egui::vec2(8.0, 8.0), + ), + 0.0, + egui::Color32::from_rgb(100, 200, 255), + ); + painter.text( + egui::pos2(rect.right() - 32.0, legend_y), + egui::Align2::LEFT_TOP, + "compute", + egui::FontId::proportional(10.0), + egui::Color32::GRAY, + ); } } }