1#import bevy_hanabi::vfx_common::{
2    ChildInfo, ChildInfoBuffer, EventBuffer, IndirectDispatch, IndirectBuffer,
3    EffectMetadata, RenderGroupIndirect, SimParams, Spawner,
4    seed, tau, pcg_hash, to_float01, frand, frand2, frand3, frand4,
5    rand_uniform_f, rand_uniform_vec2, rand_uniform_vec3, rand_uniform_vec4,
6    rand_normal_f, rand_normal_vec2, rand_normal_vec3, rand_normal_vec4, proj
7}
8
9struct Particle {
10{{ATTRIBUTES}}
11}
12
13struct ParticleBuffer {
14    particles: array<Particle>,
15}
16
17#ifdef READ_PARENT_PARTICLE
18
19struct ParentParticle {
20    {{PARENT_ATTRIBUTES}}
21}
22
23struct ParentParticleBuffer {
24    particles: array<ParentParticle>,
25}
26
27#endif
28
29{{PROPERTIES}}
30
31@group(0) @binding(0) var<uniform> sim_params : SimParams;
32
33// "particle" group @1
34@group(1) @binding(0) var<storage, read_write> particle_buffer : ParticleBuffer;
35@group(1) @binding(1) var<storage, read_write> indirect_buffer : IndirectBuffer;
36#ifdef READ_PARENT_PARTICLE
37@group(1) @binding(2) var<storage, read> parent_particle_buffer : ParentParticleBuffer;
38#endif
39
40// "spawner" group @2
41@group(2) @binding(0) var<storage, read> spawner : Spawner;
42{{PROPERTIES_BINDING}}
43
44// "metadata" group @3
45@group(3) @binding(0) var<storage, read_write> effect_metadata : EffectMetadata;
46#ifdef CONSUME_GPU_SPAWN_EVENTS
47@group(3) @binding(1) var<storage, read> child_info_buffer : ChildInfoBuffer;
48@group(3) @binding(2) var<storage, read> event_buffer : EventBuffer;
49#endif
50
51{{INIT_EXTRA}}
52
53@compute @workgroup_size(64)
54fn main(@builtin(global_invocation_id) global_invocation_id: vec3<u32>) {
55    let thread_index = global_invocation_id.x;
56
57    // Cap to max number of dead particles, copied from dead_count at the end of the
58    // previous iteration, and constant during this pass (unlike dead_count).
59    let max_spawn = atomicLoad(&effect_metadata.max_spawn);
60    if (thread_index >= max_spawn) {
61        return;
62    }
63
64    // Cap to the actual number of spawning requested by CPU or GPU, since compute shaders run
65    // in workgroup_size(64) so more threads than needed are launched (rounded up to 64).
66#ifdef CONSUME_GPU_SPAWN_EVENTS
67    let event_index = thread_index;
68    let global_child_index = effect_metadata.global_child_index;
69    let event_count = child_info_buffer.rows[global_child_index].event_count;
70    if (event_index >= u32(event_count)) {
71        return;
72    }
73#else
74    // Cap to the actual number of spawning requested by CPU (in the case of
75    // spawners) or the number of particles present in the source group (in the
76    // case of cloners).
77    let spawn_count: u32 = u32(spawner.spawn);
78    if (thread_index >= spawn_count) {
79        return;
80    }
81#endif
82
83    // Always write into ping, read from pong
84    let write_index = effect_metadata.ping;
85    let read_index = 1u - write_index;
86
87    // Recycle a dead particle from the destination group
88    let dead_index = atomicSub(&effect_metadata.dead_count, 1u) - 1u;
89    let particle_index = indirect_buffer.indices[3u * dead_index + 2u];
90
91    let particle_counter = atomicAdd(&effect_metadata.particle_counter, 1u);
92
93    // Initialize the PRNG seed
94    seed = pcg_hash(particle_index ^ spawner.seed);
95
96    // Spawner transform
97    let transform = transpose(
98        mat4x4(
99            spawner.transform[0],
100            spawner.transform[1],
101            spawner.transform[2],
102            vec4<f32>(0.0, 0.0, 0.0, 1.0)
103        )
104    );
105
106#ifdef READ_PARENT_PARTICLE
107    // Fetch parent particle which triggered this spawn
108    let parent_particle_index = event_buffer.spawn_events[event_index].particle_index;
109    let parent_particle = parent_particle_buffer.particles[parent_particle_index];
110#endif
111
112    // Initialize new particle
113    var particle = Particle();
114    {{INIT_CODE}}
115#ifdef ATTRIBUTE_PREV
116    particle.prev = 0xffffffffu;
117#endif  // ATTRIBUTE_PREV
118#ifdef ATTRIBUTE_NEXT
119    particle.next = 0xffffffffu;
120#endif  // ATTRIBUTE_NEXT
121
122    // Only add emitter's transform to CPU-spawned particles. GPU-spawned particles
123    // don't have a CPU spawner to read from (they can read the parent particle
124    // transform, which is equivalent).
125#ifndef CONSUME_GPU_SPAWN_EVENTS
126    {{SIMULATION_SPACE_TRANSFORM_PARTICLE}}
127#endif
128
129    // Count as alive
130    atomicAdd(&effect_metadata.alive_count, 1u);
131
132    // Add to alive list
133    let instance_index = atomicAdd(&effect_metadata.instance_count, 1u);
134    indirect_buffer.indices[3u * instance_index + write_index] = particle_index;
135
136    // Write back new particle
137    particle_buffer.particles[particle_index] = particle;
138}