graphic design tool
git clone https://git.lucas.co/cce-designer.git
src/collide.rs (8.3K)
1 //! The Collision node's test — is this point inside the collider, or within
2 //! Distance of its surface — on the CPU and on the GPU. The second Phase 7
3 //! step 4 operator, and the one the GPU is made for: every query walks EVERY
4 //! collider triangle (the CPU test has always been that brute-force loop),
5 //! so the work is queries x triangles, per-point, and embarrassingly
6 //! parallel, with one dispatch and no passes to chain.
7 //!
8 //! Same algorithm on both sides — the Voronoi-region point-triangle distance
9 //! and the Möller–Trumbore ray test, step for step — held together by
10 //! `collision_gpu_matches_cpu`. The layout is the GPU's: queries as a flat
11 //! `xyz` array, triangles as nine floats each, one `u32` flag out.
12
13 use cce_ui::vk::{Binding, ComputeDevice, Kernel};
14 use glam::Vec3;
15
16 /// Which test the node runs.
17 #[derive(Clone, Copy, Debug, PartialEq)]
18 pub enum Test {
19 /// Enclosed by the collider's volume: parity of a ray cast's crossings.
20 Inside,
21 /// Within this distance of the collider's surface.
22 Proximity(f32),
23 }
24
25 /// Fixed irrational-ish ray, NOT axis-aligned: the template meshes
26 /// tessellate on the axes, and a ray along one skims edge-on through whole
27 /// fans of triangles, double-counting crossings.
28 pub fn ray_dir() -> Vec3 {
29 Vec3::new(0.9174771, 0.3369154, 0.2095338).normalize()
30 }
31
32 /// The CPU test: one flag per query, walking every triangle.
33 pub fn hits_cpu(queries: &[Vec3], tris: &[[Vec3; 3]], test: Test) -> Vec<u32> {
34 let dir = ray_dir();
35 queries
36 .iter()
37 .map(|&p| match test {
38 Test::Proximity(d) => {
39 let d2 = d * d;
40 u32::from(tris.iter().any(|t| crate::geometry::point_triangle_distance_sq(p, t[0], t[1], t[2]) <= d2))
41 }
42 Test::Inside => {
43 let crossings = tris.iter().filter(|t| crate::spatial::ray_triangle(p, dir, t[0], t[1], t[2]).is_some()).count();
44 (crossings % 2 == 1) as u32
45 }
46 })
47 .collect()
48 }
49
50 /// The same test as a WGSL kernel: one invocation per query.
51 pub const COLLIDE_WGSL: &str = r#"
52 struct Params { n: u32, m: u32, mode: u32, pad: u32, d2: f32, rx: f32, ry: f32, rz: f32 }
53 @group(0) @binding(0) var<storage, read> queries: array<f32>;
54 @group(0) @binding(1) var<storage, read> tris: array<f32>;
55 @group(0) @binding(2) var<storage, read_write> hits: array<u32>;
56 @group(0) @binding(3) var<uniform> params: Params;
57
58 // Squared distance from p to triangle (a, b, c): the Voronoi-region walk.
59 fn dist_sq(p: vec3<f32>, a: vec3<f32>, b: vec3<f32>, c: vec3<f32>) -> f32 {
60 let ab = b - a;
61 let ac = c - a;
62 let ap = p - a;
63 let d1 = dot(ab, ap);
64 let d2 = dot(ac, ap);
65 if (d1 <= 0.0 && d2 <= 0.0) { return dot(ap, ap); }
66 let bp = p - b;
67 let d3 = dot(ab, bp);
68 let d4 = dot(ac, bp);
69 if (d3 >= 0.0 && d4 <= d3) { return dot(bp, bp); }
70 let vc = d1 * d4 - d3 * d2;
71 if (vc <= 0.0 && d1 >= 0.0 && d3 <= 0.0) {
72 let v = d1 / (d1 - d3);
73 let e = ap - ab * v;
74 return dot(e, e);
75 }
76 let cp = p - c;
77 let d5 = dot(ab, cp);
78 let d6 = dot(ac, cp);
79 if (d6 >= 0.0 && d5 <= d6) { return dot(cp, cp); }
80 let vb = d5 * d2 - d1 * d6;
81 if (vb <= 0.0 && d2 >= 0.0 && d6 <= 0.0) {
82 let w = d2 / (d2 - d6);
83 let e = ap - ac * w;
84 return dot(e, e);
85 }
86 let va = d3 * d6 - d5 * d4;
87 if (va <= 0.0 && (d4 - d3) >= 0.0 && (d5 - d6) >= 0.0) {
88 let w = (d4 - d3) / ((d4 - d3) + (d5 - d6));
89 let e = bp - (c - b) * w;
90 return dot(e, e);
91 }
92 let denom = 1.0 / (va + vb + vc);
93 let v = vb * denom;
94 let w = vc * denom;
95 let e = ap - ab * v - ac * w;
96 return dot(e, e);
97 }
98
99 // Möller–Trumbore, as spatial::ray_triangle: a hit strictly ahead of the origin.
100 fn ray_hits(o: vec3<f32>, d: vec3<f32>, v0: vec3<f32>, v1: vec3<f32>, v2: vec3<f32>) -> bool {
101 let edge1 = v1 - v0;
102 let edge2 = v2 - v0;
103 let h = cross(d, edge2);
104 let a = dot(edge1, h);
105 if (abs(a) < 1e-6) { return false; }
106 let f = 1.0 / a;
107 let s = o - v0;
108 let u = f * dot(s, h);
109 if (u < 0.0 || u > 1.0) { return false; }
110 let q = cross(s, edge1);
111 let v = f * dot(d, q);
112 if (v < 0.0 || u + v > 1.0) { return false; }
113 let t = f * dot(edge2, q);
114 return t > 1e-5;
115 }
116
117 @compute @workgroup_size(64)
118 fn collide(@builtin(global_invocation_id) id: vec3<u32>) {
119 let i = id.x;
120 if (i >= params.n) { return; }
121 let p = vec3<f32>(queries[i * 3u], queries[i * 3u + 1u], queries[i * 3u + 2u]);
122 var hit = 0u;
123 if (params.mode == 1u) {
124 for (var t = 0u; t < params.m; t = t + 1u) {
125 let b = t * 9u;
126 let a0 = vec3<f32>(tris[b], tris[b + 1u], tris[b + 2u]);
127 let a1 = vec3<f32>(tris[b + 3u], tris[b + 4u], tris[b + 5u]);
128 let a2 = vec3<f32>(tris[b + 6u], tris[b + 7u], tris[b + 8u]);
129 if (dist_sq(p, a0, a1, a2) <= params.d2) { hit = 1u; break; }
130 }
131 } else {
132 let d = vec3<f32>(params.rx, params.ry, params.rz);
133 var crossings = 0u;
134 for (var t = 0u; t < params.m; t = t + 1u) {
135 let b = t * 9u;
136 let a0 = vec3<f32>(tris[b], tris[b + 1u], tris[b + 2u]);
137 let a1 = vec3<f32>(tris[b + 3u], tris[b + 4u], tris[b + 5u]);
138 let a2 = vec3<f32>(tris[b + 6u], tris[b + 7u], tris[b + 8u]);
139 if (ray_hits(p, d, a0, a1, a2)) { crossings = crossings + 1u; }
140 }
141 hit = crossings & 1u;
142 }
143 hits[i] = hit;
144 }"#;
145
146 #[repr(C)]
147 #[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
148 struct Params {
149 n: u32,
150 m: u32,
151 mode: u32,
152 pad: u32,
153 d2: f32,
154 rx: f32,
155 ry: f32,
156 rz: f32,
157 }
158
159 /// The GPU test: one dispatch, one flag per query.
160 pub fn hits_gpu(dev: &mut ComputeDevice, queries: &[Vec3], tris: &[[Vec3; 3]], test: Test) -> Result<Vec<u32>, String> {
161 if queries.is_empty() {
162 return Ok(Vec::new());
163 }
164 let q: Vec<f32> = queries.iter().flat_map(|p| [p.x, p.y, p.z]).collect();
165 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();
166 let dir = ray_dir();
167 let params = Params {
168 n: queries.len() as u32,
169 m: tris.len() as u32,
170 mode: match test {
171 Test::Inside => 0,
172 Test::Proximity(_) => 1,
173 },
174 pad: 0,
175 d2: match test {
176 Test::Proximity(d) => d * d,
177 Test::Inside => 0.0,
178 },
179 rx: dir.x,
180 ry: dir.y,
181 rz: dir.z,
182 };
183 let mut hits = vec![0u32; queries.len()];
184 dev.run_over(
185 &Kernel::new(COLLIDE_WGSL, "collide"),
186 &mut [Binding::input(&q), Binding::input(&t), Binding::rw(&mut hits), Binding::uniform(¶ms)],
187 queries.len() as u32,
188 )?;
189 Ok(hits)
190 }
191
192 /// Below this many query-triangle pairs the CPU is faster; in auto mode it
193 /// takes them. From `collision_timing` in release on an Intel Iris Xe,
194 /// Proximity: 360k pairs cpu 2.7 ms / gpu 15 ms (that run pays the ~15 ms
195 /// pipeline compile); 3.6M pairs cpu 20 ms / gpu 2.2 ms; 15M pairs cpu
196 /// 76 ms / gpu 6.6 ms; 242M pairs cpu 1150 ms / gpu 52 ms. The GPU is
197 /// ahead from a few hundred thousand pairs up, and by 20x at the sizes
198 /// where the CPU test stalls the app.
199 pub const GPU_MIN_WORK: usize = 250_000;
200
201 /// The test as the node runs it: the backend chosen by `CCE_COMPUTE` and the
202 /// amount of work, a GPU failure in auto mode falling back to the CPU with
203 /// a note and in forced-GPU mode reported back for the node-error slot.
204 pub fn hits(queries: &[Vec3], tris: &[[Vec3; 3]], test: Test) -> Result<Vec<u32>, String> {
205 let work = queries.len().saturating_mul(tris.len());
206 if crate::gpu::use_gpu(work, GPU_MIN_WORK) {
207 match crate::gpu::with_any_device(|dev| hits_gpu(dev, queries, tris, test)) {
208 Ok(Ok(h)) => return Ok(h),
209 Ok(Err(e)) | Err(e) => {
210 if crate::gpu::choice() == crate::gpu::Choice::Gpu {
211 return Err(format!("CCE_COMPUTE=gpu but the collision test could not run there ({e}); tested on the CPU"));
212 }
213 note_fallback_once(&e);
214 }
215 }
216 }
217 Ok(hits_cpu(queries, tris, test))
218 }
219
220 fn note_fallback_once(e: &str) {
221 static ONCE: std::sync::Once = std::sync::Once::new();
222 ONCE.call_once(|| eprintln!("cce-designer: collision test fell back to the CPU: {e}"));
223 }