mirror of
https://github.com/storytold/effectcraft.git
synced 2026-10-29 04:26:55 -06:00
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 <noreply@anthropic.com>
This commit is contained in:
Brandon Thomas
co-authored by
Claude Opus 5.5
Claude Opus 5.5
parent
9a6389db86
commit
0e3b354520
@@ -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<std::sync::Arc<AuxChannels>> {
|
||||
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").
|
||||
|
||||
@@ -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<Vec<SimParticle>>;
|
||||
}
|
||||
|
||||
/// 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<Vec<u32>> {
|
||||
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;
|
||||
@@ -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<std::sync::Arc<PState>> {
|
||||
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<Vec<crate::psim::SimParticle>> {
|
||||
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<PState>, key: u64, t: f64, phys: &Phys) -> std::sync::Arc<PState> {
|
||||
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<PState> = 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<PState> = 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");
|
||||
|
||||
@@ -231,6 +231,38 @@ impl Pg {
|
||||
|
||||
static PG_CACHE: SimCache<PgState> = 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<std::sync::Arc<PgState>> {
|
||||
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<Vec<crate::psim::SimParticle>> {
|
||||
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<Sprite> = st
|
||||
.parts
|
||||
|
||||
@@ -102,6 +102,9 @@ pub struct GpuContext {
|
||||
uploads: Mutex<Uploads>,
|
||||
/// Advanced 3D render pipelines (built on first use).
|
||||
pub(crate) adv3d: std::sync::OnceLock<crate::adv3d::Pipes>,
|
||||
/// 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(),
|
||||
})
|
||||
}
|
||||
|
||||
|
||||
@@ -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<effectcraft_render::three_d::adv::Target> {
|
||||
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<Vec<effectcraft_effects::psim::SimParticle>> {
|
||||
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;
|
||||
|
||||
@@ -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<Checkpoint>,
|
||||
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<u8> {
|
||||
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<Vec<SimParticle>> {
|
||||
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<u8> = 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<f32> = 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<Pipes>;
|
||||
pub(crate) type StatesCell = Mutex<Checkpoints>;
|
||||
@@ -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<u32>,
|
||||
// Animation, flat, seed, particles in `prev`.
|
||||
u2: vec4<u32>,
|
||||
// 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<vec4<f32>, 6>,
|
||||
}
|
||||
|
||||
struct PState {
|
||||
// Position, age.
|
||||
p: vec4<f32>,
|
||||
// Velocity, life.
|
||||
v: vec4<f32>,
|
||||
// rnd, alive (1 / 0), id.
|
||||
m: vec4<f32>,
|
||||
}
|
||||
|
||||
@group(1) @binding(0) var<uniform> S: PSim;
|
||||
@group(1) @binding(1) var<storage, read> births: array<u32>;
|
||||
@group(1) @binding(2) var<storage, read> prev: array<PState>;
|
||||
@group(1) @binding(3) var<storage, read_write> next: array<PState>;
|
||||
|
||||
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<f32>) -> vec3<f32> {
|
||||
let l = sqrt(v.x * v.x + v.y * v.y + v.z * v.z);
|
||||
if (l > 1e-9) {
|
||||
return v / l;
|
||||
}
|
||||
return vec3<f32>(0.0, -1.0, 0.0);
|
||||
}
|
||||
|
||||
// sim::rand_dir.
|
||||
fn ps_rand_dir(i: u32, k: u32, flat: bool) -> vec3<f32> {
|
||||
if (flat) {
|
||||
let a = ps_h(i, k) * PS_TAU;
|
||||
return vec3<f32>(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<f32>(r * cos(a), r * sin(a), z);
|
||||
}
|
||||
|
||||
struct Part {
|
||||
p: vec3<f32>,
|
||||
v: vec3<f32>,
|
||||
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<f32>(jitter.x, 0.0, z);
|
||||
}
|
||||
case 6u: {
|
||||
dir = ps_norm3(vec3<f32>(jitter.x * 0.25, -1.0, jitter.z * 0.25));
|
||||
}
|
||||
case 7u: {
|
||||
dir = ps_norm3(vec3<f32>(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<f32>(sin(a), -cos(a));
|
||||
let speed = max(vel + ps_hs(id, 3u) * vel_spread, 0.0);
|
||||
var off: vec2<f32>;
|
||||
if (r >= 0.0) {
|
||||
off = vec2<f32>(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<f32>(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<f32>(p, 0.0), vec3<f32>(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<function, Part>, 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<f32>(v, 0.0);
|
||||
(*q).p = vec3<f32>(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<u32>) {
|
||||
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<f32>(q.p, q.age), vec4<f32>(q.v, q.life), vec4<f32>(q.rnd, select(0.0, 1.0, alive), 0.0, 0.0));
|
||||
}
|
||||
@@ -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<LayerPixels> {
|
||||
None
|
||||
}
|
||||
fn audio(&self, _id: u64, _start: f64, _frames: usize, _rate: u32) -> Option<Vec<f32>> {
|
||||
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<SimParticle>, Vec<SimParticle>) {
|
||||
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<u32, &SimParticle> = 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::<f32>() / 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);
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -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<LayerPixels> {
|
||||
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<three_d::adv::Target> {
|
||||
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
|
||||
|
||||
Reference in New Issue
Block a user