Directly use DispatchIndirectArgs in shaders (#506)

Switch to directly using an `array<DispatchIndirectArgs>` in the
`vfx_indirect` compute pass, instead of a `u32` array, since those
structs are already tightly packed. This makes the shader more readable.

Author: djeedai

src/render/mod.rs

@@ -2885,10 +2885,10 @@ impl EffectsMeta {
 
         // Ensure individual GpuSpawnerParams elements are properly aligned so they can
         // be addressed individually by the computer shaders.
-        let item_align = gpu_limits.storage_buffer_align().get() as u64;
+        let item_align = gpu_limits.storage_buffer_align();
         trace!(
             "Aligning storage buffers to {} bytes as device limits requires.",
-            item_align
+            item_align.get()
         );
 
         Self {
@@ -2900,7 +2900,7 @@ impl EffectsMeta {
             sim_params_uniforms: UniformBuffer::default(),
             spawner_buffer: AlignedBufferVec::new(
                 BufferUsages::STORAGE,
-                NonZeroU64::new(item_align),
+                Some(item_align.into()),
                 Some("hanabi:buffer:spawner".to_string()),
             ),
             dispatch_indirect_buffer: GpuBuffer::new(
@@ -2914,7 +2914,7 @@ impl EffectsMeta {
             ),
             effect_metadata_buffer: BufferTable::new(
                 BufferUsages::STORAGE | BufferUsages::INDIRECT,
-                Some(NonZeroU64::new(item_align).unwrap()),
+                Some(item_align.into()),
                 Some("hanabi:buffer:effect_metadata".to_string()),
             ),
             gpu_limits,
@@ -6180,7 +6180,7 @@ pub(crate) fn prepare_bind_groups(
                             resource: effect_metadata_buffer.as_entire_binding(),
                         },
                         // @group(1) @binding(1) var<storage, read_write> dispatch_indirect_buffer
-                        // : array<u32>;
+                        // : array<DispatchIndirectArgs>;
                         BindGroupEntry {
                             binding: 1,
                             resource: dispatch_indirect_buffer.as_entire_binding(),

src/render/vfx_common.wgsl

@@ -110,9 +110,6 @@ struct DispatchIndirectArgs {
     z: u32,
 }
 
-/// Stride in u32 count (4 bytes) of the DispatchIndirectArgs struct.
-const DISPATCH_INDIRECT_STRIDE: u32 = 3u;
-
 /// Indirect draw (non-indexed) dispatch struct for GPU-driven passes. The layout of this struct is dictated by WGSL.
 /// See https://docs.rs/wgpu/latest/wgpu/util/struct.DrawIndirectArgs.html.
 struct DrawIndirectArgs {
@@ -216,7 +213,6 @@ struct EffectMetadata {
     /// of an array, or sometimes individually as a single unit. In the later case,
     /// we need it to be aligned to the GPU limits of the device. That limit is only
     /// known at runtime when initializing the WebGPU device.
-    // FIXME - not anymore, but would be again with proper batching, so keep for now
     {{EFFECT_METADATA_PADDING}}
 }
 

src/render/vfx_indirect.wgsl

@@ -1,16 +1,15 @@
 #import bevy_hanabi::vfx_common::{
-    ChildInfo, ChildInfoBuffer, SimParams, Spawner,
+    ChildInfo, ChildInfoBuffer, DispatchIndirectArgs, SimParams, Spawner,
     EM_OFFSET_ALIVE_COUNT, EM_OFFSET_MAX_UPDATE, EM_OFFSET_CAPACITY,
     EM_OFFSET_MAX_SPAWN, EM_OFFSET_INDIRECT_DISPATCH_INDEX, DRAW_INDEXED_INDIRECT_STRIDE,
-    EM_OFFSET_INDIRECT_WRITE_INDEX, DISPATCH_INDIRECT_STRIDE, EFFECT_METADATA_STRIDE
+    EM_OFFSET_INDIRECT_WRITE_INDEX, EFFECT_METADATA_STRIDE
 }
 
 @group(0) @binding(0) var<uniform> sim_params : SimParams;
 
 // Tightly packed array of EffectMetadata[], accessed as u32 array.
 @group(1) @binding(0) var<storage, read_write> effect_metadata_buffer : array<u32>;
-// Tightly packed array of DispatchIndirectArgs[], accessed as u32 array.
-@group(1) @binding(1) var<storage, read_write> dispatch_indirect_buffer : array<u32>;
+@group(1) @binding(1) var<storage, read_write> dispatch_indirect_buffer : array<DispatchIndirectArgs>;
 // Tightly packed array of DrawIndexedIndirectArgs[], accessed as u32 array. This can contain
 // some DrawIndirectArgs[] instead, but in that case the stride is adjusted so all rows have
 // the same size. Since we access the instance_count, which is at the same position in both,
@@ -70,10 +69,9 @@ fn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
     // Calculate the number of workgroups (thread groups) to dispatch for the update
     // pass, which is the number of alive particles rounded up to 64 (workgroup_size).
     let indirect_dispatch_index = effect_metadata_buffer[em_base + EM_OFFSET_INDIRECT_DISPATCH_INDEX];
-    let di_base = DISPATCH_INDIRECT_STRIDE * indirect_dispatch_index;
-    dispatch_indirect_buffer[di_base] = (alive_count + 63u) >> 6u;
-    dispatch_indirect_buffer[di_base + 1u] = 1u;
-    dispatch_indirect_buffer[di_base + 2u] = 1u;
+    dispatch_indirect_buffer[indirect_dispatch_index].x = (alive_count + 63u) >> 6u;
+    dispatch_indirect_buffer[indirect_dispatch_index].y = 1u;
+    dispatch_indirect_buffer[indirect_dispatch_index].z = 1u;
 
     // Swap ping/pong buffers. The update pass always writes into ping, and both the update
     // pass and the render pass always read from pong.