facett-core/src/render/gpu/cull.wgsl
// GPU frustum-cull + LOD filter + ORDER-PRESERVING stream-compaction.
//
// One invocation per way. Visible ways' vertices are compacted into an INDEX LIST —
// `compact[dst + v] = vert_start + v` — AT THE OFFSET THEIR BAKE ORDER EARNS THEM, so
// `draw_indexed_indirect` pulls them out of geom[] in the order
// `facett_map::paint_order` put them in. geom[] itself is never rewritten and is the
// draw's vertex buffer.
//
// ── WHY AN INDEX LIST AND NOT A VERTEX COPY (GG audit row 21) ────────────────
//
// This pass used to COPY vertices — VPV words each — into a `compact` buffer that
// therefore had to be allocated at the full size of the chunk's geometry: a second
// whole copy of every chunk, resident for the chunk's whole life, so the card's bill
// for a map was ≈ 2× its geometry. MEASURED on t14s / Radeon 890M, 2026-09-06,
// 22 000 ways (74.8 MB of geometry + meta): 147.8 MB allocated, 1.975×.
//
// An index is 4 B whatever the vertex is (24 B `Vertex`, 48 B `LineVertex`), so the
// compaction now costs 1/6 or 1/12 of the geometry instead of 1/1 — and the copy no
// longer knows a stride at all, which retires the `VPV` override: ONE pipeline
// compacts every stream (LAW #5 by construction, not by an override). The order
// guarantee is unchanged: an index stream in bake order is consumed in bake order.
//
// ── 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_indexed_indirect reads as index_count.
// main one invocation per way. dst = wg_sum[wg] + way_off[i]; write the
// way's vertex indices there.
//
// 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 of the header is kept: it is what sets instance_count (1, or 2 for the road
// lane whose casing and fill are two instances of one draw) and zeroes
// first_index / base_vertex / first_instance.
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>;
// The geometry. NOT read by any pass any more — the index list refers into it and the
// DRAW reads it as its vertex buffer. The binding stays so `cull_bgl` keeps the one
// five-entry layout every chunk's bind group is minted against; a layout may carry an
// entry a shader does not name.
@group(0) @binding(1) var<storage, read> geom: array<u32>;
// THE INDEX LIST — one u32 per surviving vertex, in bake order. The draw's index buffer.
@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..5] DrawIndexedIndirectArgs — [index_count,
// instance_count, first_index, base_vertex,
// first_instance]
// ctl[5 .. 5+groups] per workgroup: its surviving vertex TOTAL
// (count_pass), then its BASE (scan_pass)
// ctl[5+groups .. 5+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 `index_count` must be
// atomic-typed; the scan slots use `atomicLoad`/`atomicStore` purely as ordinary
// loads and stores. Rust twin of the header width: `GpuMapRenderer::ctl_bytes`.
@group(0) @binding(3) var<storage, read_write> ctl: array<atomic<u32>>;
const HEADER_WORDS: u32 = 5u;
fn wg_slot(g: u32) -> u32 { return HEADER_WORDS + g; }
fn off_slot(groups: u32, i: u32) -> u32 { return HEADER_WORDS + groups + i; }
fn group_count() -> u32 { return (arrayLength(&way_meta) + 63u) / 64u; }
// `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_indexed_indirect` reads as `index_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` holds one index per
// vertex of the chunk, so an ordered compaction cannot reach its end — which is
// the point of keeping it: if it ever fires, the scan and the copy have disagreed.
if dst + m.vert_count > arrayLength(&compact) { return; }
// The way's vertex INDICES, contiguous, in bake order. Nothing here depends on
// what a vertex is — the draw pipeline's vertex layout owns that.
for (var v = 0u; v < m.vert_count; v++) {
compact[dst + v] = m.vert_start + v;
}
}