|
| 1 | +struct GridParams { |
| 2 | + grid_min: vec3<f32>, |
| 3 | + cell_size: f32, |
| 4 | + dims_x: u32, |
| 5 | + dims_y: u32, |
| 6 | + dims_z: u32, |
| 7 | + _pad: u32, |
| 8 | +} |
| 9 | + |
| 10 | +struct Params { |
| 11 | + connection_radius: f32, |
| 12 | + connection_ramp: f32, // falloff exponent |
| 13 | + line_alpha: f32, // opacity scale |
| 14 | + max_links: u32, // per-particle edge cap |
| 15 | +} |
| 16 | + |
| 17 | +// Auto-bound by apply(): names must be position/color. |
| 18 | +@group(0) @binding(0) var<storage, read> position: array<f32>; |
| 19 | +@group(0) @binding(1) var<storage, read> color: array<f32>; |
| 20 | +@group(0) @binding(2) var<storage, read_write> edge_pos: array<f32>; |
| 21 | +@group(0) @binding(3) var<storage, read_write> edge_col: array<f32>; |
| 22 | +@group(0) @binding(4) var<storage, read_write> indices: array<u32>; |
| 23 | +@group(0) @binding(5) var<storage, read_write> draw_args: array<atomic<u32>>; |
| 24 | +// Bound by grid.bind(). |
| 25 | +@group(0) @binding(6) var<storage, read> offsets: array<u32>; |
| 26 | +@group(0) @binding(7) var<storage, read> sorted: array<u32>; |
| 27 | +@group(0) @binding(8) var<uniform> params: Params; |
| 28 | +@group(0) @binding(9) var<uniform> gp: GridParams; |
| 29 | + |
| 30 | +fn cell_coords(p: vec3<f32>, grid_min: vec3<f32>, cell_size: f32, dims: vec3<u32>) -> vec3<i32> { |
| 31 | + let rel = (p - grid_min) / cell_size; |
| 32 | + return vec3<i32>( |
| 33 | + clamp(i32(floor(rel.x)), 0, i32(dims.x) - 1), |
| 34 | + clamp(i32(floor(rel.y)), 0, i32(dims.y) - 1), |
| 35 | + clamp(i32(floor(rel.z)), 0, i32(dims.z) - 1), |
| 36 | + ); |
| 37 | +} |
| 38 | + |
| 39 | +fn cell_index(c: vec3<u32>, dims: vec3<u32>) -> u32 { |
| 40 | + return c.x + c.y * dims.x + c.z * dims.x * dims.y; |
| 41 | +} |
| 42 | + |
| 43 | +fn load_pos(i: u32) -> vec3<f32> { |
| 44 | + return vec3<f32>(position[i * 3u], position[i * 3u + 1u], position[i * 3u + 2u]); |
| 45 | +} |
| 46 | + |
| 47 | +fn load_rgb(i: u32) -> vec3<f32> { |
| 48 | + return vec3<f32>(color[i * 4u], color[i * 4u + 1u], color[i * 4u + 2u]); |
| 49 | +} |
| 50 | + |
| 51 | +fn emit_vertex(slot: u32, p: vec3<f32>, rgb: vec3<f32>, a: f32) { |
| 52 | + edge_pos[slot * 3u + 0u] = p.x; |
| 53 | + edge_pos[slot * 3u + 1u] = p.y; |
| 54 | + edge_pos[slot * 3u + 2u] = p.z; |
| 55 | + edge_col[slot * 4u + 0u] = rgb.x; |
| 56 | + edge_col[slot * 4u + 1u] = rgb.y; |
| 57 | + edge_col[slot * 4u + 2u] = rgb.z; |
| 58 | + edge_col[slot * 4u + 3u] = a; |
| 59 | + indices[slot] = slot; |
| 60 | +} |
| 61 | + |
| 62 | +@compute @workgroup_size(64) |
| 63 | +fn main(@builtin(global_invocation_id) gid: vec3<u32>) { |
| 64 | + let i = gid.x; |
| 65 | + let count = arrayLength(&position) / 3u; |
| 66 | + if i >= count { return; } |
| 67 | + |
| 68 | + let cap = arrayLength(&indices); |
| 69 | + let pos = load_pos(i); |
| 70 | + let rgb = load_rgb(i); |
| 71 | + let radius = params.connection_radius; |
| 72 | + let r2 = radius * radius; |
| 73 | + let dims = vec3<u32>(gp.dims_x, gp.dims_y, gp.dims_z); |
| 74 | + let base = cell_coords(pos, gp.grid_min, gp.cell_size, dims); |
| 75 | + |
| 76 | + let reach = max(1, i32(ceil(radius / gp.cell_size))); |
| 77 | + let z0 = max(base.z - reach, 0); |
| 78 | + let z1 = min(base.z + reach, i32(gp.dims_z) - 1); |
| 79 | + let y0 = max(base.y - reach, 0); |
| 80 | + let y1 = min(base.y + reach, i32(gp.dims_y) - 1); |
| 81 | + let x0 = max(base.x - reach, 0); |
| 82 | + let x1 = min(base.x + reach, i32(gp.dims_x) - 1); |
| 83 | + |
| 84 | + var emitted = 0u; |
| 85 | + for (var cz = z0; cz <= z1; cz++) { |
| 86 | + for (var cy = y0; cy <= y1; cy++) { |
| 87 | + for (var cx = x0; cx <= x1; cx++) { |
| 88 | + let cell = cell_index(vec3<u32>(u32(cx), u32(cy), u32(cz)), dims); |
| 89 | + let start = offsets[cell]; |
| 90 | + let end = offsets[cell + 1u]; |
| 91 | + for (var s = start; s < end; s++) { |
| 92 | + let j = sorted[s]; |
| 93 | + if j <= i { continue; } // one direction per edge |
| 94 | + let pj = load_pos(j); |
| 95 | + let diff = pos - pj; |
| 96 | + let d2 = dot(diff, diff); |
| 97 | + if d2 > r2 { continue; } |
| 98 | + |
| 99 | + if emitted >= params.max_links { return; } |
| 100 | + emitted += 1u; |
| 101 | + |
| 102 | + let d = sqrt(d2); |
| 103 | + let a = pow(1.0 / (d / radius + 1.0), params.connection_ramp) * params.line_alpha; |
| 104 | + |
| 105 | + let slot = atomicAdd(&draw_args[0], 2u); |
| 106 | + if slot + 1u < cap { |
| 107 | + emit_vertex(slot, pos, rgb, a); |
| 108 | + emit_vertex(slot + 1u, pj, load_rgb(j), a); |
| 109 | + } |
| 110 | + } |
| 111 | + } |
| 112 | + } |
| 113 | + } |
| 114 | +} |
0 commit comments