diff --git a/CHANGELOG.md b/CHANGELOG.md index 84eef1e8..006fac2d 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -3,6 +3,12 @@ The format is based on [Keep a Changelog](https://keepachangelog.com/en/1.0.0/), and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0.html). +## [Unreleased] + +### Changed + +- Moved most shared WGSL code into an import module `vfx_common.wgsl`. This requires using `naga_oil` for import resolution, which in turns means `naga` and `naga_oil` are now dependencies of `bevy_hanabi` itself. + ## [0.9.0] 2023-12-26 ### Added diff --git a/Cargo.toml b/Cargo.toml index 6a8b81b6..f6160e2d 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -1,6 +1,6 @@ [package] name = "bevy_hanabi" -version = "0.9.0" +version = "0.10.0-dev" authors = ["Jerome Humbert "] edition = "2021" description = "Hanabi GPU particle system for the Bevy game engine" @@ -38,6 +38,9 @@ ron = "0.8" bitflags = "2.3" typetag = "0.2" thiserror = "1.0" +# Same versions as Bevy 0.12 (bevy_render) +naga = "0.13" +naga_oil = "0.10" [dependencies.bevy] version = "0.12" @@ -50,8 +53,6 @@ all-features = true [dev-dependencies] # Same versions as Bevy 0.12 (bevy_render) wgpu = "0.17.1" -naga = "0.13" -naga_oil = "0.10" # For procedural texture generation in examples noise = "0.8" diff --git a/src/lib.rs b/src/lib.rs index 90920fde..7ae922e6 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -1773,6 +1773,9 @@ else { return c1; } "PARTICLE_SCREEN_SPACE_SIZE".into(), ShaderDefValue::Bool(true), ); + if name == "Update" { + shader_defs.insert("RI_MAX_SPAWN_ATOMIC".into(), ShaderDefValue::Bool(true)); + } let mut composer = Composer::default(); // Import bevy_render::view for the render shader @@ -1791,6 +1794,15 @@ else { return c1; } assert!(res.is_ok()); } + // Import bevy_hanabi::vfx_common + { + let min_storage_buffer_offset_alignment = 256usize; + let common_shader = + HanabiPlugin::make_common_shader(min_storage_buffer_offset_alignment); + let res = composer.add_composable_module((&common_shader).into()); + assert!(res.is_ok()); + } + match composer.make_naga_module(NagaModuleDescriptor { source: code, file_path: "init.wgsl", diff --git a/src/plugin.rs b/src/plugin.rs index 8bb8c6c6..9efbce01 100644 --- a/src/plugin.rs +++ b/src/plugin.rs @@ -21,8 +21,9 @@ use crate::{ render::{ extract_effect_events, extract_effects, prepare_effects, prepare_resources, queue_effects, DispatchIndirectPipeline, DrawEffects, EffectAssetEvents, EffectBindGroups, EffectSystems, - EffectsMeta, ExtractedEffects, ParticlesInitPipeline, ParticlesRenderPipeline, - ParticlesUpdatePipeline, ShaderCache, SimParams, VfxSimulateDriverNode, VfxSimulateNode, + EffectsMeta, ExtractedEffects, GpuSpawnerParams, ParticlesInitPipeline, + ParticlesRenderPipeline, ParticlesUpdatePipeline, ShaderCache, SimParams, + VfxSimulateDriverNode, VfxSimulateNode, }, spawn::{self, Random}, tick_spawners, update_properties_from_asset, ParticleEffect, RemovedEffectsEvent, Spawner, @@ -46,10 +47,39 @@ pub mod simulate_graph { } } +// {626E7AD3-4E54-487E-B796-9A90E34CC1EC} +const HANABI_COMMON_TEMPLATE_HANDLE: Handle = + Handle::weak_from_u128(0x626E7AD34E54487EB7969A90E34CC1ECu128); + /// Plugin to add systems related to Hanabi. #[derive(Debug, Clone, Copy)] pub struct HanabiPlugin; +impl HanabiPlugin { + /// Create the `vfx_common.wgsl` shader with proper alignment. + /// + /// This creates a new [`Shader`] from the `vfx_common.wgsl` code, by + /// applying the given alignment for storage buffers. This produces a shader + /// ready for the specific GPU device associated with that alignment. + pub(crate) fn make_common_shader(min_storage_buffer_offset_alignment: usize) -> Shader { + let spawner_padding_code = + GpuSpawnerParams::padding_code(min_storage_buffer_offset_alignment); + let common_code = include_str!("render/vfx_common.wgsl") + .replace("{{SPAWNER_PADDING}}", &spawner_padding_code); + Shader::from_wgsl( + common_code, + std::path::Path::new(file!()) + .parent() + .unwrap() + .join(format!( + "render/vfx_common_{}.wgsl", + min_storage_buffer_offset_alignment + )) + .to_string_lossy(), + ) + } +} + impl Plugin for HanabiPlugin { fn build(&self, app: &mut App) { // Register asset @@ -94,8 +124,7 @@ impl Plugin for HanabiPlugin { let render_device = app .sub_app(RenderApp) .world - .get_resource::() - .unwrap() + .resource::() .clone(); let adapter_name = app @@ -113,6 +142,17 @@ impl Plugin for HanabiPlugin { info!("Initializing Hanabi for GPU adapter {}", adapter_name); } + // Insert the properly aligned `vfx_common.wgsl` shader into Assets, so + // that the automated Bevy shader processing finds it as an import. This is used + // for init/update/render shaders (but not the indirect one). + { + let common_shader = HanabiPlugin::make_common_shader( + render_device.limits().min_storage_buffer_offset_alignment as usize, + ); + let mut assets = app.world.resource_mut::>(); + assets.insert(HANABI_COMMON_TEMPLATE_HANDLE, common_shader); + } + let effects_meta = EffectsMeta::new(render_device); // Register the custom render pipeline diff --git a/src/render/mod.rs b/src/render/mod.rs index fd6cdc3f..70a8edb5 100644 --- a/src/render/mod.rs +++ b/src/render/mod.rs @@ -28,6 +28,7 @@ use bevy::{ utils::HashMap, }; use bitflags::bitflags; +use naga_oil::compose::{Composer, NagaModuleDescriptor}; use rand::random; use std::marker::PhantomData; use std::{ @@ -46,8 +47,8 @@ use crate::{ next_multiple_of, render::batch::{BatchInput, BatchState, Batcher, EffectBatch}, spawn::EffectSpawner, - CompiledParticleEffect, EffectProperties, EffectShader, ParticleLayout, PropertyLayout, - RemovedEffectsEvent, SimulationCondition, SimulationSpace, + CompiledParticleEffect, EffectProperties, EffectShader, HanabiPlugin, ParticleLayout, + PropertyLayout, RemovedEffectsEvent, SimulationCondition, SimulationSpace, }; mod aligned_buffer_vec; @@ -246,7 +247,7 @@ impl From<&Mat4> for GpuCompressedTransform { /// together form the spawner parameter buffer. #[repr(C)] #[derive(Debug, Default, Clone, Copy, Pod, Zeroable, ShaderType)] -struct GpuSpawnerParams { +pub(crate) struct GpuSpawnerParams { /// Transform of the effect, as a Mat4 without the last row (which is always /// (0,0,0,1) for an affine transform), stored transposed as a mat3x4 to /// avoid padding in WGSL. This is either added to emitted particles at @@ -269,6 +270,33 @@ struct GpuSpawnerParams { force_field: [GpuForceFieldSource; ForceFieldSource::MAX_SOURCES], } +impl GpuSpawnerParams { + /// Get the aligned size of this type based on the given alignment in bytes. + pub fn aligned_size(align_size: usize) -> usize { + next_multiple_of(GpuSpawnerParams::min_size().get() as usize, align_size) + } + + /// Get the WGSL padding code to append to the GPU struct to align it. + pub fn padding_code(align_size: usize) -> String { + let aligned_size = GpuSpawnerParams::aligned_size(align_size); + trace!( + "Aligning spawner params to {} bytes as device limits requires. Aligned size: {} bytes.", + align_size, + aligned_size + ); + + // We need to pad the Spawner WGSL struct based on the device padding so that we + // can use it as an array element but also has a direct struct binding. + if GpuSpawnerParams::min_size().get() as usize != aligned_size { + let padding_size = aligned_size - GpuSpawnerParams::min_size().get() as usize; + assert!(padding_size % 4 == 0); + format!("padding: array", padding_size / 4) + } else { + "".to_string() + } + } +} + // FIXME - min_storage_buffer_offset_alignment #[repr(C)] #[derive(Debug, Clone, Copy, Pod, Zeroable, ShaderType)] @@ -327,8 +355,7 @@ impl FromWorld for DispatchIndirectPipeline { // so needs the proper align. Because WGSL removed the @stride attribute, we pad // the WGSL type manually, so need to enforce min_binding_size everywhere. let item_align = render_device.limits().min_storage_buffer_offset_alignment as usize; - let spawner_aligned_size = - next_multiple_of(GpuSpawnerParams::min_size().get() as usize, item_align); + let spawner_aligned_size = GpuSpawnerParams::aligned_size(item_align); trace!( "Aligning spawner params to {} bytes as device limits requires. Size: {} bytes.", item_align, @@ -403,22 +430,49 @@ impl FromWorld for DispatchIndirectPipeline { // We need to pad the Spawner WGSL struct based on the device padding so that we // can use it as an array element but also has a direct struct binding. - let spawner_padding_code = if GpuSpawnerParams::min_size().get() as usize - != spawner_aligned_size - { - let padding_size = spawner_aligned_size - GpuSpawnerParams::min_size().get() as usize; - assert!(padding_size % 4 == 0); - format!("padding: array", padding_size / 4) - } else { - "".to_string() + let spawner_padding_code = GpuSpawnerParams::padding_code(item_align); + let indirect_code = + include_str!("vfx_indirect.wgsl").replace("{{SPAWNER_PADDING}}", &spawner_padding_code); + + // Resolve imports. Because we don't insert this shader into Bevy' pipeline + // cache, we don't get that part "for free", so we have to do it manually here. + let indirect_naga_module = { + let mut composer = Composer::default(); + + // Import bevy_hanabi::vfx_common + { + let common_shader = HanabiPlugin::make_common_shader(item_align); + let mut desc: naga_oil::compose::ComposableModuleDescriptor<'_> = + (&common_shader).into(); + desc.shader_defs.insert( + "SPAWNER_PADDING".to_string(), + naga_oil::compose::ShaderDefValue::Bool(true), + ); + let res = composer.add_composable_module(desc); + assert!(res.is_ok()); + } + + let shader_defs = default(); + + match composer.make_naga_module(NagaModuleDescriptor { + source: &indirect_code, + file_path: "vfx_indirect.wgsl", + shader_defs, + ..Default::default() + }) { + Ok(naga_module) => ShaderSource::Naga(Cow::Owned(naga_module)), + Err(compose_error) => panic!( + "Failed to compose vfx_indirect.wgsl, naga_oil returned: {}", + compose_error.emit_to_string(&composer) + ), + } }; - let indirect_code = include_str!("vfx_indirect.wgsl") - .to_string() - .replace("{{SPAWNER_PADDING}}", &spawner_padding_code); + debug!("Create indirect dispatch shader:\n{}", indirect_code); + let shader_module = render_device.create_shader_module(ShaderModuleDescriptor { label: Some("hanabi:vfx_indirect_shader"), - source: ShaderSource::Wgsl(Cow::Owned(indirect_code)), + source: indirect_naga_module, }); let pipeline = render_device.create_compute_pipeline(&RawComputePipelineDescriptor { @@ -772,7 +826,7 @@ impl SpecializedComputePipeline for ParticlesUpdatePipeline { self.render_indirect_layout.clone(), ], shader: key.shader, - shader_defs: vec![], + shader_defs: vec!["RI_MAX_SPAWN_ATOMIC".into()], entry_point: "main".into(), push_constant_ranges: Vec::new(), } @@ -2049,7 +2103,7 @@ pub(crate) struct BufferBindGroups { /// Bind group for the render graphic shader. /// /// ```wgsl - /// @binding(0) var particle_buffer : ParticlesBuffer; + /// @binding(0) var particle_buffer : ParticleBuffer; /// @binding(1) var indirect_buffer : IndirectBuffer; /// @binding(2) var dispatch_indirect : DispatchIndirect; /// ``` diff --git a/src/render/shader_cache.rs b/src/render/shader_cache.rs index 0154ae35..42a538c7 100644 --- a/src/render/shader_cache.rs +++ b/src/render/shader_cache.rs @@ -3,7 +3,7 @@ use std::hash::{Hash, Hasher}; use bevy::{ asset::{Assets, Handle}, ecs::{change_detection::ResMut, system::Resource}, - log::debug, + log::{debug, trace}, render::render_resource::Shader, utils::HashMap, }; @@ -40,10 +40,17 @@ impl ShaderCache { let mut hasher = bevy::utils::AHasher::default(); source.hash(&mut hasher); let hash = hasher.finish(); - let handle = shaders.add(Shader::from_wgsl( + let shader = Shader::from_wgsl( source.to_string(), format!("hanabi/{}_{}.wgsl", filename, hash), - )); + ); + trace!( + "Shader path={} import_path={:?} imports={:?}", + shader.path, + shader.import_path, + shader.imports + ); + let handle = shaders.add(shader); debug!("Inserted new configured shader: {:?}\n{}", handle, source); self.cache.insert(source.to_string(), handle.clone()); handle diff --git a/src/render/vfx_common.wgsl b/src/render/vfx_common.wgsl new file mode 100644 index 00000000..94c2196d --- /dev/null +++ b/src/render/vfx_common.wgsl @@ -0,0 +1,120 @@ +#define_import_path bevy_hanabi::vfx_common + +struct SimParams { + /// Delta time in seconds since last simulation tick. + delta_time: f32, + /// Time in seconds since the start of simulation. + time: f32, +//#ifdef SIM_PARAMS_INDIRECT_DATA + /// Number of effects batched together. + num_effects: u32, + /// Stride in bytes of the RenderIndirect struct. Used to calculate + /// the position of each effect's data into the buffer of a batch. + render_stride: u32, + /// Stride in bytes of the DispatchIndirect struct. Used to calculate + /// the position of each effect's data into the buffer of a batch. + dispatch_stride: u32, +//#endif +} + +struct ForceFieldSource { + position: vec3, + max_radius: f32, + min_radius: f32, + mass: f32, + force_exponent: f32, + conform_to_sphere: f32, +} + +struct Spawner { + transform: mat3x4, // transposed (row-major) + inverse_transform: mat3x4, // transposed (row-major) + spawn: i32, + seed: u32, + count: atomic, + effect_index: u32, + force_field: array, +#ifdef SPAWNER_PADDING + {{SPAWNER_PADDING}} +#endif +} + +struct IndirectBuffer { + indices: array, +} + +// Dispatch indirect array offsets. Used when accessing an array of DispatchIndirect +// as a raw array, so that we can avoid WGSL struct padding and keep data +// more compact in the render indirect buffer. Each offset corresponds to a field +// in the DispatchIndirect struct. +const DI_OFFSET_X: u32 = 0u; +const DI_OFFSET_Y: u32 = 1u; +const DI_OFFSET_Z: u32 = 2u; +const DI_OFFSET_PONG: u32 = 3u; + +/// Dispatch indirect parameters for GPU driven update compute. +struct DispatchIndirect { + x: u32, + y: u32, + z: u32, + /// Index of the ping-pong buffer of particle indices to read particles from + /// during rendering. Cached from RenderIndirect::ping after it's swapped + /// in the indirect dispatch, because the RenderIndirect struct is used by GPU + /// as an indirect draw source so cannot also be bound as regular storage + /// buffer for reading. + pong: u32, +} + +// Render indirect array offsets. Used when accessing an array of RenderIndirect +// as a raw array, so that we can avoid WGSL struct padding and keep data +// more compact in the render indirect buffer. Each offset corresponds to a field +// in the RenderIndirect struct. +const RI_OFFSET_VERTEX_COUNT: u32 = 0u; +const RI_OFFSET_INSTANCE_COUNT: u32 = 1u; +const RI_OFFSET_BASE_INDEX: u32 = 2u; +const RI_OFFSET_VERTEX_OFFSET: u32 = 3u; +const RI_OFFSET_BASE_INSTANCE: u32 = 4u; +const RI_OFFSET_ALIVE_COUNT: u32 = 5u; +const RI_OFFSET_DEAD_COUNT: u32 = 6u; +const RI_OFFSET_MAX_SPAWN: u32 = 7u; +const RI_OFFSET_PING: u32 = 8u; +const RI_OFFSET_MAX_UPDATE: u32 = 9u; + +/// Render indirect parameters for GPU driven rendering. +struct RenderIndirect { + /// Number of vertices in the particle mesh. Currently always 4 (quad mesh). + vertex_count: u32, + /// Number of mesh instances, equal to the number of particles. + instance_count: atomic, + /// Base index (always zero). + base_index: u32, + /// Vertex offset (always zero). + vertex_offset: i32, + /// Base instance (always zero). + base_instance: u32, + /// Number of particles alive after the init pass, used to calculate the number + /// of compute threads to spawn for the update pass and to cap those threads + /// via `max_update`. + alive_count: atomic, + /// Number of dead particles, decremented during the init pass as new particles + /// are spawned, and incremented during the update pass as existing particles die. + dead_count: atomic, + /// Maxmimum number of init threads to run on next frame. This is cached from + /// `dead_count` during the indirect dispatch of the previous frame, so that the + /// init compute pass can cap its thread count while also decrementing the actual + /// `dead_count` as particles are spawned. +#ifdef RI_MAX_SPAWN_ATOMIC + max_spawn: atomic, +#else + max_spawn: u32, +#endif + /// Index of the ping buffer for particle indices. Init and update compute passes + /// always write into the ping buffer and read from the pong buffer. The buffers + /// are swapped during the indirect dispatch. + ping: u32, + /// Maximum number of update threads to run. This is cached from `alive_count` + /// during the indirect dispatch, so that the update compute pass can cap its + /// thread count while also modifying the actual `alive_count` if some particle + /// dies during the update pass. + max_update: u32, +} diff --git a/src/render/vfx_indirect.wgsl b/src/render/vfx_indirect.wgsl index f8101fa5..cebb9173 100644 --- a/src/render/vfx_indirect.wgsl +++ b/src/render/vfx_indirect.wgsl @@ -1,51 +1,14 @@ - -struct SimParams { - delta_time: f32, - time: f32, - num_effects: u32, - render_stride: u32, - dispatch_stride: u32, -}; - -struct ForceFieldSource { - position: vec3, - max_radius: f32, - min_radius: f32, - mass: f32, - force_exponent: f32, - conform_to_sphere: f32, -}; - -struct Spawner { - transform: mat3x4, // transposed (row-major) - inverse_transform: mat3x4, // transposed (row-major) - spawn: i32, - seed: u32, - count: atomic, - effect_index: u32, - force_field: array, - {{SPAWNER_PADDING}} -}; +#import bevy_hanabi::vfx_common::{ + SimParams, Spawner, + DI_OFFSET_X, DI_OFFSET_PONG, + RI_OFFSET_ALIVE_COUNT, RI_OFFSET_MAX_UPDATE, RI_OFFSET_DEAD_COUNT, + RI_OFFSET_MAX_SPAWN, RI_OFFSET_INSTANCE_COUNT, RI_OFFSET_PING +} struct SpawnerBuffer { spawners: array, } -// naga doesn't support 'const' yet -// https://github.com/gfx-rs/naga/issues/1829 - -// const OFFSET_INSTANCE_COUNT: u32 = 1u; -// const OFFSET_ALIVE_COUNT: u32 = 5u; -// const OFFSET_DEAD_COUNT: u32 = 6u; -// const OFFSET_MAX_SPAWN: u32 = 7u -// const OFFSET_PING: u32 = 8u; -// const OFFSET_MAX_UPDATE: u32 = 9u; - -// const OFFSET_X: u32 = 0u; -// const OFFSET_Y: u32 = 1u; -// const OFFSET_Z: u32 = 2u; -// const OFFSET_PONG: u32 = 3u; - @group(0) @binding(0) var render_indirect_buffer : array; @group(0) @binding(1) var dispatch_indirect : array; @group(0) @binding(2) var spawner_buffer : SpawnerBuffer; @@ -76,31 +39,31 @@ fn main(@builtin(global_invocation_id) global_invocation_id: vec3) { // Calculate the number of thread groups to dispatch for the update pass, which is // the number of alive particles rounded up to 64 (workgroup_size). - let alive_count = render_indirect_buffer[ri_base + 5u]; - dispatch_indirect[di_base + 0u] = (alive_count + 63u) / 64u; + let alive_count = render_indirect_buffer[ri_base + RI_OFFSET_ALIVE_COUNT]; + dispatch_indirect[di_base + DI_OFFSET_X] = (alive_count + 63u) / 64u; // Update max_update from current value of alive_count, so that the update pass // coming next can cap its threads to this value, while also atomically modifying // alive_count itself for next frame. - render_indirect_buffer[ri_base + 9u] = alive_count; + render_indirect_buffer[ri_base + RI_OFFSET_MAX_UPDATE] = alive_count; // Copy the number of dead particles to a constant location, so that the init pass // on next frame can atomically modify dead_count in parallel yet still read its // initial value at the beginning of the init pass, and limit the number of particles // spawned to the number of dead particles to recycle. - let dead_count = render_indirect_buffer[ri_base + 6u]; - render_indirect_buffer[ri_base + 7u] = dead_count; + let dead_count = render_indirect_buffer[ri_base + RI_OFFSET_DEAD_COUNT]; + render_indirect_buffer[ri_base + RI_OFFSET_MAX_SPAWN] = dead_count; // Clear the rendering instance count, which will be upgraded by the update pass // with the particles actually alive at the end of their update (after aged). - render_indirect_buffer[ri_base + 1u] = 0u; + render_indirect_buffer[ri_base + RI_OFFSET_INSTANCE_COUNT] = 0u; // Swap ping/pong buffers - let ping = render_indirect_buffer[ri_base + 8u]; + let ping = render_indirect_buffer[ri_base + RI_OFFSET_PING]; let pong = 1u - ping; - render_indirect_buffer[ri_base + 8u] = pong; + render_indirect_buffer[ri_base + RI_OFFSET_PING] = pong; // Copy the new pong into the dispatch buffer, which will be used during rendering // to determine where to read particle indices. - dispatch_indirect[di_base + 3u] = pong; + dispatch_indirect[di_base + DI_OFFSET_PONG] = pong; } diff --git a/src/render/vfx_init.wgsl b/src/render/vfx_init.wgsl index 41ada2df..aaadcd57 100644 --- a/src/render/vfx_init.wgsl +++ b/src/render/vfx_init.wgsl @@ -1,3 +1,5 @@ +#import bevy_hanabi::vfx_common::{ForceFieldSource, IndirectBuffer, RenderIndirect, SimParams, Spawner} + struct Particle { {{ATTRIBUTES}} } @@ -6,46 +8,6 @@ struct ParticleBuffer { particles: array, } -struct SimParams { - delta_time: f32, - time: f32, -} - -struct ForceFieldSource { - position: vec3, - max_radius: f32, - min_radius: f32, - mass: f32, - force_exponent: f32, - conform_to_sphere: f32, -}; - -struct Spawner { - transform: mat3x4, // transposed (row-major) - inverse_transform: mat3x4, // transposed (row-major) - spawn: i32, - seed: u32, - count: atomic, - effect_index: u32, - force_field: array, -} - -struct IndirectBuffer { - indices: array, -} - -struct RenderIndirectBuffer { - vertex_count: u32, - instance_count: atomic, - base_index: u32, - vertex_offset: i32, - base_instance: u32, - alive_count: atomic, - dead_count: atomic, - max_spawn: u32, - ping: u32, -} - {{PROPERTIES}} @group(0) @binding(0) var sim_params : SimParams; @@ -53,7 +15,7 @@ struct RenderIndirectBuffer { @group(1) @binding(1) var indirect_buffer : IndirectBuffer; {{PROPERTIES_BINDING}} @group(2) @binding(0) var spawner : Spawner; // NOTE - same group as update -@group(3) @binding(0) var render_indirect : RenderIndirectBuffer; +@group(3) @binding(0) var render_indirect : RenderIndirect; var seed : u32 = 0u; diff --git a/src/render/vfx_render.wgsl b/src/render/vfx_render.wgsl index 75c6f30f..795adf5c 100644 --- a/src/render/vfx_render.wgsl +++ b/src/render/vfx_render.wgsl @@ -1,48 +1,14 @@ #import bevy_render::view::View +#import bevy_hanabi::vfx_common::{DispatchIndirect, ForceFieldSource, IndirectBuffer, SimParams, Spawner} struct Particle { {{ATTRIBUTES}} } -struct ParticlesBuffer { +struct ParticleBuffer { particles: array, } -struct SimParams { - delta_time: f32, - time: f32, -}; - -struct IndirectBuffer { - indices: array, -} - -struct DispatchIndirect { - x: u32, - y: u32, - z: u32, - pong: u32, -} - -struct ForceFieldSource { - position: vec3, - max_radius: f32, - min_radius: f32, - mass: f32, - force_exponent: f32, - conform_to_sphere: f32, -} - -struct Spawner { - transform: mat3x4, // transposed (row-major) - inverse_transform: mat3x4, // transposed (row-major) - spawn: i32, - seed: u32, - count: i32, - effect_index: u32, - force_field: array, -} - struct VertexOutput { @builtin(position) position: vec4, @location(0) color: vec4, @@ -53,7 +19,7 @@ struct VertexOutput { @group(0) @binding(0) var view: View; @group(0) @binding(1) var sim_params : SimParams; -@group(1) @binding(0) var particle_buffer : ParticlesBuffer; +@group(1) @binding(0) var particle_buffer : ParticleBuffer; @group(1) @binding(1) var indirect_buffer : IndirectBuffer; @group(1) @binding(2) var dispatch_indirect : DispatchIndirect; #ifdef LOCAL_SPACE_SIMULATION diff --git a/src/render/vfx_update.wgsl b/src/render/vfx_update.wgsl index 3bfd9467..33fd9d3a 100644 --- a/src/render/vfx_update.wgsl +++ b/src/render/vfx_update.wgsl @@ -1,51 +1,12 @@ +#import bevy_hanabi::vfx_common::{ForceFieldSource, IndirectBuffer, RenderIndirect, SimParams, Spawner} + struct Particle { {{ATTRIBUTES}} -}; +} struct ParticleBuffer { particles: array, -}; - -struct SimParams { - delta_time: f32, - time: f32, -}; - -struct ForceFieldSource { - position: vec3, - max_radius: f32, - min_radius: f32, - mass: f32, - force_exponent: f32, - conform_to_sphere: f32, -}; - -struct Spawner { - transform: mat3x4, // transposed (row-major) - inverse_transform: mat3x4, // transposed (row-major) - spawn: atomic, - seed: u32, - count_unused: u32, - effect_index: u32, - force_field: array, -}; - -struct IndirectBuffer { - indices: array, -}; - -struct RenderIndirectBuffer { - vertex_count: u32, - instance_count: atomic, - base_index: u32, - vertex_offset: i32, - base_instance: u32, - alive_count: atomic, - dead_count: atomic, - max_spawn: atomic, - ping: u32, - max_update: u32, -}; +} {{PROPERTIES}} @@ -54,7 +15,7 @@ struct RenderIndirectBuffer { @group(1) @binding(1) var indirect_buffer : IndirectBuffer; {{PROPERTIES_BINDING}} @group(2) @binding(0) var spawner : Spawner; // NOTE - same group as init -@group(3) @binding(0) var render_indirect : RenderIndirectBuffer; +@group(3) @binding(0) var render_indirect : RenderIndirect; var seed : u32 = 0u;