GPU test harnesses.

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"

examples/combine_mix_test.rs

@@ -0,0 +1,109 @@
+//! `combine` and `mix`, including the alias guard.
+
+use processing::prelude::*;
+
+fn f32s(bytes: &[u8]) -> Vec<f32> {
+    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<u8> {
+    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<bool> {
+    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<f32> = (0..12).map(|i| i as f32).collect();
+    let bv: Vec<f32> = (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<f32> = 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<f32> = 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<f32> = vec![1.0; 12];
+    let b2: Vec<f32> = 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();
+    }
+}

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<u32> {
+    bytes
+        .chunks_exact(4)
+        .map(|c| u32::from_le_bytes([c[0], c[1], c[2], c[3]]))
+        .collect()
+}
+
+fn f32s(bytes: &[u8]) -> Vec<f32> {
+    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<u8> {
+    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<storage, read_write> data: array<f32>;
+@compute @workgroup_size(4)
+fn main(@builtin(global_invocation_id) id: vec3<u32>) {
+    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<f32> = (0..n)
+            .map(|i| ((i * (round + 1)) % 4) as f32 + 0.5)
+            .collect();
+        let packed: Vec<f32> = 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<u32> = (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<u32> = (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<u32> = values
+        .iter()
+        .map(|v| {
+            let s = acc;
+            acc += v;
+            s
+        })
+        .collect();
+    let mut acc = 0u32;
+    let inclusive: Vec<u32> = 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(())
+}

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<f32> {
+    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<u8> {
+    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<bool> {
+    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<f32> = 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<bool> {
+    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<f32> = 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();
+    }
+}