From 0e3b354520ec7b62979f7a0362c7f1485beff9b9 Mon Sep 17 00:00:00 2001 From: Brandon Thomas Date: Sat, 3 Oct 2026 02:44:36 -0400 Subject: [PATCH] M12.7: GPU particles (CC Particle World, CC Particle Systems II, Particle Playground cannon): compute-shader simulation with GPU checkpoints, CPU as the oracle Co-Authored-By: Claude Opus 5.5 --- crates/effects/src/lib.rs | 8 + crates/effects/src/psim.rs | 106 ++++++++++ crates/effects/src/sim.rs | 88 ++++++++- crates/effects/src/sim3.rs | 76 +++++++- crates/gpu/src/context.rs | 5 + crates/gpu/src/lib.rs | 13 ++ crates/gpu/src/particles.rs | 209 ++++++++++++++++++++ crates/gpu/src/shaders/particles.wgsl | 270 ++++++++++++++++++++++++++ crates/gpu/src/tests_particles.rs | 155 +++++++++++++++ crates/render/src/lib.rs | 8 + 10 files changed, 925 insertions(+), 13 deletions(-) create mode 100644 crates/effects/src/psim.rs create mode 100644 crates/gpu/src/particles.rs create mode 100644 crates/gpu/src/shaders/particles.wgsl create mode 100644 crates/gpu/src/tests_particles.rs diff --git a/crates/effects/src/lib.rs b/crates/effects/src/lib.rs index a5b4f3465..8840a3839 100644 --- a/crates/effects/src/lib.rs +++ b/crates/effects/src/lib.rs @@ -40,6 +40,7 @@ mod ocio; pub mod paint; mod perspective; mod perspective2; +pub mod psim; pub mod puppet; pub mod roto; mod sim; @@ -67,6 +68,8 @@ use effectcraft_project::build::Ids; use effectcraft_project::{GroupKind, ParamUi, PropGroup, Property}; pub use effectcraft_raster::{AuxChannels, Image}; pub use misc::{INVERT_ALPHA, INVERT_CHANNELS}; +pub use sim::particle_state; +pub use sim3::playground_state; /// Effect categories in Effects & Presets order. pub const CATEGORIES: &[&str] = &[ @@ -260,6 +263,11 @@ pub trait EffectHost: Sync { fn aux(&self) -> Option> { None } + /// A particle simulation backend (GPU particles; see [`psim`]). `None` = simulate on the + /// CPU. + fn particles(&self) -> Option<&dyn psim::ParticleSim> { + None + } } /// Extra context the renderer may supply (all optional; `Default` is "nothing known"). diff --git a/crates/effects/src/psim.rs b/crates/effects/src/psim.rs new file mode 100644 index 000000000..aabbcdce8 --- /dev/null +++ b/crates/effects/src/psim.rs @@ -0,0 +1,106 @@ +//! Pluggable particle simulation (GPU particles, M12.7). +//! +//! The stepped particle systems (CC Particle World, CC Particle Systems II and Particle +//! Playground's Cannon without interacting forces) evolve every particle independently: births +//! come from a rate (the same `carry` arithmetic as the CPU step), and each particle integrates +//! its own motion with the same fixed step. An accelerator ([`ParticleSim`], reached through +//! [`crate::EffectHost::particles`]) can therefore simulate all of them in parallel and return +//! the state the CPU's `SimCache` would reach. The CPU stays the oracle: anything the backend +//! declines (interacting forces, caps that would drop births) simulates on the CPU. + +/// CC Particle World / CC Particle Systems II physics (see `sim::Phys`), with the per-step and +/// per-spawn constants already evaluated on the CPU (bit-identical to the CPU's). +#[derive(Clone, Copy, Debug, PartialEq)] +pub struct EngineDesc { + pub life: f32, + pub producer: [f32; 3], + pub radius: [f32; 3], + /// Animation (0 Explosive … 9 Fractal Uni). + pub anim: u32, + pub speed: f32, + pub gravity: [f32; 3], + /// Normalised direction axis. + pub axis: [f32; 3], + pub extra: f32, + /// Cone Axis spread factor `|tan(extra angle / 2)|` (≤ 20). + pub cone: f32, + /// Velocity damping per step `exp(−drag · dt)`. + pub damp: f32, + pub seed: u32, + /// 2D systems keep z = 0. + pub flat: bool, +} + +/// Particle Playground's Cannon with gravity only (see `sim3::Pg`). +#[derive(Clone, Copy, Debug, PartialEq)] +pub struct CannonDesc { + pub pos: [f32; 2], + pub barrel_angle: f32, + pub barrel_radius: f32, + pub dir_spread: f32, + pub vel: f32, + pub vel_spread: f32, + pub gravity: [f32; 2], + pub grav_spread: f32, + pub seed: u32, + /// Particles leaving `[x0, y0, x1, y1]` are removed. + pub bounds: [f32; 4], +} + +#[derive(Clone, Copy, Debug, PartialEq)] +pub enum ParticleSystem { + Engine(EngineDesc), + Cannon(CannonDesc), +} + +/// One simulation request: the state after `steps` steps of system `key`. +#[derive(Clone, Debug)] +pub struct SimRequest<'a> { + /// Identifies the system (parameters, seed); checkpoints are reused per key. + pub key: u64, + pub steps: u64, + /// Birth step of every particle id born before `steps` (ascending). + pub births: &'a [u32], + pub system: ParticleSystem, +} + +/// A live particle (ascending ids, as the CPU state keeps them). +#[derive(Clone, Copy, Debug, PartialEq)] +pub struct SimParticle { + pub p: [f32; 3], + pub v: [f32; 3], + pub age: f32, + pub life: f32, + pub rnd: f32, + pub id: u32, +} + +/// A particle simulation backend (the GPU). +pub trait ParticleSim: Send + Sync { + /// The live particles after `req.steps` steps, or `None` (simulate on the CPU). + fn simulate(&self, req: &SimRequest) -> Option>; +} + +/// Birth step of every particle born in the first `steps` steps at `rate` particles per second +/// (the CPU step's `carry` arithmetic, exactly). `None` when more than `cap` would be born (the +/// CPU drops births at its live-particle cap, which an independent simulation can't know). +pub fn births(rate: f64, steps: u64, cap: usize) -> Option> { + let per = rate / SPS; + if per * steps as f64 > cap as f64 + 2.0 { + return None; + } + let mut out = Vec::new(); + let mut carry = 0.0f64; + for i in 0..steps { + carry += per; + let n = carry.floor() as u32; + carry -= n as f64; + for _ in 0..n { + out.push(i as u32); + } + } + (out.len() <= cap).then_some(out) +} + +/// Steps per second of the stepped simulations. +pub const SPS: f64 = crate::sim::SPS; diff --git a/crates/effects/src/sim.rs b/crates/effects/src/sim.rs index 7f90af6b1..18beeb762 100644 --- a/crates/effects/src/sim.rs +++ b/crates/effects/src/sim.rs @@ -818,6 +818,60 @@ impl Phys { } } +impl Phys { + /// The backend description (per-step and per-spawn constants evaluated here, in f32). + pub(crate) fn desc(&self) -> crate::psim::EngineDesc { + let mut drag = self.resistance; + if self.anim == 3 { + drag += 3.0 * self.extra.max(0.1); + } + let dt = (1.0 / SPS) as f32; + crate::psim::EngineDesc { + life: self.life, + producer: self.producer, + radius: self.radius, + anim: self.anim, + speed: self.speed, + gravity: self.gravity, + axis: norm3(self.axis), + extra: self.extra, + cone: (self.extra_angle.to_radians() * 0.5).tan().abs().min(20.0), + damp: (-drag * dt).exp(), + seed: self.seed, + flat: self.flat, + } + } +} + +/// The state at layer time `ctx.time` from the host's particle backend (GPU particles), or +/// `None` (simulate on the CPU). +fn accelerated(ctx: &EffectCtx, key: u64, phys: &Phys) -> Option> { + let backend = ctx.env.host?.particles()?; + let steps = steps_at(ctx.time); + let births = crate::psim::births(phys.rate, steps, phys.max)?; + let req = crate::psim::SimRequest { key, steps, births: &births, system: crate::psim::ParticleSystem::Engine(phys.desc()) }; + let parts = backend.simulate(&req)?; + Some(std::sync::Arc::new(PState { + parts: parts.into_iter().map(|q| Particle { p: q.p, v: q.v, age: q.age, life: q.life, rnd: q.rnd, id: q.id }).collect(), + carry: 0.0, + next_id: births.len() as u32, + })) +} + +/// The live particles of a stepped particle effect (`ec.sim.ccparticleworld`, +/// `ec.sim.ccparticlesystems2`) at `ctx.time`: from the host's particle backend when it has +/// one (GPU particles), else from the CPU simulation. For comparing the two. +pub fn particle_state(id: &str, ctx: &EffectCtx) -> Option> { + let (phys, cache, salt) = match id { + "ec.sim.ccparticleworld" => (pw_phys(ctx), &PW_CACHE, 1), + "ec.sim.ccparticlesystems2" => (ps2_phys(ctx), &PS2_CACHE, 2), + _ => return None, + }; + let key = params_key(ctx, &Buf { img: Image::new(0, 0), offset: [0.0; 2], scale: 1.0 }, salt); + let st = accelerated(ctx, key, &phys).unwrap_or_else(|| simulate(cache, key, ctx.time, &phys)); + Some(st.parts.iter().map(|q| crate::psim::SimParticle { p: q.p, v: q.v, age: q.age, life: q.life, rnd: q.rnd, id: q.id }).collect()) +} + /// Run a particle system to layer time `t` through `cache`. pub(crate) fn simulate(cache: &SimCache, key: u64, t: f64, phys: &Phys) -> std::sync::Arc { cache.run(key, steps_at(t), PState::default, |st, i| phys.step(st, i)) @@ -859,10 +913,10 @@ fn type_sprite(kind: u32, x: f32, y: f32, r: f32, c: [f32; 4], vx: f32, vy: f32, static PW_CACHE: SimCache = SimCache::new(6); -fn particle_world(ctx: &EffectCtx, mut b: Buf) -> Buf { +/// CC Particle World's physics. +fn pw_phys(ctx: &EffectCtx) -> Phys { let pr = ctx.params; - let (lw, lh) = (ctx.layer_size[0] as f32, ctx.layer_size[1] as f32); - let phys = Phys { + Phys { rate: pr.f("birthRate").max(0.0) * 120.0, life: pr.f("longevity").max(0.0) as f32, producer: [pr.f("producer/producerX") as f32, pr.f("producer/producerY") as f32, pr.f("producer/producerZ") as f32], @@ -882,11 +936,18 @@ fn particle_world(ctx: &EffectCtx, mut b: Buf) -> Buf { seed: (pr.f("extras/randomSeed") as u32).wrapping_mul(0x9e3779b9) ^ ctx.seed, max: 40_000, flat: false, - }; + } +} + +fn particle_world(ctx: &EffectCtx, mut b: Buf) -> Buf { + let pr = ctx.params; + let (lw, lh) = (ctx.layer_size[0] as f32, ctx.layer_size[1] as f32); + let phys = pw_phys(ctx); if phys.rate <= 0.0 && ctx.time <= 0.0 { return b; } - let st = simulate(&PW_CACHE, params_key(ctx, &Buf { img: Image::new(0, 0), offset: [0.0; 2], scale: 1.0 }, 1), ctx.time, &phys); + let key = params_key(ctx, &Buf { img: Image::new(0, 0), offset: [0.0; 2], scale: 1.0 }, 1); + let st = accelerated(ctx, key, &phys).unwrap_or_else(|| simulate(&PW_CACHE, key, ctx.time, &phys)); let kind = pr.e("particle/particleType"); let birth = pr.color("particle/birthColor"); let death = pr.color("particle/deathColor"); @@ -936,13 +997,14 @@ fn particle_world(ctx: &EffectCtx, mut b: Buf) -> Buf { static PS2_CACHE: SimCache = SimCache::new(6); -fn particle_systems2(ctx: &EffectCtx, mut b: Buf) -> Buf { +/// CC Particle Systems II's physics. +fn ps2_phys(ctx: &EffectCtx) -> Phys { let pr = ctx.params; let lh = ctx.layer_size[1] as f32; let unit = lh.max(1.0); let pos = pr.v2("producer/position"); let dir = (pr.f("physics/direction") as f32).to_radians(); - let phys = Phys { + Phys { rate: pr.f("birthRate").max(0.0) * 60.0, life: pr.f("longevity").max(0.0) as f32, producer: [pos[0] as f32 / unit, pos[1] as f32 / unit, 0.0], @@ -958,8 +1020,16 @@ fn particle_systems2(ctx: &EffectCtx, mut b: Buf) -> Buf { seed: (pr.f("randomSeed") as u32).wrapping_mul(0x85ebca6b) ^ ctx.seed, max: 40_000, flat: true, - }; - let st = simulate(&PS2_CACHE, params_key(ctx, &Buf { img: Image::new(0, 0), offset: [0.0; 2], scale: 1.0 }, 2), ctx.time, &phys); + } +} + +fn particle_systems2(ctx: &EffectCtx, mut b: Buf) -> Buf { + let pr = ctx.params; + let lh = ctx.layer_size[1] as f32; + let unit = lh.max(1.0); + let phys = ps2_phys(ctx); + let key = params_key(ctx, &Buf { img: Image::new(0, 0), offset: [0.0; 2], scale: 1.0 }, 2); + let st = accelerated(ctx, key, &phys).unwrap_or_else(|| simulate(&PS2_CACHE, key, ctx.time, &phys)); let kind = pr.e("particle/particleType"); let birth = pr.color("particle/birthColor"); let death = pr.color("particle/deathColor"); diff --git a/crates/effects/src/sim3.rs b/crates/effects/src/sim3.rs index 5560917a1..f7425df7c 100644 --- a/crates/effects/src/sim3.rs +++ b/crates/effects/src/sim3.rs @@ -231,6 +231,38 @@ impl Pg { static PG_CACHE: SimCache = SimCache::new(6); +/// Particle Playground cannons with gravity only, simulated by the host's particle backend. +fn pg_accelerated(ctx: &EffectCtx, key: u64, pg: &Pg) -> Option> { + let backend = ctx.env.host?.particles()?; + let steps = steps_at(ctx.time); + let births = crate::psim::births(pg.rate, steps, 60_000)?; + let desc = crate::psim::CannonDesc { + pos: pg.pos, + barrel_angle: pg.barrel_angle, + barrel_radius: pg.barrel_radius, + dir_spread: pg.dir_spread, + vel: pg.vel, + vel_spread: pg.vel_spread, + gravity: pg.gravity, + grav_spread: pg.grav_spread, + seed: pg.seed, + bounds: pg.bounds, + }; + let req = crate::psim::SimRequest { key, steps, births: &births, system: crate::psim::ParticleSystem::Cannon(desc) }; + let parts = backend.simulate(&req)?; + Some(std::sync::Arc::new(PgState { + parts: parts + .into_iter() + .map(|q| { + let born = pg.spawn(q.id); + PgParticle { p: [q.p[0], q.p[1]], v: [q.v[0], q.v[1]], ..born } + }) + .collect(), + carry: 0.0, + next_id: births.len() as u32, + })) +} + fn image_key(img: &Image) -> u64 { use std::hash::{Hash, Hasher}; let mut h = std::collections::hash_map::DefaultHasher::new(); @@ -244,14 +276,15 @@ fn image_key(img: &Image) -> u64 { h.finish() } -fn particle_playground(ctx: &EffectCtx, mut b: Buf) -> Buf { +/// Particle Playground's producers and forces. +fn pg_of(ctx: &EffectCtx) -> Pg { let pr = ctx.params; let (lw, lh) = (ctx.layer_size[0] as f32, ctx.layer_size[1] as f32); let grav_dir = (pr.f("gravity/gravityDirection") as f32).to_radians(); let gf = pr.f("gravity/gravityForce") as f32; let wall = (pr.f("wall/wallBoundary").round() as usize).checked_sub(1).and_then(|i| ctx.env.masks.get(i)).map(|m| m.points.clone()); let mapper = if pr.b("mapperEnabled") { ctx.layer_param("persistentPropertyMapper/useLayerAsMap", true).map(|lp| MapImg::from(ctx, lp)) } else { None }; - let pg = Pg { + Pg { cannon: pr.b("cannonEnabled"), pos: { pr.v2("cannon/cannonPosition").map(|v| v as f32) }, barrel_angle: pr.f("cannon/barrelAngle") as f32, @@ -275,7 +308,34 @@ fn particle_playground(ctx: &EffectCtx, mut b: Buf) -> Buf { ], seed: (pr.f("randomSeed") as u32).wrapping_mul(0x9e37_79b9) ^ ctx.seed, bounds: [-lw * 2.0, -lh * 2.0, lw * 3.0, lh * 3.0], - }; + } +} + +/// The live particles of a Particle Playground whose cannon is the only producer and force +/// besides gravity, at `ctx.time`: from the host's particle backend when it has one (GPU +/// particles), else from the CPU simulation (`None` for other set-ups). For comparing the two. +pub fn playground_state(ctx: &EffectCtx) -> Option> { + let pg = pg_of(ctx); + let pr = ctx.params; + let grid = pr.b("gridEnabled") && pr.f("grid/particlesAcross").round() > 0.0 && pr.f("grid/particlesDown").round() > 0.0; + let exploder = pr.b("exploderEnabled") && pr.get("layerExploder/explodeLayer").and_then(Value::as_layer).is_some(); + if !pg.cannon || grid || exploder || (pg.repel != 0.0 && pg.repel_radius > 0.0) || pg.wall.is_some() || pg.mapper.is_some() { + return None; + } + let key = params_key(ctx, &Buf { img: Image::new(0, 0), offset: [0.0; 2], scale: 1.0 }, 0x7067); + let st = pg_accelerated(ctx, key, &pg).unwrap_or_else(|| PG_CACHE.run(key, steps_at(ctx.time), PgState::default, |st, _| pg.step(st))); + Some( + st.parts + .iter() + .map(|q| crate::psim::SimParticle { p: [q.p[0], q.p[1], 0.0], v: [q.v[0], q.v[1], 0.0], age: 0.0, life: 0.0, rnd: 0.0, id: q.id }) + .collect(), + ) +} + +fn particle_playground(ctx: &EffectCtx, mut b: Buf) -> Buf { + let pr = ctx.params; + let (lw, lh) = (ctx.layer_size[0] as f32, ctx.layer_size[1] as f32); + let pg = pg_of(ctx); // Initial particles: Grid and Layer Exploder. let grid_on = pr.b("gridEnabled"); let across = pr.f("grid/particlesAcross").max(0.0).round() as u32; @@ -347,7 +407,15 @@ fn particle_playground(ctx: &EffectCtx, mut b: Buf) -> Buf { } st }; - let st = PG_CACHE.run(key, steps_at(ctx.time), init, |st, _| pg.step(st)); + // A cannon with gravity only: every particle moves on its own (GPU particles). + let independent = pg.cannon + && !(grid_on && across > 0 && down > 0) + && exploder.is_none() + && !(pg.repel != 0.0 && pg.repel_radius > 0.0) + && pg.wall.is_none() + && pg.mapper.is_none(); + let gpu = if independent { pg_accelerated(ctx, key, &pg) } else { None }; + let st = gpu.unwrap_or_else(|| PG_CACHE.run(key, steps_at(ctx.time), init, |st, _| pg.step(st))); let s = b.scale as f32; let sprites: Vec = st .parts diff --git a/crates/gpu/src/context.rs b/crates/gpu/src/context.rs index e7cda93e4..7dd14ba83 100644 --- a/crates/gpu/src/context.rs +++ b/crates/gpu/src/context.rs @@ -102,6 +102,9 @@ pub struct GpuContext { uploads: Mutex, /// Advanced 3D render pipelines (built on first use). pub(crate) adv3d: std::sync::OnceLock, + /// GPU particle pipeline (built on first use) and simulation checkpoints. + pub(crate) particles: crate::particles::PipesCell, + pub(crate) particle_states: crate::particles::StatesCell, } fn tex_entry(binding: u32) -> wgpu::BindGroupLayoutEntry { @@ -263,6 +266,8 @@ impl GpuContext { dummy_buf, uploads: Mutex::new(Uploads::default()), adv3d: std::sync::OnceLock::new(), + particles: Default::default(), + particle_states: Default::default(), }) } diff --git a/crates/gpu/src/lib.rs b/crates/gpu/src/lib.rs index 9ebce2090..83f27b26d 100644 --- a/crates/gpu/src/lib.rs +++ b/crates/gpu/src/lib.rs @@ -29,6 +29,7 @@ mod fx_color; mod fx_distort; mod fx_generate; mod ops; +mod particles; mod walk; use std::sync::Arc; @@ -139,6 +140,16 @@ impl Accelerator for Gpu { fn raster_3d(&self, scene: &effectcraft_render::three_d::adv::Scene) -> Option { adv3d::render(&self.ctx, scene) } + + fn particles(&self) -> Option<&dyn effectcraft_effects::psim::ParticleSim> { + self.ctx.can_readback().then_some(self as &dyn effectcraft_effects::psim::ParticleSim) + } +} + +impl effectcraft_effects::psim::ParticleSim for Gpu { + fn simulate(&self, req: &effectcraft_effects::psim::SimRequest) -> Option> { + particles::simulate(&self.ctx, req) + } } #[cfg(test)] @@ -149,3 +160,5 @@ mod tests_3d; mod tests_adjust; #[cfg(test)] mod tests_adv3d; +#[cfg(test)] +mod tests_particles; diff --git a/crates/gpu/src/particles.rs b/crates/gpu/src/particles.rs new file mode 100644 index 000000000..669a4ba05 --- /dev/null +++ b/crates/gpu/src/particles.rs @@ -0,0 +1,209 @@ +//! GPU particles (`particles.wgsl`): the stepped particle systems of `effectcraft-effects` +//! ([`effectcraft_effects::psim`]) simulated with one invocation per particle. States stay on +//! the GPU as checkpoints (per system key, like the CPU's `SimCache`), so playback advances from +//! the previous frame's state instead of simulating from layer time 0. + +use std::sync::{Mutex, OnceLock}; + +use effectcraft_effects::psim::{ParticleSystem, SimParticle, SimRequest}; +use wgpu::util::DeviceExt; + +use crate::context::GpuContext; + +/// Bytes of one particle state (`PState` in the shader). +const STATE_BYTES: u64 = 48; +/// Checkpoints kept (all keys). +const CHECKPOINTS: usize = 24; + +pub(crate) struct Pipes { + bgl: wgpu::BindGroupLayout, + pipe: wgpu::ComputePipeline, +} + +struct Checkpoint { + key: u64, + steps: u64, + n: u32, + buf: wgpu::Buffer, + used: u64, +} + +#[derive(Default)] +pub(crate) struct Checkpoints { + list: Vec, + clock: u64, +} + +fn pipes(g: &GpuContext) -> &Pipes { + g.particles.get_or_init(|| { + let d = &g.device; + let src = [include_str!("shaders/common.wgsl"), include_str!("shaders/kernels.wgsl"), include_str!("shaders/particles.wgsl")].concat(); + let module = + d.create_shader_module(wgpu::ShaderModuleDescriptor { label: Some("effectcraft particles"), source: wgpu::ShaderSource::Wgsl(src.into()) }); + let buffer = |binding: u32, ty: wgpu::BufferBindingType| wgpu::BindGroupLayoutEntry { + binding, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { ty, has_dynamic_offset: false, min_binding_size: None }, + count: None, + }; + let bgl = d.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("effectcraft particles"), + entries: &[ + buffer(0, wgpu::BufferBindingType::Uniform), + buffer(1, wgpu::BufferBindingType::Storage { read_only: true }), + buffer(2, wgpu::BufferBindingType::Storage { read_only: true }), + buffer(3, wgpu::BufferBindingType::Storage { read_only: false }), + ], + }); + let layout = d.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: Some("effectcraft particles"), + bind_group_layouts: &[None, Some(&bgl)], + immediate_size: 0, + }); + let pipe = d.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("particles"), + layout: Some(&layout), + module: &module, + entry_point: Some("particles"), + compilation_options: Default::default(), + cache: None, + }); + Pipes { bgl, pipe } + }) +} + +/// The shader's parameter block. +fn uniforms(req: &SimRequest, n: u32, s0: u64, prev_n: u32) -> Vec { + let dt = (1.0 / effectcraft_effects::psim::SPS) as f32; + let mut u = [n, s0 as u32, req.steps as u32, 0]; + let u2; + let mut f = [[0.0f32; 4]; 6]; + match &req.system { + ParticleSystem::Engine(e) => { + u2 = [e.anim, e.flat as u32, e.seed, prev_n]; + f[0] = [e.life, e.speed, e.extra, e.cone]; + f[1] = [e.producer[0], e.producer[1], e.producer[2], 0.0]; + f[2] = [e.radius[0], e.radius[1], e.radius[2], 0.0]; + f[3] = [e.gravity[0], e.gravity[1], e.gravity[2], 0.0]; + f[4] = [e.axis[0], e.axis[1], e.axis[2], 0.0]; + f[5] = [e.damp, dt, 0.0, 0.0]; + } + ParticleSystem::Cannon(c) => { + u[3] = 1; + u2 = [0, 1, c.seed, prev_n]; + f[0] = [c.pos[0], c.pos[1], c.barrel_angle, c.barrel_radius]; + f[1] = [c.dir_spread, c.vel, c.vel_spread, c.grav_spread]; + f[2] = [c.gravity[0], c.gravity[1], dt, 0.0]; + f[3] = c.bounds; + } + } + let mut out = Vec::with_capacity(128); + for x in u.iter().chain(&u2) { + out.extend_from_slice(&x.to_le_bytes()); + } + for x in f.iter().flatten() { + out.extend_from_slice(&x.to_le_bytes()); + } + out +} + +/// Simulate `req` (see [`effectcraft_effects::psim::ParticleSim`]). +pub(crate) fn simulate(g: &GpuContext, req: &SimRequest) -> Option> { + if !g.can_readback() || req.steps > u32::MAX as u64 { + return None; + } + let n = req.births.len() as u32; + if n == 0 { + return Some(vec![]); + } + let size = n as u64 * STATE_BYTES; + if size > g.device.limits().max_storage_buffer_binding_size { + return None; + } + let p = pipes(g); + // The latest checkpoint at or before the target step. + let (s0, prev, prev_n) = { + let mut c = g.particle_states.lock().ok()?; + c.clock += 1; + let now = c.clock; + match c.list.iter_mut().filter(|k| k.key == req.key && k.steps <= req.steps).max_by_key(|k| k.steps) { + Some(k) => { + k.used = now; + (k.steps, Some(k.buf.clone()), k.n) + } + None => (0, None, 0), + } + }; + let d = &g.device; + let births: Vec = req.births.iter().flat_map(|b| b.to_le_bytes()).collect(); + let births = d.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some("births"), contents: &births, usage: wgpu::BufferUsages::STORAGE }); + let prev = prev.unwrap_or_else(|| { + d.create_buffer(&wgpu::BufferDescriptor { label: Some("particles"), size: STATE_BYTES, usage: wgpu::BufferUsages::STORAGE, mapped_at_creation: false }) + }); + let next = d.create_buffer(&wgpu::BufferDescriptor { + label: Some("particles"), + size, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }); + let ub = d.create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: Some("particles"), + contents: &uniforms(req, n, s0, prev_n), + usage: wgpu::BufferUsages::UNIFORM, + }); + let bg = d.create_bind_group(&wgpu::BindGroupDescriptor { + label: None, + layout: &p.bgl, + entries: &[ + wgpu::BindGroupEntry { binding: 0, resource: ub.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 1, resource: births.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 2, resource: prev.as_entire_binding() }, + wgpu::BindGroupEntry { binding: 3, resource: next.as_entire_binding() }, + ], + }); + let read = d.create_buffer(&wgpu::BufferDescriptor { + label: Some("particles readback"), + size, + usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + let mut enc = d.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("particles") }); + { + let mut pass = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some("particles"), timestamp_writes: None }); + pass.set_pipeline(&p.pipe); + pass.set_bind_group(1, &bg, &[]); + pass.dispatch_workgroups(n.div_ceil(64), 1, 1); + } + enc.copy_buffer_to_buffer(&next, 0, &read, 0, size); + g.queue.submit([enc.finish()]); + let (tx, rx) = std::sync::mpsc::channel(); + read.map_async(wgpu::MapMode::Read, .., move |r| { + let _ = tx.send(r); + }); + g.device.poll(wgpu::PollType::wait_indefinitely()).ok()?; + rx.recv().ok()?.ok()?; + let view = read.get_mapped_range(..).ok()?; + let f: Vec = view.chunks_exact(4).map(|c| f32::from_le_bytes([c[0], c[1], c[2], c[3]])).collect(); + drop(view); + read.unmap(); + let parts = f + .chunks_exact(12) + .enumerate() + .filter(|(_, s)| s[9] != 0.0) + .map(|(id, s)| SimParticle { p: [s[0], s[1], s[2]], v: [s[4], s[5], s[6]], age: s[3], life: s[7], rnd: s[8], id: id as u32 }) + .collect(); + // Keep the new state as a checkpoint. + if let Ok(mut c) = g.particle_states.lock() { + let now = c.clock; + c.list.retain(|k| !(k.key == req.key && k.steps == req.steps)); + c.list.push(Checkpoint { key: req.key, steps: req.steps, n, buf: next, used: now }); + while c.list.len() > CHECKPOINTS { + let Some(i) = c.list.iter().enumerate().min_by_key(|(_, k)| k.used).map(|(i, _)| i) else { break }; + c.list.remove(i); + } + } + Some(parts) +} + +pub(crate) type PipesCell = OnceLock; +pub(crate) type StatesCell = Mutex; diff --git a/crates/gpu/src/shaders/particles.wgsl b/crates/gpu/src/shaders/particles.wgsl new file mode 100644 index 000000000..694d6fb67 --- /dev/null +++ b/crates/gpu/src/shaders/particles.wgsl @@ -0,0 +1,270 @@ +// GPU particles: effects::psim systems, one invocation per particle id. Mirrors the CPU step +// (effects::sim::Phys::spawn / step, effects::sim3::Pg::spawn / step) operation for operation. +// +// Concatenated after common.wgsl and kernels.wgsl (hash_noise, value_noise, powz); bindings in +// group 1 (group 0 is unused here). + +struct PSim { + // Particles, from step, to step, system (0 engine, 1 cannon). + u: vec4, + // Animation, flat, seed, particles in `prev`. + u2: vec4, + // Engine: f[0] = (life, speed, extra, cone), f[1] = producer, f[2] = radius, + // f[3] = gravity, f[4] = axis, f[5] = (damp, dt). + // Cannon: f[0] = (pos.x, pos.y, barrel angle, barrel radius), f[1] = (dir spread, vel, + // vel spread, grav spread), f[2] = (gravity.x, gravity.y, dt), f[3] = bounds. + f: array, 6>, +} + +struct PState { + // Position, age. + p: vec4, + // Velocity, life. + v: vec4, + // rnd, alive (1 / 0), id. + m: vec4, +} + +@group(1) @binding(0) var S: PSim; +@group(1) @binding(1) var births: array; +@group(1) @binding(2) var prev: array; +@group(1) @binding(3) var next: array; + +const PS_TAU: f32 = 6.2831855; + +fn ps_h(i: u32, k: u32) -> f32 { + return hash_noise(i, k, S.u2.z); +} + +fn ps_hs(i: u32, k: u32) -> f32 { + return hash_noise(i, k, S.u2.z) * 2.0 - 1.0; +} + +// sim::norm3. +fn ps_norm3(v: vec3) -> vec3 { + let l = sqrt(v.x * v.x + v.y * v.y + v.z * v.z); + if (l > 1e-9) { + return v / l; + } + return vec3(0.0, -1.0, 0.0); +} + +// sim::rand_dir. +fn ps_rand_dir(i: u32, k: u32, flat: bool) -> vec3 { + if (flat) { + let a = ps_h(i, k) * PS_TAU; + return vec3(cos(a), sin(a), 0.0); + } + let z = ps_hs(i, k); + let a = ps_h(i, k + 1u) * PS_TAU; + let r = sqrt(max(1.0 - z * z, 0.0)); + return vec3(r * cos(a), r * sin(a), z); +} + +struct Part { + p: vec3, + v: vec3, + age: f32, + life: f32, + rnd: f32, +} + +// sim::Phys::spawn (+ the step's sub-step birth offset). +fn ps_spawn_engine(id: u32) -> Part { + let flat = S.u2.y != 0u; + let anim = S.u2.x; + let producer = S.f[1].xyz; + let radius = S.f[2].xyz; + let u = ps_rand_dir(id, 1u, flat); + var e = 1.0 / 3.0; + if (flat) { + e = 0.5; + } + let rr = powz(ps_h(id, 3u), e); + var p = producer + u * radius * rr; + if (flat) { + p.z = 0.0; + } + let ax = S.f[4].xyz; + let jitter = ps_rand_dir(id, 5u, flat); + var dir = jitter; + switch anim { + case 1u, 9u: { + dir = ps_norm3(ax + jitter * 0.08); + } + case 2u: { + let k = S.f[0].w; + dir = ps_norm3(ax + jitter * k * ps_h(id, 8u)); + } + case 5u: { + var z = jitter.z; + if (flat) { + z = 0.0; + } + dir = vec3(jitter.x, 0.0, z); + } + case 6u: { + dir = ps_norm3(vec3(jitter.x * 0.25, -1.0, jitter.z * 0.25)); + } + case 7u: { + dir = ps_norm3(vec3(1.0, jitter.y * 0.05, jitter.z * 0.05)); + } + default: {} + } + let sp = S.f[0].y * (0.5 + 0.5 * ps_h(id, 9u)); + var v = dir * sp; + if (flat) { + v.z = 0.0; + } + let life = S.f[0].x * (0.8 + 0.4 * ps_h(id, 10u)); + let dt = S.f[5].y; + let frac = ps_h(id, 12u) * dt; + return Part(p + v * frac, v, frac, life, ps_h(id, 11u)); +} + +// sim::Phys::step for one particle at step `j`. +fn ps_step_engine(q: Part, j: u32) -> Part { + let flat = S.u2.y != 0u; + let anim = S.u2.x; + let seed = S.u2.z; + let dt = S.f[5].y; + let t = f32(j) * dt; + let producer = S.f[1].xyz; + let gravity = S.f[3].xyz; + let extra = S.f[0].z; + var a = gravity; + switch anim { + case 4u: { + let dx = q.p.x - producer.x; + let dz = q.p.z - producer.z; + let k = 4.0 * extra; + a.x += -dz * k; + a.z += dx * k; + if (flat) { + let dy = q.p.y - producer.y; + a.x += -dy * k; + a.y += dx * k; + } + } + case 5u: { + let dx = q.p.x - producer.x; + let dy = q.p.y - producer.y; + let k = 6.0 * extra; + a.x += -dy * k - dx * 2.0; + a.y += dx * k - dy * 2.0; + } + case 6u: { + let n = value_noise(q.p.x * 8.0, q.p.y * 8.0, t * 3.0, seed) - 0.5; + a.x += n * 4.0 * extra; + a.y -= gravity.y * 1.5 + 0.3 * extra; + } + case 8u, 9u: { + let f = 6.0 * extra; + a.x += (value_noise(q.p.x * 5.0, q.p.y * 5.0, q.p.z * 5.0 + t, seed) - 0.5) * f; + a.y += (value_noise(q.p.x * 5.0 + 17.0, q.p.y * 5.0, q.p.z * 5.0 + t, seed) - 0.5) * f; + if (!flat) { + a.z += (value_noise(q.p.x * 5.0, q.p.y * 5.0 + 31.0, q.p.z * 5.0 + t, seed) - 0.5) * f; + } + } + default: {} + } + let damp = S.f[5].x; + var o = q; + o.v = (q.v + a * dt) * damp; + o.p = q.p + o.v * dt; + if (flat) { + o.p.z = 0.0; + o.v.z = 0.0; + } + o.age = q.age + dt; + return o; +} + +// sim3::Pg::spawn (+ the sub-step birth offset). +fn ps_spawn_cannon(id: u32) -> Part { + let pos = S.f[0].xy; + let angle = S.f[0].z; + let r = S.f[0].w; + let spread = S.f[1].x; + let vel = S.f[1].y; + let vel_spread = S.f[1].z; + let a = radians(angle + ps_hs(id, 1u) * spread * 0.5); + let d = vec2(sin(a), -cos(a)); + let speed = max(vel + ps_hs(id, 3u) * vel_spread, 0.0); + var off: vec2; + if (r >= 0.0) { + off = vec2(ps_hs(id, 2u) * r, ps_hs(id, 4u) * r); + } else { + let aa = ps_h(id, 5u) * PS_TAU; + let rr = sqrt(ps_h(id, 6u)) * -r; + off = vec2(cos(aa) * rr, sin(aa) * rr); + } + let v = d * speed; + let dt = S.f[2].z; + let frac = ps_h(id, 9u) * dt; + let p = pos + off + v * frac; + return Part(vec3(p, 0.0), vec3(v, 0.0), 0.0, 0.0, 0.0); +} + +// sim3::Pg::step for one particle (gravity only); returns false when it leaves the bounds. +fn ps_step_cannon(q: ptr, id: u32) -> bool { + let dt = S.f[2].z; + let spread = S.f[1].w; + var g = 1.0; + if (spread > 0.0) { + g = 1.0 + (ps_h(id, 7u) * 2.0 - 1.0) * spread; + } + let a = S.f[2].xy * g; + let v = ((*q).v.xy + a * dt) * 1.0; + let p = (*q).p.xy + v * dt; + (*q).v = vec3(v, 0.0); + (*q).p = vec3(p, 0.0); + let b = S.f[3]; + return p.x > b.x && p.x < b.z && p.y > b.y && p.y < b.w; +} + +@compute @workgroup_size(64, 1) +fn particles(@builtin(global_invocation_id) gid: vec3) { + let id = gid.x; + if (id >= S.u.x) { + return; + } + let s0 = S.u.y; + let s1 = S.u.z; + let cannon = S.u.w == 1u; + let b = births[id]; + var q: Part; + var j: u32; + if (b < s0 && id < S.u2.w) { + let o = prev[id]; + if (o.m.y == 0.0) { + next[id] = o; + return; + } + q = Part(o.p.xyz, o.v.xyz, o.p.w, o.v.w, o.m.x); + j = s0; + } else { + if (cannon) { + q = ps_spawn_cannon(id); + } else { + q = ps_spawn_engine(id); + } + j = b + 1u; + } + var alive = true; + for (; j < s1; j++) { + if (cannon) { + if (!ps_step_cannon(&q, id)) { + alive = false; + break; + } + } else { + q = ps_step_engine(q, j); + if (q.age >= q.life) { + alive = false; + break; + } + } + } + next[id] = PState(vec4(q.p, q.age), vec4(q.v, q.life), vec4(q.rnd, select(0.0, 1.0, alive), 0.0, 0.0)); +} diff --git a/crates/gpu/src/tests_particles.rs b/crates/gpu/src/tests_particles.rs new file mode 100644 index 000000000..e3ef97a82 --- /dev/null +++ b/crates/gpu/src/tests_particles.rs @@ -0,0 +1,155 @@ +//! GPU particles vs the CPU simulation (the oracle): CC Particle World, CC Particle Systems II +//! and Particle Playground's cannon at fixed seeds must give the same live particles (same +//! ids) at positions within tolerance, from scratch and when advancing / seeking from GPU +//! checkpoints. + +use effectcraft_effects::psim::{ParticleSim, SimParticle}; +use effectcraft_effects::{EffectCtx, EffectEnv, EffectHost, LayerPixels, Params}; +use effectcraft_keyframe::Value; + +use crate::Gpu; +use crate::tests::{gpu, n}; + +struct Host(&'static Gpu); + +impl EffectHost for Host { + fn layer(&self, _id: u64, _masks_and_effects: bool) -> Option { + None + } + fn audio(&self, _id: u64, _start: f64, _frames: usize, _rate: u32) -> Option> { + None + } + fn particles(&self) -> Option<&dyn ParticleSim> { + Some(self.0) + } +} + +fn params(id: &str, vals: &[(&str, Value)]) -> Params { + let spec = effectcraft_effects::find(id).unwrap(); + let size = [320.0, 240.0]; + let mut p = Params { values: spec.params.iter().map(|ps| (ps.id.to_string(), effectcraft_effects::default_value(ps, size))).collect() }; + for (k, v) in vals { + assert!(p.values.contains_key(*k), "{id}: no {k}"); + p.values.insert(k.to_string(), v.clone()); + } + p +} + +/// (CPU, GPU) particles of effect `id` at layer time `t`. +fn states(id: &str, p: &Params, t: f64, host: &Host) -> (Vec, Vec) { + let ctx = |env| EffectCtx { params: p, time: t, layer_size: [320.0, 240.0], seed: 5, adjustment: false, env }; + let run = + |c: &EffectCtx| if id == "ec.sim.particleplayground" { effectcraft_effects::playground_state(c) } else { effectcraft_effects::particle_state(id, c) }; + let cpu = run(&ctx(EffectEnv::default())).expect("cpu state"); + let gpu = run(&ctx(EffectEnv { host: Some(host), ..Default::default() })).expect("gpu state"); + (cpu, gpu) +} + +/// Same ids alive (up to `id_slack` particles dying a step apart), positions within `tol` +/// (world units) for 99 % of them and on average within `tol / 10`. +fn agree(label: &str, cpu: &[SimParticle], gpu: &[SimParticle], tol: f32, id_slack: usize) { + assert!(cpu.len() > 20, "{label}: the scene has particles ({})", cpu.len()); + let gi: std::collections::HashMap = gpu.iter().map(|q| (q.id, q)).collect(); + let mut missing = 0; + let mut errs = vec![]; + for q in cpu { + match gi.get(&q.id) { + Some(g) => errs.push((0..3).map(|k| (q.p[k] - g.p[k]).abs()).fold(0.0f32, f32::max)), + None => missing += 1, + } + } + let extra = gpu.len() + missing - cpu.len(); + assert!(missing + extra <= id_slack, "{label}: {missing} CPU particles missing on the GPU, {extra} extra (of {})", cpu.len()); + errs.sort_by(f32::total_cmp); + let mean = errs.iter().sum::() / errs.len() as f32; + let p99 = errs[(errs.len() * 99 / 100).min(errs.len() - 1)]; + eprintln!("{label}: {} particles, mean err {mean:.2e}, p99 {p99:.2e}, max {:.2e}", cpu.len(), errs.last().unwrap()); + assert!(p99 <= tol && mean <= tol / 10.0, "{label}: positions differ (mean {mean}, p99 {p99})"); +} + +#[test] +fn particle_world_matches_the_cpu() { + let Some(g) = gpu() else { return }; + let host = Host(g); + for anim in [0u32, 1, 2, 3, 4, 5, 6, 7, 8, 9] { + let p = params( + "ec.sim.ccparticleworld", + &[ + ("physics/animation", Value::Enum(anim)), + ("birthRate", n(2.0)), + ("longevity", n(1.2)), + ("physics/extra", n(0.6)), + ("physics/resistance", n(0.5)), + ("extras/randomSeed", n(7.0)), + ], + ); + let (cpu, gpu) = states("ec.sim.ccparticleworld", &p, 1.5, &host); + // Fractal / fire animations sample value noise along the path: rounding differences can + // grow, so they get a looser bound. + let tol = if matches!(anim, 6 | 8 | 9) { 1e-3 } else { 1e-4 }; + agree(&format!("particle world anim {anim}"), &cpu, &gpu, tol, cpu.len() / 100 + 1); + } +} + +#[test] +fn particle_systems_ii_matches_the_cpu() { + let Some(g) = gpu() else { return }; + let host = Host(g); + for anim in [0u32, 1, 2, 4, 5, 8] { + let p = params( + "ec.sim.ccparticlesystems2", + &[("physics/animation", Value::Enum(anim)), ("birthRate", n(3.0)), ("longevity", n(1.5)), ("physics/gravity", n(0.8)), ("randomSeed", n(3.0))], + ); + let (cpu, gpu) = states("ec.sim.ccparticlesystems2", &p, 2.0, &host); + let tol = if anim == 8 { 1e-3 } else { 1e-4 }; + agree(&format!("particle systems ii anim {anim}"), &cpu, &gpu, tol, cpu.len() / 100 + 1); + } +} + +#[test] +fn playground_cannon_matches_the_cpu_and_checkpoints_reuse() { + let Some(g) = gpu() else { return }; + let host = Host(g); + for (radius, spread) in [(4.0, 0.0), (-10.0, 50.0)] { + let p = params( + "ec.sim.particleplayground", + &[ + ("grid/particlesAcross", n(0.0)), + ("cannon/barrelRadius", n(radius)), + ("cannon/particlesPerSecond", n(90.0)), + ("cannon/directionRandomSpread", n(30.0)), + ("cannon/velocityRandomSpread", n(40.0)), + ("gravity/gravityForceRandomSpread", n(spread)), + ], + ); + // Advance through checkpoints (0.5 → 1.0 → 2.0), seek back (1.5), compare each. + for t in [0.5, 1.0, 2.0, 1.5] { + let (cpu, gpu) = states("ec.sim.particleplayground", &p, t, &host); + // Positions are layer pixels here. + agree(&format!("cannon r {radius} spread {spread} t {t}"), &cpu, &gpu, 0.01, cpu.len() / 100 + 1); + } + } +} + +/// Through the renderer: the GPU backend's effect host hands the simulation to the GPU, and +/// the composited frames agree with the CPU render. +#[test] +fn particle_effects_render_like_the_cpu() { + use effectcraft_project::BitDepth; + use effectcraft_time::Tick; + + use crate::tests::{Scene, check, compare_at, opts}; + for (id, vals) in [ + ("ec.sim.ccparticleworld", vec![("birthRate", n(1.0)), ("physics/animation", Value::Enum(4))]), + ("ec.sim.ccparticlesystems2", vec![("birthRate", n(2.0))]), + ("ec.sim.particleplayground", vec![("grid/particlesAcross", n(0.0)), ("cannon/particlesPerSecond", n(60.0))]), + ] { + for depth in [BitDepth::Bpc8, BitDepth::Bpc32] { + let mut s = Scene::new(depth); + let mut l = s.solid([0.1, 0.1, 0.15], 97, 61); + s.effect(&mut l, id, &vals); + s.push(l); + check(&format!("{id} {depth:?}"), compare_at(&s, opts(), Tick::from_seconds_f64(1.0)), 0.002); + } + } +} diff --git a/crates/render/src/lib.rs b/crates/render/src/lib.rs index ffb559369..44f4876a0 100644 --- a/crates/render/src/lib.rs +++ b/crates/render/src/lib.rs @@ -82,6 +82,9 @@ struct FxHost<'r, 'a, 'c> { const MAX_FX_DEPTH: usize = 8; impl EffectHost for FxHost<'_, '_, '_> { + fn particles(&self) -> Option<&dyn effectcraft_effects::psim::ParticleSim> { + self.r.active_accel().and_then(|a| a.particles()) + } fn layer(&self, id: u64, masks_and_effects: bool) -> Option { let other = self.ctx.layer(effectcraft_project::LayerId(id))?; if other.id == self.layer.id || self.r.depth > MAX_FX_DEPTH { @@ -240,6 +243,11 @@ pub trait Accelerator: Send + Sync { fn raster_3d(&self, _scene: &three_d::adv::Scene) -> Option { None } + /// A particle simulation backend (GPU particles) for the stepped particle effects, with + /// the CPU simulation's semantics. `None` = they simulate on the CPU. + fn particles(&self) -> Option<&dyn effectcraft_effects::psim::ParticleSim> { + None + } } /// Where an effect stack runs (see [`Renderer::run_effects_on`]): a CPU buffer, or an image