From d53b8627a5abffb610c2697a1fd86b3946a984b8 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?charlotte=20=F0=9F=8C=B8?= Date: Mon, 28 Sep 2026 21:32:24 -0700 Subject: [PATCH] GPU test harnesses. --- Cargo.toml | 44 +++++++++++ examples/combine_mix_test.rs | 109 ++++++++++++++++++++++++++ examples/compute_write_order.rs | 128 ++++++++++++++++++++++++++++++ examples/field_falloff_test.rs | 98 +++++++++++++++++++++++ examples/glue_test.rs | 110 ++++++++++++++++++++++++++ examples/grid_test.rs | 133 ++++++++++++++++++++++++++++++++ examples/group_compact_test.rs | 79 +++++++++++++++++++ examples/lookup_test.rs | 100 ++++++++++++++++++++++++ examples/map_test.rs | 77 ++++++++++++++++++ examples/reduce_test.rs | 71 +++++++++++++++++ examples/scan_test.rs | 68 ++++++++++++++++ examples/sort_test.rs | 91 ++++++++++++++++++++++ 12 files changed, 1108 insertions(+) create mode 100644 examples/combine_mix_test.rs create mode 100644 examples/compute_write_order.rs create mode 100644 examples/field_falloff_test.rs create mode 100644 examples/glue_test.rs create mode 100644 examples/grid_test.rs create mode 100644 examples/group_compact_test.rs create mode 100644 examples/lookup_test.rs create mode 100644 examples/map_test.rs create mode 100644 examples/reduce_test.rs create mode 100644 examples/scan_test.rs create mode 100644 examples/sort_test.rs diff --git a/Cargo.toml b/Cargo.toml index be9f10df..154537bf 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -212,6 +212,50 @@ path = "examples/camera_controllers.rs" name = "compute_readback" path = "examples/compute_readback.rs" +[[example]] +name = "scan_test" +path = "examples/scan_test.rs" + +[[example]] +name = "grid_test" +path = "examples/grid_test.rs" + +[[example]] +name = "map_test" +path = "examples/map_test.rs" + +[[example]] +name = "combine_mix_test" +path = "examples/combine_mix_test.rs" + +[[example]] +name = "lookup_test" +path = "examples/lookup_test.rs" + +[[example]] +name = "glue_test" +path = "examples/glue_test.rs" + +[[example]] +name = "sort_test" +path = "examples/sort_test.rs" + +[[example]] +name = "field_falloff_test" +path = "examples/field_falloff_test.rs" + +[[example]] +name = "group_compact_test" +path = "examples/group_compact_test.rs" + +[[example]] +name = "reduce_test" +path = "examples/reduce_test.rs" + +[[example]] +name = "compute_write_order" +path = "examples/compute_write_order.rs" + [[example]] name = "alias_spike" path = "examples/alias_spike.rs" diff --git a/examples/combine_mix_test.rs b/examples/combine_mix_test.rs new file mode 100644 index 00000000..6d3bb7df --- /dev/null +++ b/examples/combine_mix_test.rs @@ -0,0 +1,109 @@ +//! `combine` and `mix`, including the alias guard. + +use processing::prelude::*; + +fn f32s(bytes: &[u8]) -> Vec { + bytes + .chunks_exact(4) + .map(|c| f32::from_le_bytes([c[0], c[1], c[2], c[3]])) + .collect() +} +fn to_bytes(v: &[f32]) -> Vec { + v.iter().flat_map(|f| f.to_le_bytes()).collect() +} +fn approx(a: &[f32], b: &[f32]) -> bool { + a.len() == b.len() && a.iter().zip(b).all(|(x, y)| (x - y).abs() < 1e-5) +} + +fn sketch() -> error::Result { + init(Config::default())?; + let surface = surface_create_offscreen(1, 1, 1.0, TextureFormat::Rgba8Unorm)?; + let _graphics = graphics_create(surface, 1, 1, TextureFormat::Rgba8Unorm)?; + + let mut ok = true; + + let av: Vec = (0..12).map(|i| i as f32).collect(); + let bv: Vec = (0..12).map(|i| (i * 10) as f32).collect(); + let a = buffer_create_with_data(to_bytes(&av))?; + let b = buffer_create_with_data(to_bytes(&bv))?; + combine(a, a, b, 3, COMBINE_ADD, 1.0, 0.0)?; + let got = f32s(&buffer_read(a)?); + let want: Vec = av.iter().zip(&bv).map(|(x, y)| x + y).collect(); + if approx(&got, &want) { + println!(" PASS combine in-place a+=b"); + } else { + ok = false; + println!(" FAIL combine in-place: got {got:?}, want {want:?}"); + } + + let dst = buffer_create_with_data(to_bytes(&vec![0.0; 12]))?; + combine(dst, a, b, 3, COMBINE_MUL, 1.0, 0.0)?; + let got_dst = f32s(&buffer_read(dst)?); + let cur_a = f32s(&buffer_read(a)?); + let want_dst: Vec = want.iter().zip(&bv).map(|(x, y)| x * y).collect(); + if approx(&got_dst, &want_dst) && approx(&cur_a, &want) { + println!(" PASS combine out-of-place dst=a*b (inputs preserved)"); + } else { + ok = false; + println!(" FAIL combine out-of-place: dst={got_dst:?} want {want_dst:?}"); + } + + match combine(b, a, b, 3, COMBINE_ADD, 1.0, 0.0) { + Err(_) => println!(" PASS combine alias guard (dst==b rejected before dispatch)"), + Ok(()) => { + ok = false; + println!(" FAIL combine alias guard: expected Err, got Ok"); + } + } + + buffer_destroy(a)?; + buffer_destroy(b)?; + buffer_destroy(dst)?; + + let comps = 4u32; + let a2: Vec = vec![1.0; 12]; + let b2: Vec = vec![3.0; 12]; + let tv = vec![0.0f32, 0.5, 1.0]; + let am = buffer_create_with_data(to_bytes(&a2))?; + let bm = buffer_create_with_data(to_bytes(&b2))?; + let tm = buffer_create_with_data(to_bytes(&tv))?; + let dm = buffer_create_with_data(to_bytes(&vec![0.0; 12]))?; + mix(dm, am, bm, tm, comps, 1.0, 0.0, true)?; + let got_mix = f32s(&buffer_read(dm)?); + let mut want_mix = vec![1.0f32; 4]; + want_mix.extend([2.0f32; 4]); + want_mix.extend([3.0f32; 4]); + if approx(&got_mix, &want_mix) { + println!(" PASS mix out-of-place (per-particle t broadcast across components)"); + } else { + ok = false; + println!(" FAIL mix out-of-place: got {got_mix:?}, want {want_mix:?}"); + } + + mix(am, am, bm, tm, comps, 1.0, 0.0, true)?; + let got_mix_ip = f32s(&buffer_read(am)?); + if approx(&got_mix_ip, &want_mix) { + println!(" PASS mix in-place"); + } else { + ok = false; + println!(" FAIL mix in-place: got {got_mix_ip:?}, want {want_mix:?}"); + } + + buffer_destroy(am)?; + buffer_destroy(bm)?; + buffer_destroy(tm)?; + buffer_destroy(dm)?; + + Ok(ok) +} + +fn main() { + let ok = sketch().unwrap(); + if ok { + println!("combine_mix_test: ALL PASS"); + exit(0).unwrap(); + } else { + println!("combine_mix_test: FAILURES"); + exit(1).unwrap(); + } +} diff --git a/examples/compute_write_order.rs b/examples/compute_write_order.rs new file mode 100644 index 00000000..bec5f9b7 --- /dev/null +++ b/examples/compute_write_order.rs @@ -0,0 +1,128 @@ +use processing::prelude::*; +use processing_render::geometry::AttributeFormat; +use processing_render::{GridParams, grid_build, grid_create, grid_get}; + +fn main() { + match run() { + Ok(_) => exit(0).unwrap(), + Err(e) => { + eprintln!("{e:?}"); + exit(1).unwrap(); + } + } +} + +fn u32s(bytes: &[u8]) -> Vec { + bytes + .chunks_exact(4) + .map(|c| u32::from_le_bytes([c[0], c[1], c[2], c[3]])) + .collect() +} + +fn f32s(bytes: &[u8]) -> Vec { + bytes + .chunks_exact(4) + .map(|c| f32::from_le_bytes([c[0], c[1], c[2], c[3]])) + .collect() +} + +fn bytes_f32(values: &[f32]) -> Vec { + values.iter().flat_map(|v| v.to_le_bytes()).collect() +} + +fn run() -> error::Result<()> { + init(Config::default())?; + let surface = surface_create_offscreen(1, 1, 1.0, TextureFormat::Rgba8Unorm)?; + let _graphics = graphics_create(surface, 1, 1, TextureFormat::Rgba8Unorm)?; + + // partial writes after a kernel + let double = compute_create(shader_create( + r#" +@group(0) @binding(0) var data: array; +@compute @workgroup_size(4) +fn main(@builtin(global_invocation_id) id: vec3) { + data[id.x] = data[id.x] * 2.0; +} +"#, + )?)?; + let buf = buffer_create_with_data(bytes_f32(&[1.0, 2.0, 3.0, 4.0]))?; + compute_set(double, "data", shader_value::ShaderValue::Buffer(buf))?; + compute_dispatch(double, 1, 1, 1)?; + buffer_write_element(buf, 4, bytes_f32(&[99.0]))?; // asset unsynced: GPU only + assert_eq!(f32s(&buffer_read(buf)?), [2.0, 99.0, 6.0, 8.0]); + buffer_write_element(buf, 0, bytes_f32(&[7.0]))?; // asset synced: GPU and asset + assert_eq!(f32s(&buffer_read(buf)?), [7.0, 99.0, 6.0, 8.0]); + compute_dispatch(double, 1, 1, 1)?; + assert_eq!(f32s(&buffer_read(buf)?), [14.0, 198.0, 12.0, 16.0]); + // a later update must not re-upload over the kernel's output + let _ = buffer_create(4)?; + assert_eq!(f32s(&buffer_read(buf)?), [14.0, 198.0, 12.0, 16.0]); + println!("partial writes: ok"); + + // grid rebuilds must see each CPU write + let params = GridParams { + min: [0.0, 0.0, 0.0], + cell_size: 1.0, + dims: [4, 1, 1], + }; + let n = 64u32; + let grid = grid_create(params, n)?; + let position = buffer_create((n * 3 * 4) as u64)?; + for round in 0..4u32 { + let xs: Vec = (0..n) + .map(|i| ((i * (round + 1)) % 4) as f32 + 0.5) + .collect(); + let packed: Vec = xs.iter().flat_map(|&x| [x, 0.5, 0.5]).collect(); + buffer_write(position, bytes_f32(&packed))?; + grid_build(grid, position)?; + let offsets = u32s(&buffer_read(grid_get(grid)?.offsets)?); + let mut expected = [0u32; 4]; + for x in &xs { + expected[*x as usize] += 1; + } + let got: Vec = (0..4).map(|c| offsets[c + 1] - offsets[c]).collect(); + assert_eq!(got, expected, "round {round}: offsets {offsets:?}"); + } + println!("grid rebuilds after CPU writes: ok"); + + let values: Vec = (0..3000).map(|i| i % 7).collect(); + let scan = buffer_create_with_data(values.iter().flat_map(|v| v.to_le_bytes()).collect())?; + prefix_sum_u32(scan)?; + let got = u32s(&buffer_read(scan)?); + let mut acc = 0u32; + let exclusive: Vec = values + .iter() + .map(|v| { + let s = acc; + acc += v; + s + }) + .collect(); + let mut acc = 0u32; + let inclusive: Vec = values + .iter() + .map(|v| { + acc += v; + acc + }) + .collect(); + assert!( + got == exclusive || got == inclusive, + "scan mismatch: {:?}", + &got[..16] + ); + println!("prefix sum: ok"); + + // written before its GPU buffer exists + let attr = geometry_attribute_create("heat", AttributeFormat::Float)?; + let particles = particles_create(4, vec![attr])?; + let heat = particles_buffer(particles, attr)?.expect("heat buffer"); + buffer_write(heat, bytes_f32(&[1.0, 2.0, 3.0, 4.0]))?; + assert_eq!(f32s(&buffer_read(heat)?), [1.0, 2.0, 3.0, 4.0]); + compute_set(double, "data", shader_value::ShaderValue::Buffer(heat))?; + compute_dispatch(double, 1, 1, 1)?; + assert_eq!(f32s(&buffer_read(heat)?), [2.0, 4.0, 6.0, 8.0]); + println!("write before prepare: ok"); + + Ok(()) +} diff --git a/examples/field_falloff_test.rs b/examples/field_falloff_test.rs new file mode 100644 index 00000000..afb4a268 --- /dev/null +++ b/examples/field_falloff_test.rs @@ -0,0 +1,98 @@ +//! Shared `falloff` WESL helper against CPU falloff at known distances. + +use bevy::prelude::Entity; +use processing::prelude::*; + +const RADIUS: f32 = 2.0; + +fn f32s(b: &[u8]) -> Vec { + b.chunks_exact(4) + .map(|c| f32::from_le_bytes([c[0], c[1], c[2], c[3]])) + .collect() +} +fn to_bytes(v: &[f32]) -> Vec { + v.iter().flat_map(|f| f.to_le_bytes()).collect() +} +fn approx(a: &[f32], b: &[f32]) -> bool { + a.len() == b.len() && a.iter().zip(b).all(|(x, y)| (x - y).abs() < 1e-5) +} + +const DISTS: [f32; 4] = [0.0, 0.5, 1.0, 1.5]; + +fn cpu_falloff(d: f32, mode: u32) -> f32 { + let n = 1.0 - d / RADIUS; + match mode { + 1 => n, + 2 => n * n * (3.0 - 2.0 * n), + 3 => n * n, + 4 => n * n * n, + 5 => RADIUS / (d + RADIUS), + _ => 1.0, + } +} + +fn run_mode(field: Entity, weight: Entity, mode: u32) -> error::Result { + compute_set(field, "falloff_mode", shader_value::ShaderValue::UInt(mode))?; + compute_dispatch(field, 1, 1, 1)?; + let got = f32s(&buffer_read(weight)?); + let want: Vec = DISTS.iter().map(|&d| cpu_falloff(d, mode)).collect(); + if approx(&got, &want) { + println!(" PASS falloff mode {mode}: {got:?}"); + Ok(true) + } else { + println!(" FAIL falloff mode {mode}: got {got:?}, want {want:?}"); + Ok(false) + } +} + +fn sketch() -> error::Result { + init(Config::default())?; + let surface = surface_create_offscreen(1, 1, 1.0, TextureFormat::Rgba8Unorm)?; + let _graphics = graphics_create(surface, 1, 1, TextureFormat::Rgba8Unorm)?; + + let positions: Vec = DISTS.iter().flat_map(|&d| [d, 0.0, 0.0]).collect(); + let pos = buffer_create_with_data(to_bytes(&positions))?; + let weight = buffer_create_with_data(to_bytes(&vec![0.0; DISTS.len()]))?; + + let field = particles_kernel_field()?; + compute_set(field, "position", shader_value::ShaderValue::Buffer(pos))?; + compute_set(field, "weight", shader_value::ShaderValue::Buffer(weight))?; + compute_set( + field, + "center", + shader_value::ShaderValue::Float3([0.0, 0.0, 0.0]), + )?; + compute_set(field, "radius", shader_value::ShaderValue::Float(RADIUS))?; + + let mut ok = true; + for mode in [0u32, 1, 2, 3, 4, 5] { + ok &= run_mode(field, weight, mode)?; + } + + // the other importers must compile too + for (name, r) in [ + ("attract", particles_kernel_attract()), + ("vortex", particles_kernel_vortex()), + ("impulse", particles_kernel_impulse()), + ] { + match r { + Ok(_) => println!(" PASS {name} compiles (falloff import)"), + Err(e) => { + ok = false; + println!(" FAIL {name} compile: {e}"); + } + } + } + Ok(ok) +} + +fn main() { + let ok = sketch().unwrap(); + if ok { + println!("field_falloff_test: ALL PASS"); + exit(0).unwrap(); + } else { + println!("field_falloff_test: FAILURES"); + exit(1).unwrap(); + } +} diff --git a/examples/glue_test.rs b/examples/glue_test.rs new file mode 100644 index 00000000..6a9843d6 --- /dev/null +++ b/examples/glue_test.rs @@ -0,0 +1,110 @@ +//! Glue algebra verbs: `reduce_components`, `extract`, `pack`, `generate`. + +use processing::prelude::*; + +fn f32s(bytes: &[u8]) -> Vec { + bytes + .chunks_exact(4) + .map(|c| f32::from_le_bytes([c[0], c[1], c[2], c[3]])) + .collect() +} +fn to_bytes(v: &[f32]) -> Vec { + v.iter().flat_map(|f| f.to_le_bytes()).collect() +} +fn approx(a: &[f32], b: &[f32]) -> bool { + a.len() == b.len() && a.iter().zip(b).all(|(x, y)| (x - y).abs() < 1e-4) +} + +fn sketch() -> error::Result { + init(Config::default())?; + let surface = surface_create_offscreen(1, 1, 1.0, TextureFormat::Rgba8Unorm)?; + let _graphics = graphics_create(surface, 1, 1, TextureFormat::Rgba8Unorm)?; + + let mut ok = true; + + let vel = vec![ + 3.0f32, 4.0, 0.0, 1.0, 2.0, 2.0, 0.0, 0.0, 7.0, 5.0, 0.0, 0.0, + ]; + + let v = buffer_create_with_data(to_bytes(&vel))?; + let speed = buffer_create_with_data(to_bytes(&vec![0.0; 4]))?; + reduce_components(speed, v, 3, REDUCE_LENGTH)?; + let got = f32s(&buffer_read(speed)?); + if approx(&got, &[5.0, 3.0, 7.0, 5.0]) { + println!(" PASS reduce_components LENGTH (speed=|velocity|): {got:?}"); + } else { + ok = false; + println!(" FAIL reduce LENGTH: got {got:?}, want [5,3,7,5]"); + } + + let y = buffer_create_with_data(to_bytes(&vec![0.0; 4]))?; + extract(y, v, 3, 1)?; + let got_y = f32s(&buffer_read(y)?); + if approx(&got_y, &[4.0, 2.0, 0.0, 0.0]) { + println!(" PASS extract .y: {got_y:?}"); + } else { + ok = false; + println!(" FAIL extract: got {got_y:?}, want [4,2,0,0]"); + } + buffer_destroy(v)?; + buffer_destroy(speed)?; + buffer_destroy(y)?; + + let xs = buffer_create_with_data(to_bytes(&[1.0, 2.0, 3.0]))?; + let ys = buffer_create_with_data(to_bytes(&[10.0, 20.0, 30.0]))?; + let zs = buffer_create_with_data(to_bytes(&[100.0, 200.0, 300.0]))?; + let packed = buffer_create_with_data(to_bytes(&vec![0.0; 9]))?; + pack(packed, &[xs, ys, zs])?; + let got_p = f32s(&buffer_read(packed)?); + if approx( + &got_p, + &[1.0, 10.0, 100.0, 2.0, 20.0, 200.0, 3.0, 30.0, 300.0], + ) { + println!(" PASS pack x/y/z -> Float3: {got_p:?}"); + } else { + ok = false; + println!(" FAIL pack: got {got_p:?}"); + } + buffer_destroy(xs)?; + buffer_destroy(ys)?; + buffer_destroy(zs)?; + buffer_destroy(packed)?; + + let g1 = buffer_create_with_data(to_bytes(&vec![0.0; 16]))?; + let g2 = buffer_create_with_data(to_bytes(&vec![0.0; 16]))?; + let g3 = buffer_create_with_data(to_bytes(&vec![0.0; 16]))?; + generate(g1, 2, GEN_UNIFORM, 42, 1.0, 0.0)?; + generate(g2, 2, GEN_UNIFORM, 42, 1.0, 0.0)?; // same seed + generate(g3, 2, GEN_UNIFORM, 43, 1.0, 0.0)?; // different seed + let a = f32s(&buffer_read(g1)?); + let b = f32s(&buffer_read(g2)?); + let c = f32s(&buffer_read(g3)?); + let in_range = a.iter().all(|x| (0.0..1.0).contains(x)); + let deterministic = approx(&a, &b); + let seed_varies = !approx(&a, &c); + let varied = a.windows(2).any(|w| (w[0] - w[1]).abs() > 1e-6); + if in_range && deterministic && seed_varies && varied { + println!(" PASS generate uniform (in-range, deterministic, seed-sensitive)"); + } else { + ok = false; + println!( + " FAIL generate: in_range={in_range}, deterministic={deterministic}, seed_varies={seed_varies}, varied={varied}" + ); + } + buffer_destroy(g1)?; + buffer_destroy(g2)?; + buffer_destroy(g3)?; + + Ok(ok) +} + +fn main() { + let ok = sketch().unwrap(); + if ok { + println!("glue_test: ALL PASS"); + exit(0).unwrap(); + } else { + println!("glue_test: FAILURES"); + exit(1).unwrap(); + } +} diff --git a/examples/grid_test.rs b/examples/grid_test.rs new file mode 100644 index 00000000..7f5e1573 --- /dev/null +++ b/examples/grid_test.rs @@ -0,0 +1,133 @@ +//! Spatial hash grid (`grid_create` / `grid_build`) against a CPU reference. + +use std::collections::BTreeSet; + +use processing::prelude::*; + +const DIMS: [u32; 3] = [4, 4, 4]; +const CELL_SIZE: f32 = 1.0; +const MIN: [f32; 3] = [0.0, 0.0, 0.0]; + +fn bytes_to_u32s(b: &[u8]) -> Vec { + b.chunks_exact(4) + .map(|c| u32::from_le_bytes([c[0], c[1], c[2], c[3]])) + .collect() +} + +fn cpu_cell(p: [f32; 3]) -> u32 { + let rel = [ + (p[0] - MIN[0]) / CELL_SIZE, + (p[1] - MIN[1]) / CELL_SIZE, + (p[2] - MIN[2]) / CELL_SIZE, + ]; + let clampi = |v: f32, m: u32| (v.floor() as i32).clamp(0, m as i32 - 1) as u32; + let ix = clampi(rel[0], DIMS[0]); + let iy = clampi(rel[1], DIMS[1]); + let iz = clampi(rel[2], DIMS[2]); + ix + iy * DIMS[0] + iz * DIMS[0] * DIMS[1] +} + +fn sketch() -> error::Result { + init(Config::default())?; + let surface = surface_create_offscreen(1, 1, 1.0, TextureFormat::Rgba8Unorm)?; + let _graphics = graphics_create(surface, 1, 1, TextureFormat::Rgba8Unorm)?; + + let capacity: u32 = 40; + let num_cells = (DIMS[0] * DIMS[1] * DIMS[2]) as usize; + + let cells: Vec<[u32; 3]> = (0..capacity) + .map(|i| [(i * 7) % 4, (i * 13) % 4, (i * 5) % 4]) + .collect(); + let positions: Vec<[f32; 3]> = cells + .iter() + .map(|c| [c[0] as f32 + 0.5, c[1] as f32 + 0.5, c[2] as f32 + 0.5]) + .collect(); + + let pos_bytes: Vec = positions + .iter() + .flat_map(|p| p.iter().flat_map(|f| f.to_le_bytes())) + .collect(); + let position = buffer_create_with_data(pos_bytes)?; + + let params = GridParams { + min: MIN, + cell_size: CELL_SIZE, + dims: DIMS, + }; + let grid = grid_create(params, capacity)?; + grid_build(grid, position)?; + + let offsets = bytes_to_u32s(&buffer_read(grid_get(grid)?.offsets)?); + let sorted = bytes_to_u32s(&buffer_read(grid_get(grid)?.sorted)?); + + let mut counts = vec![0u32; num_cells]; + for i in 0..capacity as usize { + counts[cpu_cell(positions[i]) as usize] += 1; + } + let mut expected_offsets = vec![0u32; num_cells + 1]; + for c in 0..num_cells { + expected_offsets[c + 1] = expected_offsets[c] + counts[c]; + } + + let mut ok = true; + + if offsets != expected_offsets { + ok = false; + let i = (0..offsets.len()) + .find(|&i| offsets.get(i) != expected_offsets.get(i)) + .unwrap_or(0); + println!( + " FAIL offsets: first mismatch at cell {i}: got {:?}, want {:?}", + offsets.get(i), + expected_offsets.get(i) + ); + } else { + println!(" PASS offsets (total={})", expected_offsets[num_cells]); + } + + let mut buckets_ok = true; + for c in 0..num_cells { + let (s, e) = ( + expected_offsets[c] as usize, + expected_offsets[c + 1] as usize, + ); + let got: BTreeSet = sorted[s..e].iter().copied().collect(); + let want: BTreeSet = (0..capacity) + .filter(|&i| cpu_cell(positions[i as usize]) as usize == c) + .collect(); + if got != want { + buckets_ok = false; + println!(" FAIL cell {c}: got {got:?}, want {want:?}"); + } + } + if buckets_ok { + println!(" PASS buckets ({num_cells} cells)"); + } else { + ok = false; + } + + grid_destroy(grid)?; + if matches!( + grid_build(grid, position), + Err(error::ProcessingError::GridNotFound) + ) { + println!(" PASS destroyed grid is gone"); + } else { + ok = false; + println!(" FAIL destroyed grid still builds"); + } + + buffer_destroy(position)?; + Ok(ok) +} + +fn main() { + let ok = sketch().unwrap(); + if ok { + println!("grid_test: ALL PASS"); + exit(0).unwrap(); + } else { + println!("grid_test: FAILURES"); + exit(1).unwrap(); + } +} diff --git a/examples/group_compact_test.rs b/examples/group_compact_test.rs new file mode 100644 index 00000000..31f33070 --- /dev/null +++ b/examples/group_compact_test.rs @@ -0,0 +1,79 @@ +//! Stream compaction (`compact`) and the group predicate chain. + +use bevy::prelude::Entity; +use processing::prelude::*; + +fn u32s(b: &[u8]) -> Vec { + b.chunks_exact(4) + .map(|c| u32::from_le_bytes([c[0], c[1], c[2], c[3]])) + .collect() +} +fn f32_bytes(v: &[f32]) -> Vec { + v.iter().flat_map(|f| f.to_le_bytes()).collect() +} + +fn check(label: &str, got_count: u32, got: &[u32], want: &[u32]) -> bool { + if got_count as usize == want.len() && &got[..want.len()] == want { + println!( + " PASS {label}: count={got_count}, indices={:?}", + &got[..want.len()] + ); + true + } else { + println!(" FAIL {label}: count={got_count} indices={got:?}, want {want:?}"); + false + } +} + +fn compact_flags(flags: &[f32]) -> error::Result<(u32, Vec)> { + let f = buffer_create_with_data(f32_bytes(flags))?; + let idx = buffer_create((flags.len().max(1) as u64) * 4)?; + let count = compact(f, idx)?; + let out = u32s(&buffer_read(idx)?); + buffer_destroy(f)?; + buffer_destroy(idx)?; + Ok((count, out)) +} + +fn sketch() -> error::Result { + init(Config::default())?; + let surface = surface_create_offscreen(1, 1, 1.0, TextureFormat::Rgba8Unorm)?; + let _graphics = graphics_create(surface, 1, 1, TextureFormat::Rgba8Unorm)?; + + let mut ok = true; + + let (c, out) = compact_flags(&[0.0, 1.0, 1.0, 0.0, 1.0, 0.0, 0.0, 1.0])?; + ok &= check("compact mixed", c, &out, &[1, 2, 4, 7]); + + let (c, out) = compact_flags(&[0.0, 0.0, 0.0])?; + ok &= check("compact none", c, &out, &[]); + + let (c, out) = compact_flags(&[1.0, 1.0, 1.0, 1.0])?; + ok &= check("compact all", c, &out, &[0, 1, 2, 3]); + + let src_vals = [0.1f32, 0.5, 0.9, 0.3, 0.7]; + let src = buffer_create_with_data(f32_bytes(&src_vals))?; + let flags = buffer_create((src_vals.len() as u64) * 4)?; + let idx = buffer_create((src_vals.len() as u64) * 4)?; + map(flags, src, 1, MAP_GREATER, 0.4, 0.0)?; + let count = compact(flags, idx)?; + let out = u32s(&buffer_read(idx)?); + ok &= check("group(>0.4)+compact", count, &out, &[1, 2, 4]); + for b in [src, flags, idx] { + let _: Entity = b; + buffer_destroy(b)?; + } + + Ok(ok) +} + +fn main() { + let ok = sketch().unwrap(); + if ok { + println!("group_compact_test: ALL PASS"); + exit(0).unwrap(); + } else { + println!("group_compact_test: FAILURES"); + exit(1).unwrap(); + } +} diff --git a/examples/lookup_test.rs b/examples/lookup_test.rs new file mode 100644 index 00000000..e30771c8 --- /dev/null +++ b/examples/lookup_test.rs @@ -0,0 +1,100 @@ +//! `lookup`: per-particle texture sampling, 1-D and 2-D. + +use bevy::prelude::Entity; +use bevy::render::render_resource::Extent3d; +use processing::prelude::*; + +const RED: [u8; 4] = [255, 0, 0, 255]; +const GREEN: [u8; 4] = [0, 255, 0, 255]; +const BLUE: [u8; 4] = [0, 0, 255, 255]; +const WHITE: [u8; 4] = [255, 255, 255, 255]; + +fn f32s(bytes: &[u8]) -> Vec { + bytes + .chunks_exact(4) + .map(|c| f32::from_le_bytes([c[0], c[1], c[2], c[3]])) + .collect() +} +fn to_bytes(v: &[f32]) -> Vec { + v.iter().flat_map(|f| f.to_le_bytes()).collect() +} +fn approx(a: &[f32], b: &[f32]) -> bool { + a.len() == b.len() && a.iter().zip(b).all(|(x, y)| (x - y).abs() < 1e-4) +} + +fn ramp_image(width: u32, height: u32, texels: &[[u8; 4]]) -> error::Result { + let data: Vec = texels.iter().flatten().copied().collect(); + let img = image_create( + Extent3d { + width, + height, + depth_or_array_layers: 1, + }, + data, + TextureFormat::Rgba8Unorm, + )?; + image_set_sampler(img, 1, 0, 0)?; // nearest, clamp + Ok(img) +} + +fn sketch() -> error::Result { + init(Config::default())?; + let surface = surface_create_offscreen(1, 1, 1.0, TextureFormat::Rgba8Unorm)?; + let _graphics = graphics_create(surface, 1, 1, TextureFormat::Rgba8Unorm)?; + + let mut ok = true; + + let ramp = ramp_image(4, 1, &[RED, GREEN, BLUE, WHITE])?; + let t = vec![0.125f32, 0.375, 0.625, 0.875]; + let op_in = buffer_create_with_data(to_bytes(&t))?; + let dst = buffer_create_with_data(to_bytes(&vec![0.0; 16]))?; + lookup(dst, op_in, ramp, 1, 1.0, 0.0, 1.0, 0.0, 1.0)?; + let got = f32s(&buffer_read(dst)?); + let want = vec![ + 1.0, 0.0, 0.0, 1.0, // red + 0.0, 1.0, 0.0, 1.0, // green + 0.0, 0.0, 1.0, 1.0, // blue + 1.0, 1.0, 1.0, 1.0, // white + ]; + if approx(&got, &want) { + println!(" PASS lookup 1-D ramp"); + } else { + ok = false; + println!(" FAIL lookup 1-D: got {got:?}"); + } + buffer_destroy(op_in)?; + buffer_destroy(dst)?; + + let tex2 = ramp_image(2, 2, &[RED, GREEN, BLUE, WHITE])?; + let uv = vec![ + 0.25f32, 0.25, // red + 0.75, 0.25, // green + 0.25, 0.75, // blue + 0.75, 0.75, // white + ]; + let op_in2 = buffer_create_with_data(to_bytes(&uv))?; + let dst2 = buffer_create_with_data(to_bytes(&vec![0.0; 16]))?; + lookup(dst2, op_in2, tex2, 2, 1.0, 0.0, 1.0, 0.0, 1.0)?; + let got2 = f32s(&buffer_read(dst2)?); + if approx(&got2, &want) { + println!(" PASS lookup 2-D texture"); + } else { + ok = false; + println!(" FAIL lookup 2-D: got {got2:?}"); + } + buffer_destroy(op_in2)?; + buffer_destroy(dst2)?; + + Ok(ok) +} + +fn main() { + let ok = sketch().unwrap(); + if ok { + println!("lookup_test: ALL PASS"); + exit(0).unwrap(); + } else { + println!("lookup_test: FAILURES"); + exit(1).unwrap(); + } +} diff --git a/examples/map_test.rs b/examples/map_test.rs new file mode 100644 index 00000000..9a89df13 --- /dev/null +++ b/examples/map_test.rs @@ -0,0 +1,77 @@ +//! `map`: in-place and out-of-place. + +use processing::prelude::*; + +fn f32s(bytes: &[u8]) -> Vec { + bytes + .chunks_exact(4) + .map(|c| f32::from_le_bytes([c[0], c[1], c[2], c[3]])) + .collect() +} +fn to_bytes(v: &[f32]) -> Vec { + v.iter().flat_map(|f| f.to_le_bytes()).collect() +} +fn approx(a: &[f32], b: &[f32]) -> bool { + a.len() == b.len() && a.iter().zip(b).all(|(x, y)| (x - y).abs() < 1e-5) +} + +fn sketch() -> error::Result { + init(Config::default())?; + let surface = surface_create_offscreen(1, 1, 1.0, TextureFormat::Rgba8Unorm)?; + let _graphics = graphics_create(surface, 1, 1, TextureFormat::Rgba8Unorm)?; + + let n_particles = 5u32; + let components = 3u32; + let input: Vec = (0..(n_particles * components)).map(|i| i as f32).collect(); + + let mut ok = true; + + let a = buffer_create_with_data(to_bytes(&input))?; + map(a, a, components, MAP_AFFINE, 2.0, 1.0)?; + let got = f32s(&buffer_read(a)?); + let want: Vec = input.iter().map(|x| x * 2.0 + 1.0).collect(); + if approx(&got, &want) { + println!( + " PASS map in-place, all {} components: {:?}", + got.len(), + got + ); + } else { + ok = false; + println!(" FAIL map in-place: got {got:?}, want {want:?}"); + } + buffer_destroy(a)?; + + let neg: Vec = (0..(n_particles * components)) + .map(|i| -(i as f32) - 0.5) + .collect(); + let a2 = buffer_create_with_data(to_bytes(&neg))?; + let dst = buffer_create_with_data(to_bytes(&vec![0.0; neg.len()]))?; + map(dst, a2, components, MAP_ABS, 0.0, 0.0)?; + let got_dst = f32s(&buffer_read(dst)?); + let got_a = f32s(&buffer_read(a2)?); + let want_dst: Vec = neg.iter().map(|x| x.abs()).collect(); + if approx(&got_dst, &want_dst) && approx(&got_a, &neg) { + println!(" PASS map out-of-place: dst={got_dst:?}, a preserved"); + } else { + ok = false; + println!( + " FAIL map out-of-place: dst={got_dst:?} (want {want_dst:?}), a={got_a:?} (want {neg:?})" + ); + } + buffer_destroy(a2)?; + buffer_destroy(dst)?; + + Ok(ok) +} + +fn main() { + let ok = sketch().unwrap(); + if ok { + println!("map_test: ALL PASS"); + exit(0).unwrap(); + } else { + println!("map_test: FAILURES"); + exit(1).unwrap(); + } +} diff --git a/examples/reduce_test.rs b/examples/reduce_test.rs new file mode 100644 index 00000000..75d3d613 --- /dev/null +++ b/examples/reduce_test.rs @@ -0,0 +1,71 @@ +//! GPU->CPU `reduce` (sum / min / max), single- and multi-level. + +use bevy::prelude::Entity; +use processing::prelude::*; + +fn to_bytes(v: &[f32]) -> Vec { + v.iter().flat_map(|f| f.to_le_bytes()).collect() +} + +fn check(label: &str, got: f32, want: f32) -> bool { + if (got - want).abs() <= 1e-3 * want.abs().max(1.0) { + println!(" PASS {label}: {got}"); + true + } else { + println!(" FAIL {label}: got {got}, want {want}"); + false + } +} + +fn reduce_all(vals: &[f32]) -> error::Result<(f32, f32, f32)> { + let b = buffer_create_with_data(to_bytes(vals))?; + let sum = reduce(b, REDUCE_OP_SUM)?; + let mn = reduce(b, REDUCE_OP_MIN)?; + let mx = reduce(b, REDUCE_OP_MAX)?; + let _: Entity = b; + buffer_destroy(b)?; + Ok((sum, mn, mx)) +} + +fn sketch() -> error::Result { + init(Config::default())?; + let surface = surface_create_offscreen(1, 1, 1.0, TextureFormat::Rgba8Unorm)?; + let _graphics = graphics_create(surface, 1, 1, TextureFormat::Rgba8Unorm)?; + + let mut ok = true; + + let (s, mn, mx) = reduce_all(&vec![1.0; 1000])?; + ok &= check("ones/1000 sum", s, 1000.0); + ok &= check("ones/1000 min", mn, 1.0); + ok &= check("ones/1000 max", mx, 1.0); + + let ramp: Vec = (0..300).map(|i| i as f32).collect(); + let (s, mn, mx) = reduce_all(&ramp)?; + ok &= check("ramp/300 sum", s, 300.0 * 299.0 / 2.0); + ok &= check("ramp/300 min", mn, 0.0); + ok &= check("ramp/300 max", mx, 299.0); + + let (s, mn, mx) = reduce_all(&vec![1.0; 70_000])?; + ok &= check("ones/70k sum", s, 70_000.0); + ok &= check("ones/70k min", mn, 1.0); + ok &= check("ones/70k max", mx, 1.0); + + let big: Vec = (0..70_000).map(|i| i as f32).collect(); + let b = buffer_create_with_data(to_bytes(&big))?; + ok &= check("ramp/70k min", reduce(b, REDUCE_OP_MIN)?, 0.0); + ok &= check("ramp/70k max", reduce(b, REDUCE_OP_MAX)?, 69_999.0); + buffer_destroy(b)?; + + Ok(ok) +} + +fn main() { + let ok = sketch().unwrap(); + if ok { + println!("reduce_test: ALL PASS"); + exit(0).unwrap(); + } else { + println!("reduce_test: FAILURES"); + exit(1).unwrap(); + } +} diff --git a/examples/scan_test.rs b/examples/scan_test.rs new file mode 100644 index 00000000..fa0c635b --- /dev/null +++ b/examples/scan_test.rs @@ -0,0 +1,68 @@ +//! GPU exclusive prefix-sum (`prefix_sum_u32`) against a CPU scan. + +use processing::prelude::*; + +fn u32s_to_bytes(v: &[u32]) -> Vec { + v.iter().flat_map(|x| x.to_le_bytes()).collect() +} + +fn bytes_to_u32s(b: &[u8]) -> Vec { + b.chunks_exact(4) + .map(|c| u32::from_le_bytes([c[0], c[1], c[2], c[3]])) + .collect() +} + +fn run_case(n: usize) -> error::Result { + let input: Vec = (0..n).map(|i| (i % 7 + 1) as u32).collect(); + + let buf = buffer_create_with_data(u32s_to_bytes(&input))?; + prefix_sum_u32(buf)?; + let out = bytes_to_u32s(&buffer_read(buf)?); + buffer_destroy(buf)?; + + let mut expected = vec![0u32; n]; + let mut acc = 0u32; + for i in 0..n { + expected[i] = acc; + acc += input[i]; + } + + if out.len() != expected.len() { + println!(" FAIL n={n}: length {} != {}", out.len(), expected.len()); + return Ok(false); + } + if let Some(i) = (0..n).find(|&i| out[i] != expected[i]) { + println!( + " FAIL n={n}: first mismatch at index {i}: got {}, want {}", + out[i], expected[i] + ); + return Ok(false); + } + println!(" PASS n={n} (total={acc})"); + Ok(true) +} + +fn sketch() -> error::Result { + init(Config::default())?; + let surface = surface_create_offscreen(1, 1, 1.0, TextureFormat::Rgba8Unorm)?; + let _graphics = graphics_create(surface, 1, 1, TextureFormat::Rgba8Unorm)?; + + // empty, block boundary, two levels + let cases = [1usize, 5, 255, 256, 257, 1000, 65_536, 70_000, 300_000]; + let mut all_ok = true; + for &n in &cases { + all_ok &= run_case(n)?; + } + Ok(all_ok) +} + +fn main() { + let ok = sketch().unwrap(); + if ok { + println!("scan_test: ALL PASS"); + exit(0).unwrap(); + } else { + println!("scan_test: FAILURES"); + exit(1).unwrap(); + } +} diff --git a/examples/sort_test.rs b/examples/sort_test.rs new file mode 100644 index 00000000..159296fd --- /dev/null +++ b/examples/sort_test.rs @@ -0,0 +1,91 @@ +//! GPU bitonic sort (`bitonic_sort_by_key`): key order and payload permutation. + +use processing::prelude::*; + +fn u32s(b: &[u8]) -> Vec { + b.chunks_exact(4) + .map(|c| u32::from_le_bytes([c[0], c[1], c[2], c[3]])) + .collect() +} +fn f32s(b: &[u8]) -> Vec { + b.chunks_exact(4) + .map(|c| f32::from_le_bytes([c[0], c[1], c[2], c[3]])) + .collect() +} +fn f32_bytes(v: &[f32]) -> Vec { + v.iter().flat_map(|f| f.to_le_bytes()).collect() +} +fn u32_bytes(v: &[u32]) -> Vec { + v.iter().flat_map(|x| x.to_le_bytes()).collect() +} + +fn rnd(seed: u32) -> f32 { + let mut x = seed.wrapping_mul(747796405).wrapping_add(2891336453); + x = ((x >> ((x >> 28).wrapping_add(4))) ^ x).wrapping_mul(277803737); + x = (x >> 22) ^ x; + (x as f32) / (u32::MAX as f32) +} + +fn run_case(n: u32) -> error::Result { + let orig_keys: Vec = (0..n) + .map(|i| rnd(i.wrapping_mul(2654435761).wrapping_add(7))) + .collect(); + let idx: Vec = (0..n).collect(); + + let keys = buffer_create_with_data(f32_bytes(&orig_keys))?; + let payload = buffer_create_with_data(u32_bytes(&idx))?; + bitonic_sort_by_key(keys, payload)?; + let sorted_keys = f32s(&buffer_read(keys)?); + let perm = u32s(&buffer_read(payload)?); + buffer_destroy(keys)?; + buffer_destroy(payload)?; + + if let Some(i) = (1..n as usize).find(|&i| sorted_keys[i] < sorted_keys[i - 1]) { + println!( + " FAIL n={n}: not ascending at {i}: {} < {}", + sorted_keys[i], + sorted_keys[i - 1] + ); + return Ok(false); + } + let mut seen = vec![false; n as usize]; + for &p in &perm { + if p >= n || seen[p as usize] { + println!(" FAIL n={n}: payload not a permutation (bad/dup index {p})"); + return Ok(false); + } + seen[p as usize] = true; + } + if let Some(i) = (0..n as usize).find(|&i| orig_keys[perm[i] as usize] != sorted_keys[i]) { + println!( + " FAIL n={n}: payload[{i}] mismatch: orig[{}]={} != {}", + perm[i], orig_keys[perm[i] as usize], sorted_keys[i] + ); + return Ok(false); + } + println!(" PASS n={n}"); + Ok(true) +} + +fn sketch() -> error::Result { + init(Config::default())?; + let surface = surface_create_offscreen(1, 1, 1.0, TextureFormat::Rgba8Unorm)?; + let _graphics = graphics_create(surface, 1, 1, TextureFormat::Rgba8Unorm)?; + + let mut ok = true; + for &n in &[2u32, 8, 64, 1024, 4096, 65536] { + ok &= run_case(n)?; + } + Ok(ok) +} + +fn main() { + let ok = sketch().unwrap(); + if ok { + println!("sort_test: ALL PASS"); + exit(0).unwrap(); + } else { + println!("sort_test: FAILURES"); + exit(1).unwrap(); + } +}