Search across packages, versions, files, declarations, and comments. Use multiple terms to narrow results (AND), and add @vertex, @fragment, or @compute to filter by shader stage.
1#define_import_path bevy_hanabi::vfx_common
3struct SimParams {
4 /// Delta time in seconds since last simulation tick.
5 delta_time: f32,
6 /// Time in seconds since the start of simulation.
7 time: f32,
8 /// Virtual delta time in seconds since last simulation tick.
9 virtual_delta_time: f32,
10 /// Virtual time in seconds since the start of simulation.
11 virtual_time: f32,
12 /// Real delta time in seconds since last simulation tick.
13 real_delta_time: f32,
14 /// Real time in seconds since the start of simulation.
15 real_time: f32,
16 /// Total number of effects to update this frame. Used by the indirect
17 /// compute pipeline to cap the compute thread to the actual number of
18 /// effects to process.
19 num_effects: u32,
22struct Spawner {
23 /// Compressed transform of the emitter.
24 transform: mat3x4<f32>, // transposed (row-major)
25 /// Inverse compressed transform of the emitter.
26 inverse_transform: mat3x4<f32>, // transposed (row-major)
27 /// Number of particles to spawn this frame, as calculated by the CPU Spawner.
28 ///
29 /// This is only used if the effect is not a child effect (driven by GPU events).
30 spawn: i32,
31 /// PRNG seed for this effect instance. Currently this can change each time the
32 /// effect is recompiled, and cannot be set deterministically (TODO).
33 seed: u32,
34 /// Index of the ping-pong buffer of particle indices to read particles from
35 /// during rendering. Cached from RenderIndirect::ping after it's swapped
36 /// in the indirect dispatch, because the RenderIndirect struct is used by GPU
37 /// as an indirect draw source so cannot also be bound as regular storage
38 /// buffer for reading.
39 render_pong: u32,
40 /// Index of the [`GpuEffectMetadata`] for this effect.
41 effect_metadata_index: u32,
42#ifdef SPAWNER_PADDING
43 {{SPAWNER_PADDING}}
44#endif
47const SPAWNER_OFFSET_PONG: u32 = 27u;
49struct IndirectBuffer {
50 indices: array<u32>,
53/// A event emitted by another effect during its update pass, to trigger the spawning
54/// of one or more particle in this effect.
55struct SpawnEvent {
56 /// The particle index in the parent effect buffer of the source particle which
57 /// triggered the event. This is used to inherit attributes like position or velocity.
58 particle_index: u32,
61/// Append buffer populated during the Update pass of the previous frame by a parent effect,
62/// and read back by its child effect(s) during the Init pass of the next frame.
63struct EventBuffer {
64 /// The spawn events themselves.
65 spawn_events: array<SpawnEvent>,
68/// Info about a single child of a parent effect.
69struct ChildInfo {
70 /// Index of the effect's IndirectDispatch entry in the global init indirect dispatch array.
71 init_indirect_dispatch_index: u32,
72 /// Number of events in the associated event buffer.
73#ifdef CHILD_INFO_IS_ATOMIC
74 event_count: atomic<i32>,
75#else
76 event_count: i32,
77#endif
80/// Buffer storing all the ChildInfo structs for several effects.
81struct ChildInfoBuffer {
82 /// The child info structs themselves.
83 rows: array<ChildInfo>,
86/// Indirect compute dispatch struct for GPU-driven passes. The layout of this struct is dictated by WGSL.
87struct IndirectDispatch {
88 /// Number of workgroups. Each workgroup has exactly 64 threads.
89 x: u32,
90 /// Unused; always 1.
91 y: u32,
92 /// Unused; always 1.
93 z: u32,
96/// Stride in u32 count (4 bytes) of the IndirectDispatch struct.
97const DISPATCH_INDIRECT_STRIDE: u32 = 3u;
99// Effect metadata offsets. Used when accessing a tightly packed array of EffectMetadata
100// as a raw array<u32>, so that we can avoid WGSL struct padding and keep data more compact
101// in the GPU buffer. Each offset corresponds to a field in the EffectMetadata struct.
102const EM_OFFSET_VERTEX_COUNT: u32 = 0u;
103const EM_OFFSET_INSTANCE_COUNT: u32 = 1u;
104const EM_OFFSET_FIRST_INDEX_OR_VERTEX_OFFSET: u32 = 2u;
105const EM_OFFSET_VERTEX_OFFSET_OR_BASE_INSTANCE: u32 = 3u;
106const EM_OFFSET_BASE_INSTANCE: u32 = 4u;
107const EM_OFFSET_ALIVE_COUNT: u32 = 5u;
108const EM_OFFSET_MAX_UPDATE: u32 = 6u;
109const EM_OFFSET_DEAD_COUNT: u32 = 7u;
110const EM_OFFSET_MAX_SPAWN: u32 = 8u;
111const EM_OFFSET_PING: u32 = 9u;
112//const EM_OFFSET_SPAWNER_INDEX: u32 = 10u;
113const EM_OFFSET_INDIRECT_DISPATCH_INDEX: u32 = 11u;
115/// Draw indirect parameters for GPU-driven rendering, and additional effect data.
116struct EffectMetadata {
117 /// Number of vertices in the particle mesh.
118 vertex_count: u32,
119 /// Number of mesh instances, equal to the number of particles.
120 instance_count: atomic<u32>,
121 /// First index (if indexed) or vertex offset (if non-indexed).
122 first_index_or_vertex_offset: u32,
123 /// Vertex offset (if indexed) or base instance (if non-indexed).
124 vertex_offset_or_base_instance: i32,
125 /// Base instance (if indexed).
126 base_instance: u32,
128 /// Number of particles alive after the init pass, used to calculate the number
129 /// of compute threads to spawn for the update pass and to cap those threads
130 /// via `max_update`.
131 alive_count: atomic<u32>,
132 /// Maximum number of update threads to run. This is cached from `alive_count`
133 /// during the indirect dispatch, so that the update compute pass can cap its
134 /// thread count while also modifying the actual `alive_count` if some particle
135 /// dies during the update pass.
136 max_update: u32,
137 /// Number of dead particles, decremented during the init pass as new particles
138 /// are spawned, and incremented during the update pass as existing particles die.
139 dead_count: atomic<u32>,
140 /// Maxmimum number of init threads to run on next frame. This is cached from
141 /// `dead_count` during the indirect dispatch of the previous frame, so that the
142 /// init compute pass can cap its thread count while also decrementing the actual
143 /// `dead_count` as particles are spawned.
144 max_spawn: atomic<u32>,
145 /// Index of the ping buffer for particle indices. Init and update compute passes
146 /// always write into the ping buffer and read from the pong buffer. The buffers
147 /// are swapped (ping = 1 - ping) during the indirect dispatch.
148 ping: u32,
149 /// Unused. TODO remove.
150 spawner_index: u32,
151 /// Index of the [`GpuDispatchIndirect`] struct inside the global
152 /// [`EffectsMeta::dispatch_indirect_buffer`].
153 indirect_dispatch_index: u32,
154 /// Index of the [`GpuRenderIndirect`] struct inside the global
155 /// [`EffectsMeta::render_group_dispatch_buffer`].
156 indirect_render_index: u32,
157 /// Offset (in u32 count) of the init indirect dispatch struct inside its
158 /// buffer. This avoids having to align those 16-byte structs to the GPU
159 /// alignment (at least 32 bytes, even 256 bytes on some).
160 init_indirect_dispatch_index: u32,
161 /// Index of this effect into its parent's ChildInfo array
162 /// ([`EffectChildren::effect_cache_ids`] and its associated GPU
163 /// array). This starts at zero for the first child of each effect, and is
164 /// only unique per parent, not globally. Only available if this effect is a
165 /// child of another effect (i.e. if it has a parent).
166 local_child_index: u32,
167 /// For children, global index of the ChildInfo into the shared array.
168 global_child_index: u32,
169 /// For parents, base index of the their first ChildInfo into the shared array.
170 base_child_index: u32,
172 /// Particle stride, in number of u32.
173 particle_stride: u32,
174 /// Offset from the particle start to the first sort key, in number of u32.
175 sort_key_offset: u32,
176 /// Offset from the particle start to the second sort key, in number of u32.
177 sort_key2_offset: u32,
179 /// Atomic counter incremented each time a particle spawns. Useful for
180 /// things like RIBBON_ID or any other use where a unique value is needed.
181 /// The value loops back after some time, but unless some particle lives
182 /// forever there's little chance of repetition.
183 particle_counter: atomic<u32>,
185 /// Padding for storage buffer alignment. This struct is sometimes bound as part
186 /// of an array, or sometimes individually as a single unit. In the later case,
187 /// we need it to be aligned to the GPU limits of the device. That limit is only
188 /// known at runtime when initializing the WebGPU device.
189 // FIXME - not anymore, but would be again with proper batching, so keep for now
190 {{EFFECT_METADATA_PADDING}}
193/// Stride, in u32 count, between elements of an array<EffectMetadata>.
194const EFFECT_METADATA_STRIDE: u32 = {{EFFECT_METADATA_STRIDE}} / 4u;
196var<private> seed : u32 = 0u;
198const tau: f32 = 6.283185307179586476925286766559;
200// Rand: PCG
201// https://www.reedbeta.com/blog/hash-functions-for-gpu-rendering/
202fn pcg_hash(input: u32) -> u32 {
203 var state: u32 = input * 747796405u + 2891336453u;
204 var word: u32 = ((state >> ((state >> 28u) + 4u)) ^ state) * 277803737u;
205 return (word >> 22u) ^ word;
208fn to_float01(u: u32) -> f32 {
209 // Note: could generate only 24 bits of randomness
210 return bitcast<f32>((u & 0x007fffffu) | 0x3f800000u) - 1.;
213// Random floating-point number in [0:1]
214fn frand() -> f32 {
215 seed = pcg_hash(seed);
216 return to_float01(pcg_hash(seed));
219// Random floating-point number in [0:1]^2
220fn frand2() -> vec2<f32> {
221 seed = pcg_hash(seed);
222 var x = to_float01(seed);
223 seed = pcg_hash(seed);
224 var y = to_float01(seed);
225 return vec2<f32>(x, y);
228// Random floating-point number in [0:1]^3
229fn frand3() -> vec3<f32> {
230 seed = pcg_hash(seed);
231 var x = to_float01(seed);
232 seed = pcg_hash(seed);
233 var y = to_float01(seed);
234 seed = pcg_hash(seed);
235 var z = to_float01(seed);
236 return vec3<f32>(x, y, z);
239// Random floating-point number in [0:1]^4
240fn frand4() -> vec4<f32> {
241 // Each rand() produces 32 bits, and we need 24 bits per component,
242 // so can get away with only 3 calls.
243 var r0 = pcg_hash(seed);
244 var r1 = pcg_hash(r0);
245 var r2 = pcg_hash(r1);
246 seed = r2;
247 var x = to_float01(r0);
248 var r01 = (r0 & 0xff000000u) >> 8u | (r1 & 0x0000ffffu);
249 var y = to_float01(r01);
250 var r12 = (r1 & 0xffff0000u) >> 8u | (r2 & 0x000000ffu);
251 var z = to_float01(r12);
252 var r22 = r2 >> 8u;
253 var w = to_float01(r22);
254 return vec4<f32>(x, y, z, w);
257fn rand_uniform_f(a: f32, b: f32) -> f32 {
258 return a + frand() * (b - a);
261fn rand_uniform_vec2(a: vec2<f32>, b: vec2<f32>) -> vec2<f32> {
262 return a + frand2() * (b - a);
265fn rand_uniform_vec3(a: vec3<f32>, b: vec3<f32>) -> vec3<f32> {
266 return a + frand3() * (b - a);
269fn rand_uniform_vec4(a: vec4<f32>, b: vec4<f32>) -> vec4<f32> {
270 return a + frand4() * (b - a);
273// Normal distribution computed using Box-Muller transform
274fn rand_normal_f(mean: f32, std_dev: f32) -> f32 {
275 var u = frand();
276 var v = frand();
277 var r = sqrt(-2.0 * log(u));
278 return mean + std_dev * r * cos(tau * v);
281fn rand_normal_vec2(mean: vec2f, std_dev: vec2f) -> vec2f {
282 var u = frand();
283 var v = frand2();
284 var r = sqrt(-2.0 * log(u));
285 return mean + std_dev * r * cos(tau * v);
288fn rand_normal_vec3(mean: vec3f, std_dev: vec3f) -> vec3f {
289 var u = frand();
290 var v = frand3();
291 var r = sqrt(-2.0 * log(u));
292 return mean + std_dev * r * cos(tau * v);
295fn rand_normal_vec4(mean: vec4f, std_dev: vec4f) -> vec4f {
296 var u = frand();
297 var v = frand4();
298 var r = sqrt(-2.0 * log(u));
299 return mean + std_dev * r * cos(tau * v);
302fn proj(u: vec3<f32>, v: vec3<f32>) -> vec3<f32> {
303 return dot(v, u) / dot(u,u) * u;
Contents