graphic design tool
git clone https://git.lucas.co/cce-designer.git
feat: the Collision test runs on the GPU — the second Phase 7 step 4 operator
src/collide.rs is the node's brute-force test — every query against every
collider triangle, the Voronoi-region distance for Proximity and a
Möller–Trumbore parity cast for Inside — as one algorithm on both
backends, held together by collision_gpu_matches_cpu (zero disagreements
over 6k queries x 1.7k triangles in both modes). The resolver runs the
test as one batch over every element the type asks about, where it used
to hand select_elements a closure that asked one point at a time; the
batch is what goes to the GPU whole, one dispatch, no passes to chain.
Measured in release on an Intel Iris Xe: 3.6M pairs cpu 20 ms / gpu
2.2 ms; 15M pairs cpu 76 ms / gpu 6.6 ms; 242M pairs cpu 1150 ms / gpu
52 ms. GPU_MIN_WORK = 250k pairs is the auto threshold.
Neighbour's Diffuse and Relax's Repel stay on the CPU deliberately, and
CLAUDE.md records why: a single gather per evaluation cannot amortise a
submission's round trip, and Repel's per-pass spatial grid does not chain.
Co-Authored-By: Claude Fable 5.1 <[email protected]>
CLAUDE.md | 27 +++++++
shapeshifter.md | 10 ++-
src/collide.rs | 223 ++++++++++++++++++++++++++++++++++++++++++++++++++++++++
src/geometry.rs | 87 ++++++++++++++++------
src/main.rs | 110 ++++++++++++++++++++++++++++
5 files changed, 434 insertions(+), 23 deletions(-)
diff --git a/CLAUDE.md b/CLAUDE.md
index 4b3ed06..1132571 100644
--- a/CLAUDE.md
+++ b/CLAUDE.md
@@ -823,6 +823,33 @@ GPU is a win for large meshes and a loss for the ones most projects have,
which is the honest state of step 4 and the reason auto does not simply
mean GPU.
+**Collision is the second operator (`src/collide.rs`), and the one the
+GPU is made for.** The node's test has always been a brute-force loop —
+every query against every collider triangle, the Voronoi-region distance
+for Proximity and a Möller–Trumbore parity cast for Inside — so the work
+is queries x triangles, per-point, one dispatch, no passes to chain. The
+resolver now runs the test as ONE batch over every element the type asks
+about (points, or primitive centroids), where it used to hand
+`select_elements` a closure that asked one point at a time; that batch is
+what can go to the GPU whole. Same algorithm step for step on both sides,
+held by `collision_gpu_matches_cpu` (zero disagreements over 6k queries x
+1.7k triangles in both modes; a knife-edge query at the threshold may
+round either way and is tolerated only there). `collision_timing` in
+release, Proximity: 3.6M pairs cpu 20 ms / gpu 2.2 ms; 15M pairs cpu 76
+ms / gpu 6.6 ms; 242M pairs cpu 1150 ms / gpu 52 ms. `GPU_MIN_WORK`
+(250k pairs) is the auto threshold.
+
+**Two per-point operators are deliberately NOT on the GPU, and the
+measurements above say why.** Neighbour's Diffuse and Concentrate are a
+single gather per evaluation — one pass, then the rest of the graph runs
+on the CPU before the next frame's pass — so there is nothing to chain
+into one submission, and a single pass costs ~0.5 ms of round trip against
+a CPU gather that takes less than that on any mesh a project has. Relax's
+Repel rebuilds a spatial grid every pass, which is the part that does not
+fit a chained submission; a GPU-side grid is a project of its own, and a
+brute-force O(n^2) pass that the CPU twin would then have to match is a
+regression for every CPU user. Both stay native until a workload asks.
+
### The volume representation
`src/volume.rs` is a dense signed distance field — `Volume { origin, voxel,
diff --git a/shapeshifter.md b/shapeshifter.md
index 7b577db..1e4c696 100644
--- a/shapeshifter.md
+++ b/shapeshifter.md
@@ -711,8 +711,14 @@ Nothing else in the workspace loads OpenCL, so the ICD bug leaves with it.
> submission, ping-ponging on the device — because a pass submitted on its
> own lost to the CPU at every size; with it the GPU wins from the tens of
> thousands of points up (135k: 14 ms against 25 ms) and the auto
-> threshold sits at 32k. Next: the rest of the per-point set (Repel,
-> Diffuse, the collision response), which share the pattern.
+> threshold sits at 32k. Collision followed (`src/collide.rs`): the
+> node's brute-force queries x triangles test as one dispatch, zero
+> disagreements against the CPU, 9x faster at 3.6M pairs and 22x at 242M.
+> Diffuse and Repel stay on the CPU on purpose — a single gather per
+> evaluation cannot amortise a submission, and Repel's per-pass spatial
+> grid does not chain — so step 4's shape is settled: the GPU takes the
+> operators whose work is large per submission, and the measurements
+> in CLAUDE.md say which those are.
**Step 4 — GPU compute, through the renderer.** Scripts do not run on the
GPU: no embedded language compiles to GPU code, and none should. Parallel work
diff --git a/src/collide.rs b/src/collide.rs
new file mode 100644
index 0000000..c2c50f9
--- /dev/null
+++ b/src/collide.rs
@@ -0,0 +1,223 @@
+//! The Collision node's test — is this point inside the collider, or within
+//! Distance of its surface — on the CPU and on the GPU. The second Phase 7
+//! step 4 operator, and the one the GPU is made for: every query walks EVERY
+//! collider triangle (the CPU test has always been that brute-force loop),
+//! so the work is queries x triangles, per-point, and embarrassingly
+//! parallel, with one dispatch and no passes to chain.
+//!
+//! Same algorithm on both sides — the Voronoi-region point-triangle distance
+//! and the Möller–Trumbore ray test, step for step — held together by
+//! `collision_gpu_matches_cpu`. The layout is the GPU's: queries as a flat
+//! `xyz` array, triangles as nine floats each, one `u32` flag out.
+
+use cce_ui::vk::{Binding, ComputeDevice, Kernel};
+use glam::Vec3;
+
+/// Which test the node runs.
+#[derive(Clone, Copy, Debug, PartialEq)]
+pub enum Test {
+ /// Enclosed by the collider's volume: parity of a ray cast's crossings.
+ Inside,
+ /// Within this distance of the collider's surface.
+ Proximity(f32),
+}
+
+/// Fixed irrational-ish ray, NOT axis-aligned: the template meshes
+/// tessellate on the axes, and a ray along one skims edge-on through whole
+/// fans of triangles, double-counting crossings.
+pub fn ray_dir() -> Vec3 {
+ Vec3::new(0.9174771, 0.3369154, 0.2095338).normalize()
+}
+
+/// The CPU test: one flag per query, walking every triangle.
+pub fn hits_cpu(queries: &[Vec3], tris: &[[Vec3; 3]], test: Test) -> Vec<u32> {
+ let dir = ray_dir();
+ queries
+ .iter()
+ .map(|&p| match test {
+ Test::Proximity(d) => {
+ let d2 = d * d;
+ u32::from(tris.iter().any(|t| crate::geometry::point_triangle_distance_sq(p, t[0], t[1], t[2]) <= d2))
+ }
+ Test::Inside => {
+ let crossings = tris.iter().filter(|t| crate::spatial::ray_triangle(p, dir, t[0], t[1], t[2]).is_some()).count();
+ (crossings % 2 == 1) as u32
+ }
+ })
+ .collect()
+}
+
+/// The same test as a WGSL kernel: one invocation per query.
+pub const COLLIDE_WGSL: &str = r#"
+struct Params { n: u32, m: u32, mode: u32, pad: u32, d2: f32, rx: f32, ry: f32, rz: f32 }
+@group(0) @binding(0) var<storage, read> queries: array<f32>;
+@group(0) @binding(1) var<storage, read> tris: array<f32>;
+@group(0) @binding(2) var<storage, read_write> hits: array<u32>;
+@group(0) @binding(3) var<uniform> params: Params;
+
+// Squared distance from p to triangle (a, b, c): the Voronoi-region walk.
+fn dist_sq(p: vec3<f32>, a: vec3<f32>, b: vec3<f32>, c: vec3<f32>) -> f32 {
+ let ab = b - a;
+ let ac = c - a;
+ let ap = p - a;
+ let d1 = dot(ab, ap);
+ let d2 = dot(ac, ap);
+ if (d1 <= 0.0 && d2 <= 0.0) { return dot(ap, ap); }
+ let bp = p - b;
+ let d3 = dot(ab, bp);
+ let d4 = dot(ac, bp);
+ if (d3 >= 0.0 && d4 <= d3) { return dot(bp, bp); }
+ let vc = d1 * d4 - d3 * d2;
+ if (vc <= 0.0 && d1 >= 0.0 && d3 <= 0.0) {
+ let v = d1 / (d1 - d3);
+ let e = ap - ab * v;
+ return dot(e, e);
+ }
+ let cp = p - c;
+ let d5 = dot(ab, cp);
+ let d6 = dot(ac, cp);
+ if (d6 >= 0.0 && d5 <= d6) { return dot(cp, cp); }
+ let vb = d5 * d2 - d1 * d6;
+ if (vb <= 0.0 && d2 >= 0.0 && d6 <= 0.0) {
+ let w = d2 / (d2 - d6);
+ let e = ap - ac * w;
+ return dot(e, e);
+ }
+ let va = d3 * d6 - d5 * d4;
+ if (va <= 0.0 && (d4 - d3) >= 0.0 && (d5 - d6) >= 0.0) {
+ let w = (d4 - d3) / ((d4 - d3) + (d5 - d6));
+ let e = bp - (c - b) * w;
+ return dot(e, e);
+ }
+ let denom = 1.0 / (va + vb + vc);
+ let v = vb * denom;
+ let w = vc * denom;
+ let e = ap - ab * v - ac * w;
+ return dot(e, e);
+}
+
+// Möller–Trumbore, as spatial::ray_triangle: a hit strictly ahead of the origin.
+fn ray_hits(o: vec3<f32>, d: vec3<f32>, v0: vec3<f32>, v1: vec3<f32>, v2: vec3<f32>) -> bool {
+ let edge1 = v1 - v0;
+ let edge2 = v2 - v0;
+ let h = cross(d, edge2);
+ let a = dot(edge1, h);
+ if (abs(a) < 1e-6) { return false; }
+ let f = 1.0 / a;
+ let s = o - v0;
+ let u = f * dot(s, h);
+ if (u < 0.0 || u > 1.0) { return false; }
+ let q = cross(s, edge1);
+ let v = f * dot(d, q);
+ if (v < 0.0 || u + v > 1.0) { return false; }
+ let t = f * dot(edge2, q);
+ return t > 1e-5;
+}
+
+@compute @workgroup_size(64)
+fn collide(@builtin(global_invocation_id) id: vec3<u32>) {
+ let i = id.x;
+ if (i >= params.n) { return; }
+ let p = vec3<f32>(queries[i * 3u], queries[i * 3u + 1u], queries[i * 3u + 2u]);
+ var hit = 0u;
+ if (params.mode == 1u) {
+ for (var t = 0u; t < params.m; t = t + 1u) {
+ let b = t * 9u;
+ let a0 = vec3<f32>(tris[b], tris[b + 1u], tris[b + 2u]);
+ let a1 = vec3<f32>(tris[b + 3u], tris[b + 4u], tris[b + 5u]);
+ let a2 = vec3<f32>(tris[b + 6u], tris[b + 7u], tris[b + 8u]);
+ if (dist_sq(p, a0, a1, a2) <= params.d2) { hit = 1u; break; }
+ }
+ } else {
+ let d = vec3<f32>(params.rx, params.ry, params.rz);
+ var crossings = 0u;
+ for (var t = 0u; t < params.m; t = t + 1u) {
+ let b = t * 9u;
+ let a0 = vec3<f32>(tris[b], tris[b + 1u], tris[b + 2u]);
+ let a1 = vec3<f32>(tris[b + 3u], tris[b + 4u], tris[b + 5u]);
+ let a2 = vec3<f32>(tris[b + 6u], tris[b + 7u], tris[b + 8u]);
+ if (ray_hits(p, d, a0, a1, a2)) { crossings = crossings + 1u; }
+ }
+ hit = crossings & 1u;
+ }
+ hits[i] = hit;
+}"#;
+
+#[repr(C)]
+#[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
+struct Params {
+ n: u32,
+ m: u32,
+ mode: u32,
+ pad: u32,
+ d2: f32,
+ rx: f32,
+ ry: f32,
+ rz: f32,
+}
+
+/// The GPU test: one dispatch, one flag per query.
+pub fn hits_gpu(dev: &mut ComputeDevice, queries: &[Vec3], tris: &[[Vec3; 3]], test: Test) -> Result<Vec<u32>, String> {
+ if queries.is_empty() {
+ return Ok(Vec::new());
+ }
+ let q: Vec<f32> = queries.iter().flat_map(|p| [p.x, p.y, p.z]).collect();
+ let t: Vec<f32> = tris.iter().flat_map(|t| [t[0].x, t[0].y, t[0].z, t[1].x, t[1].y, t[1].z, t[2].x, t[2].y, t[2].z]).collect();
+ let dir = ray_dir();
+ let params = Params {
+ n: queries.len() as u32,
+ m: tris.len() as u32,
+ mode: match test {
+ Test::Inside => 0,
+ Test::Proximity(_) => 1,
+ },
+ pad: 0,
+ d2: match test {
+ Test::Proximity(d) => d * d,
+ Test::Inside => 0.0,
+ },
+ rx: dir.x,
+ ry: dir.y,
+ rz: dir.z,
+ };
+ let mut hits = vec![0u32; queries.len()];
+ dev.run_over(
+ &Kernel::new(COLLIDE_WGSL, "collide"),
+ &mut [Binding::input(&q), Binding::input(&t), Binding::rw(&mut hits), Binding::uniform(¶ms)],
+ queries.len() as u32,
+ )?;
+ Ok(hits)
+}
+
+/// Below this many query-triangle pairs the CPU is faster; in auto mode it
+/// takes them. From `collision_timing` in release on an Intel Iris Xe,
+/// Proximity: 360k pairs cpu 2.7 ms / gpu 15 ms (that run pays the ~15 ms
+/// pipeline compile); 3.6M pairs cpu 20 ms / gpu 2.2 ms; 15M pairs cpu
+/// 76 ms / gpu 6.6 ms; 242M pairs cpu 1150 ms / gpu 52 ms. The GPU is
+/// ahead from a few hundred thousand pairs up, and by 20x at the sizes
+/// where the CPU test stalls the app.
+pub const GPU_MIN_WORK: usize = 250_000;
+
+/// The test as the node runs it: the backend chosen by `CCE_COMPUTE` and the
+/// amount of work, a GPU failure in auto mode falling back to the CPU with
+/// a note and in forced-GPU mode reported back for the node-error slot.
+pub fn hits(queries: &[Vec3], tris: &[[Vec3; 3]], test: Test) -> Result<Vec<u32>, String> {
+ let work = queries.len().saturating_mul(tris.len());
+ if crate::gpu::use_gpu(work, GPU_MIN_WORK) {
+ match crate::gpu::with_any_device(|dev| hits_gpu(dev, queries, tris, test)) {
+ Ok(Ok(h)) => return Ok(h),
+ Ok(Err(e)) | Err(e) => {
+ if crate::gpu::choice() == crate::gpu::Choice::Gpu {
+ return Err(format!("CCE_COMPUTE=gpu but the collision test could not run there ({e}); tested on the CPU"));
+ }
+ note_fallback_once(&e);
+ }
+ }
+ }
+ Ok(hits_cpu(queries, tris, test))
+}
+
+fn note_fallback_once(e: &str) {
+ static ONCE: std::sync::Once = std::sync::Once::new();
+ ONCE.call_once(|| eprintln!("cce-designer: collision test fell back to the CPU: {e}"));
+}
diff --git a/src/geometry.rs b/src/geometry.rs
index fb66a4a..fbe1b3d 100644
--- a/src/geometry.rs
+++ b/src/geometry.rs
@@ -1499,7 +1499,7 @@ fn apply_group(
/// Squared distance from `p` to triangle `(a, b, c)` — closest point via the
/// Voronoi-region walk (Ericson, Real-Time Collision Detection §5.1.5).
-fn point_triangle_distance_sq(p: Vec3, a: Vec3, b: Vec3, c: Vec3) -> f32 {
+pub(crate) fn point_triangle_distance_sq(p: Vec3, a: Vec3, b: Vec3, c: Vec3) -> f32 {
let ab = b - a;
let ac = c - a;
let ap = p - a;
@@ -1598,29 +1598,74 @@ pub fn resolve_collision_geometry_with_errors(
let method = node_param_str(target, "Method", "Inside").to_lowercase();
let distance = node_param_f32(target, "Distance", 0.05).max(0.0);
- // Fixed irrational-ish direction, NOT axis-aligned: the template meshes
- // tessellate on the axes, and a ray along one skims edge-on through
- // whole fans of triangles, double-counting crossings.
- let ray_dir = Vec3::new(0.9174771, 0.3369154, 0.2095338).normalize();
- let hit = |pt: Vec3| -> bool {
- if method == "proximity" {
- let d2 = distance * distance;
- tris.iter().any(|t| point_triangle_distance_sq(pt, t[0], t[1], t[2]) <= d2)
- } else {
- let crossings = tris
- .iter()
- .filter(|t| crate::spatial::ray_triangle(pt, ray_dir, t[0], t[1], t[2]).is_some())
- .count();
- crossings % 2 == 1
- }
- };
+ let test = if method == "proximity" { crate::collide::Test::Proximity(distance) } else { crate::collide::Test::Inside };
- // Collision has no Mode parameter, so `select_elements` takes its Box
- // branch and applies `hit` per element — which is the whole difference
- // between this node and Group.
+ // The test runs as ONE batch over every element the type asks about —
+ // points, or primitive centroids — so it can go to the GPU whole
+ // (`crate::collide`, the second Phase 7 step 4 operator); the
+ // per-element closure `select_elements` takes would have asked one
+ // point at a time. Edges test both endpoints, as before.
let etype = node_param_str(target, "Element Type", "Points").to_lowercase();
let invert = node_param_str(target, "Invert", "false") == "true";
- let (member, prim_member) = select_elements(&geom, &etype, target, hit, invert);
+ let queries: Vec<Vec3> = if etype == "primitives" {
+ (0..geom.num_prims())
+ .map(|prim| {
+ let pts = geom.prim_points(prim);
+ if pts.is_empty() {
+ Vec3::ZERO
+ } else {
+ pts.iter().map(|&p| geom.pos(p as usize)).sum::<Vec3>() / pts.len() as f32
+ }
+ })
+ .collect()
+ } else {
+ (0..geom.num_points()).map(|p| geom.pos(p)).collect()
+ };
+ let flags = match crate::collide::hits(&queries, &tris, test) {
+ Ok(f) => f,
+ Err(e) => {
+ if ocl_error.is_none() {
+ *ocl_error = Some(format!("{}: {e}", target.name));
+ }
+ crate::collide::hits_cpu(&queries, &tris, test)
+ }
+ };
+ let mut member = vec![false; geom.num_points()];
+ let mut prim_member = vec![false; geom.num_prims()];
+ match etype.as_str() {
+ "primitives" => {
+ for prim in 0..geom.num_prims() {
+ if flags[prim] != 0 {
+ prim_member[prim] = true;
+ for &p in geom.prim_points(prim) {
+ member[p as usize] = true;
+ }
+ }
+ }
+ }
+ "edges" => {
+ for e in geom.edges() {
+ let (a, b) = (e[0] as usize, e[1] as usize);
+ if flags[a] != 0 && flags[b] != 0 {
+ member[a] = true;
+ member[b] = true;
+ }
+ }
+ }
+ _ => {
+ for p in 0..geom.num_points() {
+ member[p] = flags[p] != 0;
+ }
+ }
+ }
+ if invert {
+ for m in member.iter_mut() {
+ *m = !*m;
+ }
+ for m in prim_member.iter_mut() {
+ *m = !*m;
+ }
+ }
let group_name = node_param_str(target, "Group Name", "collisions").trim().to_string();
let highlight = node_param_str(target, "Highlight", "true") == "true";
diff --git a/src/main.rs b/src/main.rs
index ffc5786..289085b 100644
--- a/src/main.rs
+++ b/src/main.rs
@@ -14,6 +14,7 @@ pub mod wrangle;
pub mod shapes;
pub mod gpu;
pub mod springs;
+pub mod collide;
// Root-level aliases some modules import via `crate::` paths.
#[allow(unused_imports)]
@@ -10739,4 +10740,113 @@ mod tests {
println!("{:>7} points, 16 passes: cpu {cpu_ms:8.2} ms gpu {gpu_ms:8.2} ms ({:?})", sys.n, ran.map(|r| r.is_ok()));
}
}
+
+ // ---- the collision test (src/collide.rs): CPU and GPU, one algorithm ----
+
+ /// A collider and a cloud of queries around and through it.
+ fn collision_fixture(rows: usize, cols: usize, queries: usize) -> (Vec<Vec3>, Vec<[Vec3; 3]>) {
+ let collider = crate::geometry::sphere_detail(Vec3::new(0.1, 0.55, -0.05), 0.7, rows, cols);
+ let tris: Vec<[Vec3; 3]> = collider
+ .triangulate(|pos, _| Vec3::from(pos))
+ .chunks_exact(3)
+ .map(|t| [t[0], t[1], t[2]])
+ .collect();
+ let pts: Vec<Vec3> = (0..queries)
+ .map(|i| {
+ let f = i as f32;
+ Vec3::new((f * 0.371).sin() * 1.1 + 0.1, (f * 0.173).cos() * 1.1 + 0.55, (f * 0.529).sin() * 1.1 - 0.05)
+ })
+ .collect();
+ (pts, tris)
+ }
+
+ /// The CPU test is the node's original loop, step for step: inside is
+ /// containment, proximity is a band, both against a sphere whose
+ /// geometry is known.
+ #[test]
+ fn collision_cpu_tests_containment_and_a_surface_band() {
+ use crate::collide::{hits_cpu, Test};
+ let (pts, tris) = collision_fixture(16, 24, 2000);
+ let centre = Vec3::new(0.1, 0.55, -0.05);
+ let inside = hits_cpu(&pts, &tris, Test::Inside);
+ let band = hits_cpu(&pts, &tris, Test::Proximity(0.05));
+ let (mut n_in, mut n_band) = (0, 0);
+ for (i, p) in pts.iter().enumerate() {
+ let r = (*p - centre).length();
+ if r < 0.7 - 0.02 {
+ assert_eq!(inside[i], 1, "point {i} at r={r} is enclosed");
+ } else if r > 0.7 + 0.02 {
+ assert_eq!(inside[i], 0, "point {i} at r={r} is outside");
+ }
+ if (r - 0.7).abs() < 0.05 - 0.02 {
+ assert_eq!(band[i], 1, "point {i} at r={r} is within the band");
+ } else if (r - 0.7).abs() > 0.05 + 0.02 {
+ assert_eq!(band[i], 0, "point {i} at r={r} is outside the band");
+ }
+ n_in += inside[i];
+ n_band += band[i];
+ }
+ assert!(n_in > 0 && n_in < pts.len() as u32 && n_band > 0, "the fixture straddles the collider: {n_in} in, {n_band} in band");
+ }
+
+ /// The cross-check: the GPU's flags are the CPU's, for both tests. A
+ /// query on a knife edge of the threshold may round either way, so a
+ /// disagreement is tolerated only there. Skips where there is no Vulkan.
+ #[test]
+ fn collision_gpu_matches_cpu() {
+ use crate::collide::{hits_cpu, hits_gpu, Test};
+ let (pts, tris) = collision_fixture(24, 36, 6000);
+ for test in [Test::Inside, Test::Proximity(0.05)] {
+ let cpu = hits_cpu(&pts, &tris, test);
+ let gpu = match crate::gpu::with_any_device(|dev| hits_gpu(dev, &pts, &tris, test)) {
+ Err(e) => {
+ println!("skipping collision_gpu_matches_cpu: {e}");
+ return;
+ }
+ Ok(r) => r.expect("the collision kernel runs"),
+ };
+ assert_eq!(cpu.len(), gpu.len());
+ let mut disagreements = 0;
+ for i in 0..cpu.len() {
+ if cpu[i] != gpu[i] {
+ let d = tris
+ .iter()
+ .map(|t| crate::geometry::point_triangle_distance_sq(pts[i], t[0], t[1], t[2]).sqrt())
+ .fold(f32::INFINITY, f32::min);
+ let margin = match test {
+ Test::Proximity(r) => (d - r).abs(),
+ Test::Inside => d,
+ };
+ assert!(margin < 1e-4, "{test:?}: query {i} differs (cpu {} gpu {}) with margin {margin}", cpu[i], gpu[i]);
+ disagreements += 1;
+ }
+ }
+ println!("collision {test:?}: {} queries x {} triangles, {disagreements} knife-edge disagreements", pts.len(), tris.len());
+ assert!(cpu.iter().any(|&h| h == 1) && cpu.iter().any(|&h| h == 0));
+ }
+ }
+
+ /// Where the auto threshold sits: the test timed over sizes. Ignored,
+ /// a measurement: `cargo test --release -p cce-designer collision_timing -- --ignored --nocapture`.
+ #[test]
+ #[ignore]
+ fn collision_timing() {
+ use crate::collide::{hits_cpu, hits_gpu, Test};
+ for (rows, cols, queries) in [(16, 24, 500), (16, 24, 5000), (32, 48, 5000), (64, 96, 20000)] {
+ let (pts, tris) = collision_fixture(rows, cols, queries);
+ let t = std::time::Instant::now();
+ let cpu = hits_cpu(&pts, &tris, Test::Proximity(0.05));
+ let cpu_ms = t.elapsed().as_secs_f64() * 1e3;
+ let t = std::time::Instant::now();
+ let gpu = crate::gpu::with_any_device(|dev| hits_gpu(dev, &pts, &tris, Test::Proximity(0.05)));
+ let gpu_ms = t.elapsed().as_secs_f64() * 1e3;
+ let same = gpu.as_ref().map(|g| g.as_ref().map(|g| *g == cpu).unwrap_or(false)).unwrap_or(false);
+ println!(
+ "{:>6} queries x {:>6} tris = {:>10} pairs: cpu {cpu_ms:9.2} ms gpu {gpu_ms:8.2} ms same={same}",
+ pts.len(),
+ tris.len(),
+ pts.len() * tris.len()
+ );
+ }
+ }
}