From bc50b98d8834532c3a63fbfc37561ccaeedfa571 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?charlotte=20=F0=9F=8C=B8?= Date: Mon, 28 Sep 2026 21:56:40 -0700 Subject: [PATCH] Flocking from generic ops. --- crates/processing_ffi/src/lib.rs | 8 -- .../processing_pyo3/examples/flocking_duck.py | 2 +- .../processing_pyo3/examples/flocking_gpu.py | 76 ++++++----- crates/processing_pyo3/src/particles.rs | 30 ----- crates/processing_render/src/lib.rs | 4 +- .../processing_render/src/particles/emit.rs | 28 ---- .../src/particles/kernels/flock.wgsl | 125 ------------------ .../src/particles/kernels/mod.rs | 16 --- crates/processing_render/src/particles/mod.rs | 10 +- 9 files changed, 51 insertions(+), 248 deletions(-) delete mode 100644 crates/processing_render/src/particles/kernels/flock.wgsl diff --git a/crates/processing_ffi/src/lib.rs b/crates/processing_ffi/src/lib.rs index 6eec696..f78a866 100644 --- a/crates/processing_ffi/src/lib.rs +++ b/crates/processing_ffi/src/lib.rs @@ -3694,14 +3694,6 @@ pub extern "C" fn processing_particles_kernel_impulse() -> u64 { .unwrap_or(0) } -#[unsafe(no_mangle)] -pub extern "C" fn processing_particles_kernel_flock() -> u64 { - error::clear_error(); - error::check(particles_kernel_flock) - .map(|e| e.to_bits()) - .unwrap_or(0) -} - #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_orient() -> u64 { error::clear_error(); diff --git a/crates/processing_pyo3/examples/flocking_duck.py b/crates/processing_pyo3/examples/flocking_duck.py index e60ff4f..f53d103 100644 --- a/crates/processing_pyo3/examples/flocking_duck.py +++ b/crates/processing_pyo3/examples/flocking_duck.py @@ -1,4 +1,4 @@ -# Flocking inside a duck: the GPU boids from flocking_gpu.py, seeded from +# Flocking inside a duck: GPU boids seeded from # the vertices of the Duck glTF mesh. Each boid remembers its spawn vertex # in a `home` attribute; a homing force that is negligible near home but # grows quadratically with distance lets the boids swirl and flock locally diff --git a/crates/processing_pyo3/examples/flocking_gpu.py b/crates/processing_pyo3/examples/flocking_gpu.py index ab8ec57..4b33a5a 100644 --- a/crates/processing_pyo3/examples/flocking_gpu.py +++ b/crates/processing_pyo3/examples/flocking_gpu.py @@ -1,13 +1,10 @@ -# GPU flocking: the boids from flocking.py, moved entirely onto the GPU. -# Positions and velocities live in particle attribute buffers, two compute -# kernels update them each frame, and the flock renders instanced — nothing -# is ever read back to the CPU. Brute-force O(N²) neighbor search is trivial -# for a GPU at this scale; a spatial hash grid is the next step past ~100k. +# GPU flocking: the boids from flocking.py, built from generic particle ops +# over each boid's neighbors. Nothing is read back to the CPU. from mewnala import * from math import cos, sin from random import uniform -BOID_COUNT = 10000 +BOID_COUNT = 100000 BOUND = 30.0 # half-extent of the wrapping box NEIGHBOR_DIST = 5.0 SEPARATION_DIST = 2.5 @@ -15,12 +12,12 @@ MAX_FORCE = 6.0 # units per second² DT = 1.0 / 60.0 -# Pass 1: every boid reads the whole flock's state and writes only its -# steering force. Splitting the read from the write mirrors the CPU -# example's two loops — no boid sees a half-updated neighbor. - -# Pass 2: integrate the steering force, wrap at the box edges, and point -# each instanced boid along its velocity via the rotation quaternion. +# Separation steers away from the summed offsets, so its speed is negative. +RULES = [ + ("separation", -MAX_SPEED, 1.5, "close"), + ("alignment", MAX_SPEED, 1.0, "near"), + ("cohesion", MAX_SPEED, 1.0, "near"), +] p = None boid = None @@ -60,15 +57,7 @@ def setup(): directional_light((0.95, 0.9, 0.85), 800.0) - p = create_particles( - BOID_COUNT, - attributes=[ - Attribute.position(), - Attribute.rotation(), - Attribute.color(), - Attribute.velocity(), - ], - ) + p = create_particles(BOID_COUNT) positions = [] velocities = [] @@ -98,6 +87,38 @@ def setup(): ) +# Reynolds steering. `mask` zeroes it for boids with no neighbors in range, +# which would otherwise brake. +def steer(desired, speed, mask): + p.apply(MAP, desired, op=NORMALIZE, length=speed) + p.apply(COMBINE, desired, "velocity", op=SUB) + p.apply(MAP, desired, op=LIMIT, max_length=MAX_FORCE * DT) + p.apply(COMBINE, desired, mask, op=MUL) + + +def flock(): + p.apply(FIND_NEIGHBORS, grid=grid, radius=NEIGHBOR_DIST) + # what each rule steers toward, summed over the neighbors + p.apply(NEIGHBOR, "position", out="separation", op=SUM, + relative=True, radius=SEPARATION_DIST, falloff=INVERSE_SQUARE) + p.apply(NEIGHBOR, "velocity", out="alignment", op=SUM) + p.apply(NEIGHBOR, "position", out="cohesion", op=SUM, relative=True) + # which boids have any neighbors for each rule + p.apply(NEIGHBOR, out="close", op=COUNT, radius=SEPARATION_DIST) + p.apply(NEIGHBOR, out="near", op=COUNT) + p.apply(MAP, "close", op=GREATER, threshold=0) + p.apply(MAP, "near", op=GREATER, threshold=0) + + for i, (rule, speed, weight, mask) in enumerate(RULES): + steer(rule, speed, mask) + if i == 0: + p.apply(MAP, rule, out="force", op=AFFINE, scale=weight) + else: + p.apply(COMBINE, "force", rule, op="add", b_scale=weight) + p.apply(COMBINE, "velocity", "force", op="add") + p.apply(MAP, "velocity", op=LIMIT, min_length=MAX_SPEED * 0.25, max_length=MAX_SPEED) + + def draw(): global title_last_time, title_last_frame @@ -114,20 +135,11 @@ def draw(): camera_look_at(0.0, 0.0, 0.0) background(10, 12, 18) + flock() + material(mat) particles(p, boid) - p.flock( - grid, - sep_distance=SEPARATION_DIST, - neighbor_distance=NEIGHBOR_DIST, - weight_separation=1.5, - weight_alignment=1.0, - weight_cohesion=1.0, - max_speed=MAX_SPEED, - max_force=MAX_FORCE * DT, - min_speed=MAX_SPEED * 0.25, - ) p.apply(INTEGRATE, dt=DT) p.apply(BOUNDS_BOX, aabb_min=[-BOUND] * 3, aabb_max=[BOUND] * 3, mode=2) p.apply(ORIENT, forward=[0.0, 0.0, 1.0], up=[0.0, 1.0, 0.0]) diff --git a/crates/processing_pyo3/src/particles.rs b/crates/processing_pyo3/src/particles.rs index ea270ed..788f0d6 100644 --- a/crates/processing_pyo3/src/particles.rs +++ b/crates/processing_pyo3/src/particles.rs @@ -150,18 +150,6 @@ impl Drop for Grid { } } -static FLOCK_COMPUTE: std::sync::Mutex> = std::sync::Mutex::new(None); - -fn flock_compute() -> PyResult { - let mut guard = FLOCK_COMPUTE.lock().unwrap(); - if let Some(e) = *guard { - return Ok(e); - } - let e = particles_kernel_flock().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; - *guard = Some(e); - Ok(e) -} - static PHYSICS_COMPUTES: std::sync::Mutex>> = std::sync::Mutex::new(None); @@ -1113,24 +1101,6 @@ impl Particles { Ok(Grid { entity }) } - #[pyo3(signature = (grid, **kwargs))] - pub fn flock(&self, grid: &Grid, kwargs: Option<&Bound<'_, PyDict>>) -> PyResult<()> { - let flock = flock_compute()?; - if let Some(kwargs) = kwargs { - crate::compute::set_compute_kwargs(flock, kwargs)?; - } - let cell = grid.cell_size()?; - let neighbor_distance = kw_f32(kwargs, "neighbor_distance", cell)?.min(cell); - compute_set( - flock, - "neighbor_distance", - shader_value::ShaderValue::Float(neighbor_distance), - ) - .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; - particles_flock(self.entity, flock, grid.entity) - .map_err(|e| PyRuntimeError::new_err(format!("{e}"))) - } - #[staticmethod] pub fn orient() -> PyResult { let entity = diff --git a/crates/processing_render/src/lib.rs b/crates/processing_render/src/lib.rs index d215ea8..6bbfe67 100644 --- a/crates/processing_render/src/lib.rs +++ b/crates/processing_render/src/lib.rs @@ -41,11 +41,11 @@ pub use particles::{ particles_apply, particles_attribute_add, particles_attributes, particles_buffer, particles_capacity, particles_connectivity_indirect, particles_create, particles_create_from_geometry, particles_destroy, particles_emit, particles_emit_gpu, - particles_ensure_attribute, particles_flock, particles_gather, particles_kernel_age, + particles_ensure_attribute, particles_gather, particles_kernel_age, particles_kernel_attr_combine, particles_kernel_attr_linear, particles_kernel_attr_lookup1d, particles_kernel_attr_lookup2d, particles_kernel_attr_mix, particles_kernel_attract, particles_kernel_bounds_box, particles_kernel_bounds_geometry, particles_kernel_bounds_sphere, - particles_kernel_drag, particles_kernel_field, particles_kernel_flock, particles_kernel_force, + particles_kernel_drag, particles_kernel_field, particles_kernel_force, particles_kernel_impulse, particles_kernel_integrate, particles_kernel_noise, particles_kernel_orient, particles_kernel_transform, particles_kernel_vortex, particles_reset_indices, particles_scatter_create, particles_scatter_volume_create, diff --git a/crates/processing_render/src/particles/emit.rs b/crates/processing_render/src/particles/emit.rs index fb5f30e..51cd91a 100644 --- a/crates/processing_render/src/particles/emit.rs +++ b/crates/processing_render/src/particles/emit.rs @@ -143,34 +143,6 @@ pub fn particles_emit( }) } -pub fn particles_flock( - particles_entity: Entity, - flock_entity: Entity, - grid: Entity, -) -> error::Result<()> { - let position = app_mut(|app| { - let world = app.world(); - let field = world - .get::(particles_entity) - .ok_or(error::ProcessingError::ParticlesNotFound)?; - for (&attr_entity, &buf_entity) in &field.buffers { - let attr = world - .get::(attr_entity) - .ok_or(error::ProcessingError::InvalidEntity)?; - if attr.name == "position" { - return Ok(buf_entity); - } - } - Err(error::ProcessingError::InvalidArgument( - "particles_flock requires a `position` attribute".to_string(), - )) - })?; - - grid_build(grid, position)?; - compute_set(flock_entity, "grid", ShaderValue::Grid(grid))?; - particles_apply(particles_entity, flock_entity) -} - static NEIGHBOR_COMPUTE: Mutex> = Mutex::new(None); fn neighbor_compute() -> error::Result { diff --git a/crates/processing_render/src/particles/kernels/flock.wgsl b/crates/processing_render/src/particles/kernels/flock.wgsl deleted file mode 100644 index b1e75ec..0000000 --- a/crates/processing_render/src/particles/kernels/flock.wgsl +++ /dev/null @@ -1,125 +0,0 @@ -import processing::particles::{Grid, cell_coords, cell_index}; - -struct FlockParams { - sep_distance: f32, - neighbor_distance: f32, - weight_separation: f32, - weight_alignment: f32, - weight_cohesion: f32, - max_speed: f32, - max_force: f32, - min_speed: f32, -} - -@group(0) @binding(0) var position: array; -@group(0) @binding(1) var velocity: array; -@group(0) @binding(2) var grid_offsets: array; -@group(0) @binding(3) var grid_sorted: array; -@group(0) @binding(4) var fp: FlockParams; -@group(0) @binding(5) var grid: Grid; - -fn limit_mag(v: vec3, m: f32) -> vec3 { - let len2 = dot(v, v); - if len2 > m * m { return v * (m * inverseSqrt(len2)); } - return v; -} - -fn steer_toward(desired: vec3, vel: vec3, max_speed: f32, max_force: f32) -> vec3 { - let m2 = dot(desired, desired); - if m2 < 0.00000001 { return vec3(0.0); } - return limit_mag(desired * (max_speed * inverseSqrt(m2)) - vel, max_force); -} - -fn load_pos(i: u32) -> vec3 { - return vec3(position[i * 3u], position[i * 3u + 1u], position[i * 3u + 2u]); -} - -fn load_vel(i: u32) -> vec3 { - return vec3(velocity[i * 3u], velocity[i * 3u + 1u], velocity[i * 3u + 2u]); -} - -@compute @workgroup_size(64) -fn main(@builtin(global_invocation_id) gid: vec3) { - let i = gid.x; - let count = arrayLength(&position) / 3u; - if i >= count { return; } - - let pos = load_pos(i); - let vel = load_vel(i); - - let sep_d2 = fp.sep_distance * fp.sep_distance; - let neighbor_d2 = fp.neighbor_distance * fp.neighbor_distance; - - let dims = grid.dims; - let base = cell_coords(pos, grid.origin, grid.cell_size, dims); - let bx = base.x; - let by = base.y; - let bz = base.z; - - var sep_steer = vec3(0.0); - var sep_count = 0u; - var ali_sum = vec3(0.0); - var coh_sum = vec3(0.0); - var flock_count = 0u; - - for (var dz = -1; dz <= 1; dz++) { - let cz = bz + dz; - if cz < 0 || cz >= i32(grid.dims.z) { continue; } - for (var dy = -1; dy <= 1; dy++) { - let cy = by + dy; - if cy < 0 || cy >= i32(grid.dims.y) { continue; } - for (var dx = -1; dx <= 1; dx++) { - let cx = bx + dx; - if cx < 0 || cx >= i32(grid.dims.x) { continue; } - - let cell = cell_index(vec3(u32(cx), u32(cy), u32(cz)), dims); - let start = grid_offsets[cell]; - let end = grid_offsets[cell + 1u]; - for (var s = start; s < end; s++) { - let j = grid_sorted[s]; - if j == i { continue; } - - let diff = pos - load_pos(j); - let d2 = dot(diff, diff); - if d2 > 0.000001 && d2 < neighbor_d2 { - if d2 < sep_d2 { - sep_steer += diff / d2; - sep_count += 1u; - } - ali_sum += load_vel(j); - coh_sum += diff; - flock_count += 1u; - } - } - } - } - } - - var force = vec3(0.0); - if sep_count > 0u { - force += steer_toward(sep_steer / f32(sep_count), vel, - fp.max_speed, fp.max_force) * fp.weight_separation; - } - if flock_count > 0u { - force += steer_toward(ali_sum / f32(flock_count), vel, - fp.max_speed, fp.max_force) * fp.weight_alignment; - force += steer_toward(-coh_sum / f32(flock_count), vel, - fp.max_speed, fp.max_force) * fp.weight_cohesion; - } - - var new_vel = vel + force; - let speed2 = dot(new_vel, new_vel); - let max_speed_sq = fp.max_speed * fp.max_speed; - if speed2 > max_speed_sq { - new_vel = new_vel * (fp.max_speed * inverseSqrt(speed2)); - } else if fp.min_speed > 0.0 { - let min_speed_sq = fp.min_speed * fp.min_speed; - if speed2 < min_speed_sq && speed2 > 0.0 { - new_vel = new_vel * sqrt(min_speed_sq / speed2); - } - } - - velocity[i * 3u] = new_vel.x; - velocity[i * 3u + 1u] = new_vel.y; - velocity[i * 3u + 2u] = new_vel.z; -} diff --git a/crates/processing_render/src/particles/kernels/mod.rs b/crates/processing_render/src/particles/kernels/mod.rs index df9199f..7cb8c83 100644 --- a/crates/processing_render/src/particles/kernels/mod.rs +++ b/crates/processing_render/src/particles/kernels/mod.rs @@ -42,7 +42,6 @@ impl Plugin for ParticlesKernelsPlugin { embedded_asset!(app, "bounds_sphere.wgsl"); embedded_asset!(app, "bounds_box.wgsl"); embedded_asset!(app, "impulse.wgsl"); - embedded_asset!(app, "flock.wgsl"); embedded_asset!(app, "orient.wgsl"); embedded_asset!(app, "field.wgsl"); embedded_asset!(app, "attr_linear.wgsl"); @@ -255,21 +254,6 @@ pub fn particles_kernel_impulse() -> error::Result { Ok(entity) } -pub fn particles_kernel_flock() -> error::Result { - let shader = shader_load("embedded://processing_render/particles/kernels/flock.wgsl")?; - let entity = compute_create(shader)?; - set_requires(entity, &["position", "velocity"])?; - compute_set(entity, "sep_distance", ShaderValue::Float(1.2))?; - compute_set(entity, "neighbor_distance", ShaderValue::Float(2.5))?; - compute_set(entity, "weight_separation", ShaderValue::Float(1.5))?; - compute_set(entity, "weight_alignment", ShaderValue::Float(1.0))?; - compute_set(entity, "weight_cohesion", ShaderValue::Float(1.0))?; - compute_set(entity, "max_speed", ShaderValue::Float(0.1))?; - compute_set(entity, "max_force", ShaderValue::Float(0.003))?; - compute_set(entity, "min_speed", ShaderValue::Float(0.02))?; - Ok(entity) -} - pub fn particles_kernel_orient() -> error::Result { let shader = shader_load("embedded://processing_render/particles/kernels/orient.wgsl")?; let entity = compute_create(shader)?; diff --git a/crates/processing_render/src/particles/mod.rs b/crates/processing_render/src/particles/mod.rs index f0b2031..e6013b4 100644 --- a/crates/processing_render/src/particles/mod.rs +++ b/crates/processing_render/src/particles/mod.rs @@ -21,9 +21,7 @@ pub use algebra::{ REDUCE_SUMSQ, combine, extract, generate, lookup, map, mix, pack, reduce_components, }; pub use compact::compact; -pub use emit::{ - particles_apply, particles_emit, particles_emit_gpu, particles_flock, particles_gather, -}; +pub use emit::{particles_apply, particles_emit, particles_emit_gpu, particles_gather}; pub use grid::{Grid, GridParams, grid_build, grid_create, grid_destroy, grid_get}; pub use kernels::{ BOUNDS_CLAMP, BOUNDS_REFLECT, BOUNDS_SOFT, BOUNDS_WRAP, COMBINE_ADD, COMBINE_DIV, COMBINE_MAX, @@ -33,9 +31,9 @@ pub use kernels::{ particles_kernel_attr_lookup1d, particles_kernel_attr_lookup2d, particles_kernel_attr_mix, particles_kernel_attract, particles_kernel_bounds_box, particles_kernel_bounds_geometry, particles_kernel_bounds_sphere, particles_kernel_drag, particles_kernel_field, - particles_kernel_flock, particles_kernel_force, particles_kernel_impulse, - particles_kernel_integrate, particles_kernel_noise, particles_kernel_orient, - particles_kernel_transform, particles_kernel_vortex, + particles_kernel_force, particles_kernel_impulse, particles_kernel_integrate, + particles_kernel_noise, particles_kernel_orient, particles_kernel_transform, + particles_kernel_vortex, }; pub use neighbors::{ particles_find_neighbors, particles_neighbor_lists, particles_neighbor_reduce,