// GPU frustum-cull + LOD filter + stream-compaction.
//
// One invocation per way. Visible ways' vertices are claimed atomically
// from compact[] and copied from geom[]. The caller resets draw_count to 0
// (by writing [0, 1, 0, 0] to the indirect_buf) before dispatching.
//
// Vertex layout is NOT baked in: `VPV` (u32 words per vertex) is a pipeline-
// overridable constant, so ONE shader compacts both vertex streams (LAW #5 — a
// second cull.wgsl differing only in a stride would be a twin):
//
// VPV = 6 → 24 B `Vertex` (pos.xy + col.rgba as f32 bits) — the hairline lane
// VPV = 12 → 48 B `LineVertex` (pos, bisector, side, half_px, miter, inner_lim, col)
// — the pretty quad lane, which is what actually reaches the screen
//
// Using raw u32 arrays avoids WGSL struct-alignment padding (vec4 needs align 16)
// AND is what makes the stride a mere loop bound rather than a type.
struct WayMeta {
bbox_min: vec2<f32>, // offset 0
bbox_max: vec2<f32>, // offset 8
vert_start: u32, // offset 16 — index into geom[]
vert_count: u32, // offset 20
lod: u32, // offset 24 (0=country 1=region 2=city)
_pad: u32, // offset 28
} // stride 32
struct ViewportCull {
min: vec2<f32>, // offset 0
max: vec2<f32>, // offset 8
lod: u32, // offset 16 — show ways with lod <= this
_p0: u32, // offset 20
_p1: u32, // offset 24
_p2: u32, // offset 28
} // size 32
@group(0) @binding(0) var<storage, read> way_meta: array<WayMeta>;
@group(0) @binding(1) var<storage, read> geom: array<u32>;
@group(0) @binding(2) var<storage, read_write> compact: array<u32>;
@group(0) @binding(3) var<storage, read_write> draw_count: atomic<u32>;
@group(0) @binding(4) var<uniform> vp: ViewportCull;
// u32 words per vertex. Pipeline-overridable; defaults to the 24 B `Vertex` so an
// existing pipeline that sets no constants behaves exactly as before.
override VPV: u32 = 6u;
@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) id: vec3<u32>) {
let i = id.x;
if i >= arrayLength(&way_meta) { return; }
let m = way_meta[i];
// Frustum cull — reject if bbox completely outside viewport
if m.bbox_max.x < vp.min.x || m.bbox_min.x > vp.max.x { return; }
if m.bbox_max.y < vp.min.y || m.bbox_min.y > vp.max.y { return; }
// Cumulative LOD: country ways visible at all zoom levels, city only at high zoom
if m.lod > vp.lod { return; }
// Atomically claim vert_count vertex slots in compact[]
let dst = atomicAdd(&draw_count, m.vert_count);
let cap = arrayLength(&compact) / VPV;
if dst + m.vert_count > cap { return; }
// Copy vertices word-by-word (avoids struct-alignment issues). The word loop is
// bounded by the overridable VPV, so the same code moves a 24 B `Vertex` or a
// 48 B `LineVertex` without knowing what either is.
for (var v = 0u; v < m.vert_count; v++) {
let s = (m.vert_start + v) * VPV;
let d = (dst + v) * VPV;
for (var w = 0u; w < VPV; w++) {
compact[d + w] = geom[s + w];
}
}
}