// GPU frustum-cull + LOD filter + ORDER-PRESERVING stream-compaction.
//
// One invocation per way. Visible ways' vertices are copied from geom[] into
// compact[] AT THE OFFSET THEIR BAKE ORDER EARNS THEM, so the rasteriser
// consumes them in the order `facett_map::paint_order` put them in.
//
// ── WHY THIS IS THREE PASSES AND NOT ONE ────────────────────────────────────
//
// It used to be one, and the claim was `let dst = atomicAdd(&draw_count, n)`.
// That is a correct COMPACTION and a broken PAINTER: the destination offset is
// whatever order the workgroups happened to reach the atomic in, which is not
// the bake order and is not even stable between two dispatches of the same
// input. MEASURED on oden / RTX 4090, 2026-08-22, 2 000 ways / 24 000 quad
// vertices, whole-world viewport so nothing is culled at all:
//
// * compacted stream vs BAKED stream: 24 000 of 24 000 vertices differ (100 %)
// * run-to-run, same buffers, same dispatch: 21 888 of 24 000 differ (91.2 %)
//
// So every overlap the GPU lane draws — one building over its neighbour, a road
// over a lake, a bridge over a river — was decided per frame by the GPU's
// scheduler. That is the deepest layer of "the map draws in the order things
// were read in": under the read order there was no order at all. A `paint_order`
// sort on the CPU side cannot reach the screen through it.
//
// The cure is the standard ordered compaction: an exclusive prefix sum over the
// visible ways' vertex counts, so a way's destination is a function of how many
// vertices precede it in BAKE order and of nothing else.
//
// count_pass one invocation per way. Frustum + LOD test; the surviving count
// is scanned inside the workgroup (Hillis–Steele over 64 lanes)
// into way_off[i] = the way's offset WITHIN its workgroup, and the
// workgroup's total lands in wg_sum[wg]. Culled ways get CULLED.
// scan_pass one invocation, total. Turns wg_sum[] from per-workgroup totals
// into per-workgroup BASES (a serial exclusive scan — one pass over
// ceil(ways/64) words, ~1 600 iterations for a 100 000-way chunk)
// and writes the grand total into draw_count, which is what
// draw_indirect reads as vertex_count.
// main one invocation per way. dst = wg_sum[wg] + way_off[i]; copy.
//
// The three run as three dispatches in ONE compute pass. WebGPU orders dispatches
// within a pass and makes a dispatch's storage writes visible to the next, which
// is exactly the dependency chain above.
//
// draw_count is no longer accumulated — scan_pass STORES it — but the caller's
// reset to [0, 1, 0, 0] is kept: it is what makes instance_count 1.
//
// 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)
flags: u32, // offset 28 — bit 0: this way is an AREA (fill/outline)
} // stride 32
// `WayMeta.flags` bit 0. A way's fill triangles and its ring outline carry it; a road
// does not. It is the only thing that lets the minimum-extent gate below apply to the
// building footprints it is for without also deleting the street network.
const FLAG_FILL: u32 = 1u;
struct ViewportCull {
min: vec2<f32>, // offset 0
max: vec2<f32>, // offset 8
lod: u32, // offset 16 — show ways with lod <= this
// The minimum Mercator-unit extent an AREA way must reach to draw. 0 disables the
// gate. Computed ONCE per frame by `facett_map::style::min_fill_extent_merc`; the
// shader is handed the answer rather than re-deriving the rule (stinky `kalm-3`).
min_fill_extent: f32, // 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(4) var<uniform> vp: ViewportCull;
// ── binding 3: THE CONTROL BLOCK — the indirect args AND the scan scratch, in
// ONE buffer, because `max_storage_buffers_per_shader_stage` is **4** on the
// downlevel device this crate's tests deliberately open (and on plenty of real
// hardware). Two extra bindings for the scan would have made six and the
// bind-group layout would not validate at all — MEASURED: wgpu 29 refuses
// `cull_bgl` outright with "limit is 4, count was 6".
//
// ctl[0..4] DrawIndirectArgs — [vertex_count,
// instance_count, first_vertex, first_instance]
// ctl[4 .. 4+groups] per workgroup: its surviving vertex TOTAL
// (count_pass), then its BASE (scan_pass)
// ctl[4+groups .. 4+groups+ways] per way: its offset inside its workgroup,
// or CULLED
//
// `groups` is `ceil(ways / 64)` and every pass derives it the same way, so the
// three regions cannot drift. The array is `atomic<u32>` throughout because WGSL
// cannot mix atomic and plain elements in one array and `vertex_count` must be
// atomic-typed; the scan slots use `atomicLoad`/`atomicStore` purely as ordinary
// loads and stores.
@group(0) @binding(3) var<storage, read_write> ctl: array<atomic<u32>>;
fn wg_slot(g: u32) -> u32 { return 4u + g; }
fn off_slot(groups: u32, i: u32) -> u32 { return 4u + groups + i; }
fn group_count() -> u32 { return (arrayLength(&way_meta) + 63u) / 64u; }
// 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;
// `way_off[i]` for a way the cull rejected. Not a valid offset: it is larger than
// any chunk can hold, so a copy that ignored this marker would fail the cap test
// rather than write somewhere plausible.
const CULLED: u32 = 0xFFFFFFFFu;
// THE visibility test — one body, asked by count_pass and by nothing else. `main`
// does not re-ask it: it reads `way_off[i]`, so the two passes cannot disagree
// about which ways are in (which would corrupt every offset after the first
// disagreement).
fn visible(m: WayMeta) -> bool {
// Frustum cull — reject if bbox completely outside viewport
if m.bbox_max.x < vp.min.x || m.bbox_min.x > vp.max.x { return false; }
if m.bbox_max.y < vp.min.y || m.bbox_min.y > vp.max.y { return false; }
// Cumulative LOD: country ways visible at all zoom levels, city only at high zoom
if m.lod > vp.lod { return false; }
// ── MINIMUM PROJECTED EXTENT, areas only (stinky `kalm-3`) ────────────────
// MEASURED on two settled z12 panes over Stockholm: 37 378 and 34 708 building
// ways of an 80 000-way budget — 43-47 % of everything the server may send — and
// at camera 12.7 a 20 m house spans 1.67 px. The discrete LOD ladder cannot
// express "when it is big enough": a building's tier is a property of the
// building, its size on screen is a property of the camera.
//
// Lines are exempt by construction (they carry no FLAG_FILL), so no road network
// is ever broken by this — a motorway is thin and long and its bbox extent says
// nothing about whether it is worth drawing.
if (m.flags & FLAG_FILL) != 0u && vp.min_fill_extent > 0.0 {
let ext = max(m.bbox_max.x - m.bbox_min.x, m.bbox_max.y - m.bbox_min.y);
if ext < vp.min_fill_extent { return false; }
}
return true;
}
// Scratch for the in-workgroup Hillis–Steele scan. 64 words, one per lane.
var<workgroup> scan: array<u32, 64>;
@compute @workgroup_size(64)
fn count_pass(
@builtin(global_invocation_id) gid: vec3<u32>,
@builtin(local_invocation_id) lid: vec3<u32>,
@builtin(workgroup_id) wid: vec3<u32>,
) {
let i = gid.x;
let n = arrayLength(&way_meta);
// NO early return before the barriers below: a lane that leaves the workgroup
// makes `workgroupBarrier` non-uniform, which is undefined. An out-of-range or
// culled lane contributes 0 and walks the scan with everyone else.
var cnt = 0u;
var vis = false;
if i < n {
let m = way_meta[i];
vis = visible(m);
if vis { cnt = m.vert_count; }
}
scan[lid.x] = cnt;
workgroupBarrier();
// Inclusive scan over the 64 lanes. Read into a temporary, barrier, write,
// barrier — the read of a neighbour's slot must not race the write to ours.
for (var off = 1u; off < 64u; off = off << 1u) {
var add = 0u;
if lid.x >= off { add = scan[lid.x - off]; }
workgroupBarrier();
scan[lid.x] = scan[lid.x] + add;
workgroupBarrier();
}
let incl = scan[lid.x];
let groups = group_count();
if i < n {
var v = CULLED;
if vis { v = incl - cnt; }
atomicStore(&ctl[off_slot(groups, i)], v);
}
// Lane 63 always exists (the workgroup is fixed-size even for a partial tail),
// and after an inclusive scan it holds the workgroup's total.
if lid.x == 63u { atomicStore(&ctl[wg_slot(wid.x)], incl); }
}
@compute @workgroup_size(1)
fn scan_pass() {
let groups = group_count();
var run = 0u;
for (var g = 0u; g < groups; g++) {
let slot = wg_slot(g);
let total = atomicLoad(&ctl[slot]);
atomicStore(&ctl[slot], run); // in place: total -> exclusive base
run = run + total;
}
// What `draw_indirect` reads as `vertex_count`. STORED, not accumulated: the
// count is now a sum this pass computed, not a race the workgroups ran.
atomicStore(&ctl[0], run);
}
@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(workgroup_id) wid: vec3<u32>) {
let i = gid.x;
if i >= arrayLength(&way_meta) { return; }
let groups = group_count();
let local = atomicLoad(&ctl[off_slot(groups, i)]);
if local == CULLED { return; }
let m = way_meta[i];
let dst = atomicLoad(&ctl[wg_slot(wid.x)]) + local;
// The capacity guard the atomic version had. `compact` is allocated at the FULL
// geom size, so an ordered compaction cannot reach it — which is the point of
// keeping it: if it ever fires, the scan and the copy have disagreed.
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];
}
}
}