Bind entire metadata buffer with dynamic indexing (#503)

Remove the individual binding of a single `GpuEffectMetadata` entry per
effect, and instead bind the entire buffer. Use dynamic indexing from
the `GpuSpawner` to fetch the effect's metadata entry in the buffer.

This allows collapsing all bind groups using the same buffer into a
single one, and lays some foundations for batching effects.

Author: djeedai

src/lib.rs

@@ -949,13 +949,12 @@ impl EffectShaderSource {
                 "@group(3) @binding({binding_index}) var<storage, read_write> event_buffer_{i} : EventBuffer;\n"));
             emit_event_buffer_append_funcs_code.push_str(&format!(
                 r##"/// Append one or more spawn events to the event buffer.
-fn append_spawn_events_{0}(particle_index: u32, count: u32) {{
+fn append_spawn_events_{0}(base_child_index: u32, particle_index: u32, count: u32) {{
     // Optimize this case.
     if (count == 0u) {{
         return;
     }}
 
-    let base_child_index = effect_metadata.base_child_index;
     let capacity = arrayLength(&event_buffer_{0}.spawn_events);
     let base = min(u32(atomicAdd(&child_info_buffer.rows[base_child_index + {0}].event_count, i32(count))), capacity);
     let capped_count = min(count, capacity - base);

src/modifier/mod.rs

@@ -674,11 +674,11 @@ impl EmitSpawnEventModifier {
 
         let cond = match self.condition {
             EventEmitCondition::Always => format!(
-                "if (is_alive) {{ append_spawn_events_{channel_index}(particle_index, {}); }}",
+                "if (is_alive) {{ append_spawn_events_{channel_index}((*effect_metadata).base_child_index, particle_index, {}); }}",
                 count_var
             ),
             EventEmitCondition::OnDie => format!(
-                "if (was_alive && !is_alive) {{ append_spawn_events_{channel_index}(particle_index, {}); }}",
+                "if (was_alive && !is_alive) {{ append_spawn_events_{channel_index}((*effect_metadata).base_child_index, particle_index, {}); }}",
                 count_var
             ),
         };

src/render/mod.rs

@@ -2817,11 +2817,6 @@ impl GpuLimits {
     pub fn effect_metadata_offset(&self, buffer_index: u32) -> u64 {
         self.effect_metadata_aligned_size.get() as u64 * buffer_index as u64
     }
-
-    /// Byte alignment for [`GpuEffectMetadata`].
-    pub fn effect_metadata_size(&self) -> NonZeroU64 {
-        NonZeroU64::new(self.effect_metadata_aligned_size.get() as u64).unwrap()
-    }
 }
 
 /// Global render world resource containing the GPU data to draw all the
@@ -4838,15 +4833,13 @@ impl From<&ConsumeEventBuffers<'_>> for ConsumeEventKey {
 struct InitMetadataBindGroupKey {
     pub slab_id: SlabId,
     pub effect_metadata_buffer: BufferId,
-    pub effect_metadata_offset: u32,
     pub consume_event_key: Option<ConsumeEventKey>,
 }
 
 #[derive(Debug, Clone, PartialEq, Eq, Hash)]
 struct UpdateMetadataBindGroupKey {
     pub slab_id: SlabId,
     pub effect_metadata_buffer: BufferId,
-    pub effect_metadata_offset: u32,
     pub child_info_buffer_id: Option<BufferId>,
     pub event_buffers_keys: Vec<BindingKey>,
 }
@@ -4942,33 +4935,29 @@ impl EffectBindGroups {
     pub(self) fn get_or_create_init_metadata(
         &mut self,
         effect_batch: &EffectBatch,
-        gpu_limits: &GpuLimits,
         render_device: &RenderDevice,
         layout: &BindGroupLayout,
         effect_metadata_buffer: &Buffer,
         consume_event_buffers: Option<ConsumeEventBuffers>,
     ) -> Result<&BindGroup, ()> {
         assert!(effect_batch.metadata_table_id.is_valid());
 
-        let effect_metadata_offset =
-            gpu_limits.effect_metadata_offset(effect_batch.metadata_table_id.0) as u32;
         let key = InitMetadataBindGroupKey {
             slab_id: effect_batch.slab_id,
             effect_metadata_buffer: effect_metadata_buffer.id(),
-            effect_metadata_offset,
             consume_event_key: consume_event_buffers.as_ref().map(Into::into),
         };
 
         let make_entry = || {
             let mut entries = Vec::with_capacity(3);
             entries.push(
-                // @group(3) @binding(0) var<storage, read_write> effect_metadata : EffectMetadata;
+                // @group(3) @binding(0) var<storage, read_write> effect_metadatas : array<EffectMetadata>;
                 BindGroupEntry {
                     binding: 0,
                     resource: BindingResource::Buffer(BufferBinding {
                         buffer: effect_metadata_buffer,
-                        offset: key.effect_metadata_offset as u64,
-                        size: Some(gpu_limits.effect_metadata_size()),
+                        offset: 0,
+                        size: None,
                     }),
                 },
             );
@@ -5038,7 +5027,6 @@ impl EffectBindGroups {
     pub(self) fn get_or_create_update_metadata(
         &mut self,
         effect_batch: &EffectBatch,
-        gpu_limits: &GpuLimits,
         render_device: &RenderDevice,
         layout: &BindGroupLayout,
         effect_metadata_buffer: &Buffer,
@@ -5067,23 +5055,20 @@ impl EffectBindGroups {
         let key = UpdateMetadataBindGroupKey {
             slab_id: effect_batch.slab_id,
             effect_metadata_buffer: effect_metadata_buffer.id(),
-            effect_metadata_offset: gpu_limits
-                .effect_metadata_offset(effect_batch.metadata_table_id.0)
-                as u32,
             child_info_buffer_id,
             event_buffers_keys,
         };
 
         let make_entry = || {
             let mut entries = Vec::with_capacity(2 + event_buffers.len());
-            // @group(3) @binding(0) var<storage, read_write> effect_metadata :
-            // EffectMetadata;
+            // @group(3) @binding(0) var<storage, read_write> effect_metadatas :
+            // array<EffectMetadata>;
             entries.push(BindGroupEntry {
                 binding: 0,
                 resource: BindingResource::Buffer(BufferBinding {
                     buffer: effect_metadata_buffer,
-                    offset: key.effect_metadata_offset as u64,
-                    size: Some(gpu_limits.effect_metadata_aligned_size.into()),
+                    offset: 0,
+                    size: None,
                 }),
             });
             if emits_gpu_spawn_events {
@@ -6372,7 +6357,6 @@ pub(crate) fn prepare_bind_groups(
             if effect_bind_groups
                 .get_or_create_init_metadata(
                     effect_batch,
-                    &effects_meta.gpu_limits,
                     &render_device,
                     init_metadata_layout,
                     effects_meta.effect_metadata_buffer.buffer().unwrap(),
@@ -6397,7 +6381,6 @@ pub(crate) fn prepare_bind_groups(
             if effect_bind_groups
                 .get_or_create_update_metadata(
                     effect_batch,
-                    &effects_meta.gpu_limits,
                     &render_device,
                     update_metadata_layout,
                     effects_meta.effect_metadata_buffer.buffer().unwrap(),
@@ -7242,16 +7225,16 @@ impl Node for VfxSimulateNode {
                 }
 
                 // Compute dynamic offsets
-                let spawner_index = effect_batch.spawner_base;
+                let spawner_base = effect_batch.spawner_base;
                 let spawner_aligned_size = effects_meta.spawner_buffer.aligned_size();
                 assert!(spawner_aligned_size >= GpuSpawnerParams::min_size().get() as usize);
-                let spawner_offset = spawner_index * spawner_aligned_size as u32;
+                let spawner_offset = spawner_base * spawner_aligned_size as u32;
                 let property_offset = effect_batch.property_offset;
 
                 trace!(
                     "record commands for update pipeline of effect {:?} spawner_base={}",
                     effect_batch.handle,
-                    spawner_index,
+                    spawner_base,
                 );
 
                 // Setup update pass

src/render/vfx_init.wgsl

@@ -42,7 +42,7 @@ struct ParentParticleBuffer {
 {{PROPERTIES_BINDING}}
 
 // "metadata" group @3
-@group(3) @binding(0) var<storage, read_write> effect_metadata : EffectMetadata;
+@group(3) @binding(0) var<storage, read_write> effect_metadatas : array<EffectMetadata>;
 #ifdef CONSUME_GPU_SPAWN_EVENTS
 @group(3) @binding(1) var<storage, read> child_info_buffer : ChildInfoBuffer;
 @group(3) @binding(2) var<storage, read> event_buffer : EventBuffer;
@@ -56,7 +56,8 @@ fn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
 
     // Cap to max number of dead particles, copied from (capacity - alive_count) at the end
     // of the previous iteration, and constant during this pass (unlike alive_count).
-    let max_spawn = atomicLoad(&effect_metadata.max_spawn);
+    let effect_metadata = &effect_metadatas[spawner.effect_metadata_index];
+    let max_spawn = atomicLoad(&(*effect_metadata).max_spawn);
     if (thread_index >= max_spawn) {
         return;
     }
@@ -65,7 +66,7 @@ fn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
     // in workgroup_size(64) so more threads than needed are launched (rounded up to 64).
 #ifdef CONSUME_GPU_SPAWN_EVENTS
     let event_index = thread_index;
-    let global_child_index = effect_metadata.global_child_index;
+    let global_child_index = (*effect_metadata).global_child_index;
     let event_count = child_info_buffer.rows[global_child_index].event_count;
     if (event_index >= u32(event_count)) {
         return;
@@ -81,14 +82,14 @@ fn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
 #endif
 
     // Count as alive, and recycle a dead particle slot to store the newly spawned particle
-    let alive_index = atomicAdd(&effect_metadata.alive_count, 1u);
-    let dead_index = effect_metadata.capacity - alive_index - 1u;
+    let alive_index = atomicAdd(&(*effect_metadata).alive_count, 1u);
+    let dead_index = (*effect_metadata).capacity - alive_index - 1u;
     let particle_index = indirect_buffer.rows[dead_index].dead_index;
 
     // Bump the particle counter each time we allocate a particle. This generates a unique
     // particle ID used for various purposes (but not directly by the simulation). We still
     // store it in a variable, because the INIT_CODE might access it.
-    let particle_counter = atomicAdd(&effect_metadata.particle_counter, 1u);
+    let particle_counter = atomicAdd(&(*effect_metadata).particle_counter, 1u);
 
     // Initialize the PRNG seed
     seed = pcg_hash(particle_index ^ spawner.seed);
@@ -127,7 +128,7 @@ fn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
 #endif
 
     // Append to alive list of indirect buffer.
-    let write_index = effect_metadata.indirect_write_index;
+    let write_index = (*effect_metadata).indirect_write_index;
     indirect_buffer.rows[alive_index].particle_index[write_index] = particle_index;
 
     // Write back new particle

src/render/vfx_update.wgsl

@@ -43,7 +43,7 @@ struct ParentParticleBuffer {
 {{PROPERTIES_BINDING}}
 
 // "metadata" group @3
-@group(3) @binding(0) var<storage, read_write> effect_metadata : EffectMetadata;
+@group(3) @binding(0) var<storage, read_write> effect_metadatas : array<EffectMetadata>;
 #ifdef EMITS_GPU_SPAWN_EVENTS
 {{EMIT_EVENT_BUFFER_BINDINGS}}
 #endif
@@ -59,7 +59,8 @@ fn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
     let thread_index = global_invocation_id.x;
 
     // Cap at maximum number of alive particles.
-    if (thread_index >= effect_metadata.max_update) {
+    let effect_metadata = &effect_metadatas[spawner.effect_metadata_index];
+    if (thread_index >= (*effect_metadata).max_update) {
         return;
     }
 
@@ -83,18 +84,18 @@ fn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
     // Check if alive
     if (!is_alive) {
         // Save dead index
-        let alive_index = atomicSub(&effect_metadata.alive_count, 1u);
-        let dead_index = effect_metadata.capacity - alive_index;
+        let alive_index = atomicSub(&((*effect_metadata).alive_count), 1u);
+        let dead_index = (*effect_metadata).capacity - alive_index;
         indirect_buffer.rows[dead_index].dead_index = particle_index;
 
         // Also increment copy of dead count, which was updated in dispatch indirect
         // pass just before, and need to remain correct after this pass. We wouldn't have
         // to do that here if we had a per-effect pass between update and the next init.
-        atomicAdd(&effect_metadata.max_spawn, 1u);
+        atomicAdd(&((*effect_metadata).max_spawn), 1u);
     } else {
         // Increment visible particle count (in the absence of any GPU culling), and write
         // the indirection index for later rendering.
-        let indirect_index = atomicAdd(&draw_indirect_buffer[effect_metadata.indirect_render_index].instance_count, 1u);
+        let indirect_index = atomicAdd(&draw_indirect_buffer[(*effect_metadata).indirect_render_index].instance_count, 1u);
         indirect_buffer.rows[indirect_index].particle_index[write_index] = particle_index;
     }
 }