// **GPU label collision — the three compute passes** (GFX_V2 §3.B, item 4).
//
// This is the DEVICE TRANSCRIPTION of `facett_core::label_grid::resolve_into`. Every
// constant, every hash step and every comparison below has a named twin in that
// module, and `cpu_and_gpu_lanes_resolve_the_same_labels` runs both over the same
// fixture and asserts the surviving sets are equal. If you change one side, change the
// other and watch that test.
//
// cs_clear_frame — empty every region, once per frame
// then LABEL_ROUNDS iterations of:
// cs_clear_round — empty the per-round claim + name-bid regions
// cs_claim — settle whoever last round's ink blocked, then atomicMax this
// priority into every cell the padded screen box covers
// cs_name_bid — labels that still own all their cells compete on their NAME
// cs_emit — an owner whose name is clear draws: appends its glyph quads,
// bumps `instance_count` by atomicAdd, and marks its cells
// occupied. An owner whose name is taken RETIRES, releasing its
// pixels for the next round.
//
// ── Why rounds, and why the passes are separate ──────────────────────────────────
//
// One claim/emit pair suppresses a label against any higher-priority claimant —
// including one that itself lost. A(100) beats B(90), then C(80), overlapping only B,
// loses to B's dead bid. On the real Liechtenstein clip that cascade turned 48 lettered
// roads into 3. Occupancy — written only by labels that actually DRAW — is what fixes
// it, and rounds are what let the freed pixels be taken. The greedy CPU pass got chain
// resolution free from its visit order; that is precisely the property a device cannot
// have, so it is bought back with iteration.
//
// The four passes are separate dispatches, not one kernel with barriers, and that is
// what makes the whole thing race-free rather than merely usually-right:
//
// * `cs_claim` is the only pass that READS occupancy; `cs_emit` is the only pass that
// WRITES it. They never overlap, so no invocation can see a half-written occupancy.
// * `cs_name_bid` writes name bids; `cs_emit` reads them. Same separation.
// * the per-label state word is written only by its own invocation.
//
// ── Why the doc's own sketch is three passes and not one ─────────────────────────
//
// GFX_V2 §3.B carries a `cs_collide` sketch that FUSES claim and emit:
//
// let prev = atomicMax(&spatial_grid[cell_idx], label.priority);
// if (label.priority >= prev) { ... visible_labels[i] = label; }
//
// That form is racy, and not subtly: `prev` is only the maximum of the labels that
// happened to arrive FIRST. With priorities 1, 2, 3 arriving in that order, label 1
// sees prev 0 and emits, label 2 sees prev 1 and emits, label 3 sees prev 2 and
// emits — all three draw on top of each other. It is correct only for exactly two
// labels per cell, or when the highest priority happens to arrive first. The doc's
// numbered list (clear / test / indirect dispatch) already describes the correct
// three-pass shape; only the code sketch collapses it. The barrier between claim and
// emit is what makes the answer order-independent, which is the whole reason the rule
// in `label_grid` was designed around a per-label `grid[cell] == priority` check
// rather than around the return value of `atomicMax`.
// ── Constants — mirrored in facett_core::label_grid ──────────────────────────────
const GRID_W: u32 = 256u;
const GRID_H: u32 = 192u;
const GRID_CELLS: u32 = GRID_W * GRID_H; // 49152
const NAME_SLOTS: u32 = 16384u;
const MAX_VISIBLE_LABELS: u32 = 48u;
const HALO_TAPS: u32 = 4u;
const INSTANCES_PER_GLYPH: u32 = HALO_TAPS + 1u;
// Per-label resolution state, mirroring `label_grid::LS_*`.
const LS_PENDING: u32 = 0u;
const LS_WON: u32 = 1u;
const LS_BLOCKED: u32 = 2u;
const LS_RETIRED: u32 = 3u;
// Candidate flag bits.
const FLAG_SCREEN_SPACE: u32 = 1u; // `pos` is already screen px, do not project
const FLAG_BLOCKER: u32 = 2u; // reserves cells, draws nothing
// ── Buffers ─────────────────────────────────────────────────────────────────────
// One label offered to the pass. 48 B; mirrors `LabelCandidate`.
struct Candidate {
pos: vec2<f32>, // 0 origin-local Mercator, or screen px (see flags)
half_px: vec2<f32>, // 8 padded half extent, screen px
priority: u32, // 16 strictly unique, higher wins
name_hash: u32, // 20
glyph_start: u32, // 24 index into glyph_src
glyph_count: u32, // 28
lod: u32, // 32
flags: u32, // 36
_pad0: u32, // 40
_pad1: u32, // 44
} // stride 48
// One glyph of a label, positioned RELATIVE to the label's box centre. 32 B.
struct GlyphSrc {
off_min: vec2<f32>, // 0 label-local px
off_max: vec2<f32>, // 8
uv_min: vec2<f32>, // 16 atlas UV (0..1)
uv_max: vec2<f32>, // 24
} // stride 32
// What the draw consumes: an absolute screen-px quad + its atlas UV + its colour.
// 48 B; mirrors `LabelGlyphInstance` and the vertex attributes in `label_draw.wgsl`.
struct GlyphOut {
rect_min: vec2<f32>,
rect_max: vec2<f32>,
uv_min: vec2<f32>,
uv_max: vec2<f32>,
color: vec4<f32>,
}
struct LabelUniforms {
// Origin-local -> screen px, the same transform `draw.wgsl` applies to vertices.
zoom: vec2<f32>, // 0
ref_pos: vec2<f32>, // 8
screen_center: vec2<f32>, // 16
viewport: vec2<f32>, // 24 screen px
inv_cell: vec2<f32>, // 32 GRID_W/viewport.x, GRID_H/viewport.y
inv_repeat: f32, // 40 1 / REPEAT_CELL_PX
lod: u32, // 44
ink: vec4<f32>, // 48 (align 16)
halo: vec4<f32>, // 64
} // size 80
// ── The one state buffer: draw args, then counters, then the grid ───────────────
//
// `state` is a flat `array<atomic<u32>>` holding, in order:
//
// word 0 vertex_count ) the DrawIndirectArgs `draw_indirect` reads at
// word 1 instance_count ) offset 0 — word 1 is what pass 3 atomicAdds onto
// word 2 first_vertex )
// word 3 first_instance )
// word 4..7 labels_won / labels_drawn / glyph_overflow / pad
// word 8.. claim grid, occupancy grid, name bids, won names,
// then ONE word per candidate holding its LS_* state
//
// Raw u32 words rather than a struct, for cull.wgsl's stated reason: WGSL would pad a
// struct to `vec4` alignment, and a flat array makes the offsets plain arithmetic the
// Rust side mirrors word for word. cull.wgsl already binds its `indirect_buf` this way
// (`draw_count: atomic<u32>` IS `vertex_count` at offset 0); this extends the same
// trick rather than inventing a second convention.
//
// One buffer instead of three also keeps the pass inside
// `max_storage_buffers_per_shader_stage = 4` — the floor `Limits::downlevel_defaults()`
// still enforces. Splitting the counters out cost a fifth binding and made the pass
// undeployable on anything but a full-limits device, for no gain: the counters are read
// and reset in lockstep with the count they gate.
const W_VERTEX_COUNT: u32 = 0u;
const W_INSTANCE_COUNT: u32 = 1u;
const W_LABELS_WON: u32 = 4u;
const W_LABELS_DRAWN: u32 = 5u;
const W_OVERFLOW: u32 = 6u;
const GRID_BASE: u32 = 8u;
// The four grid regions, mirroring `label_grid::{CLAIM,OCC,NAME_CLAIM,NAME_WON}_BASE`.
// CLAIM and NAME_CLAIM are emptied every round; OCC and NAME_WON persist for the frame.
const CLAIM_BASE: u32 = GRID_BASE;
const OCC_BASE: u32 = GRID_BASE + GRID_CELLS;
const NAME_CLAIM_BASE: u32 = GRID_BASE + 2u * GRID_CELLS;
const NAME_WON_BASE: u32 = GRID_BASE + 2u * GRID_CELLS + NAME_SLOTS;
const GRID_END: u32 = GRID_BASE + 2u * GRID_CELLS + 2u * NAME_SLOTS;
// One word per candidate, holding its LS_* state. The buffer is sized to
// GRID_END + candidate count at upload, so this region always exists.
const STATE_BASE: u32 = GRID_END;
@group(0) @binding(0) var<storage, read> cands: array<Candidate>;
@group(0) @binding(1) var<storage, read> glyph_src: array<GlyphSrc>;
@group(0) @binding(2) var<storage, read_write> state: array<atomic<u32>>;
@group(0) @binding(3) var<storage, read_write> out_glyphs: array<GlyphOut>;
@group(0) @binding(4) var<uniform> U: LabelUniforms;
// ── Shared geometry (the twin of `label_grid::footprint`) ────────────────────────
// A label's inclusive cell range plus its coarse repeat cell. `ok == false` means "not
// a candidate this frame" — wrong LOD, off-screen, or a non-finite box.
struct Footprint {
ok: bool,
cx0: u32,
cy0: u32,
cx1: u32,
cy1: u32,
ncx: i32,
ncy: i32,
}
fn label_center_px(c: Candidate) -> vec2<f32> {
if (c.flags & FLAG_SCREEN_SPACE) != 0u {
return c.pos;
}
// The exact form of `MapTransform::project` / `draw.wgsl`'s vs_main: subtract
// before scaling, so a street-level zoom does not differ two ~4e7 f32.
return (c.pos - U.ref_pos) * U.zoom + U.screen_center;
}
fn span_cells(lo: f32, hi: f32, inv: f32, n: u32) -> vec2<u32> {
let top = f32(n - 1u);
let a = u32(clamp(floor(lo * inv), 0.0, top));
let b = u32(clamp(floor(hi * inv), 0.0, top));
return vec2<u32>(min(a, b), max(a, b));
}
fn footprint_of(c: Candidate) -> Footprint {
var f: Footprint;
f.ok = false;
if c.lod > U.lod {
return f;
}
let ctr = label_center_px(c);
let h = abs(c.half_px);
// A NaN fails every comparison, so this rejects non-finite boxes too; `clamp` on a
// NaN is undefined in WGSL, which is precisely what must not reach `span_cells`.
if !(ctr.x == ctr.x && ctr.y == ctr.y && h.x == h.x && h.y == h.y) {
return f;
}
let lo = ctr - h;
let hi = ctr + h;
// Reject off-screen BEFORE clamping. Clamping first would pile every off-screen
// label into the edge cells, where it would suppress the visible labels that
// genuinely live there (the CPU twin's `lod_and_offscreen` arm pins this).
if hi.x < 0.0 || hi.y < 0.0 || lo.x > U.viewport.x || lo.y > U.viewport.y {
return f;
}
let sx = span_cells(lo.x, hi.x, U.inv_cell.x, GRID_W);
let sy = span_cells(lo.y, hi.y, U.inv_cell.y, GRID_H);
f.ok = true;
f.cx0 = sx.x;
f.cx1 = sx.y;
f.cy0 = sy.x;
f.cy1 = sy.y;
f.ncx = i32(floor(ctr.x * U.inv_repeat));
f.ncy = i32(floor(ctr.y * U.inv_repeat));
return f;
}
// `label_grid::repeat_slot`, integer for integer so the two lanes cannot diverge.
fn repeat_slot(name_hash: u32, ncx: i32, ncy: i32) -> u32 {
var h = name_hash ^ (bitcast<u32>(ncx) * 0x9E3779B9u) ^ (bitcast<u32>(ncy) * 0x85EBCA6Bu);
h ^= h >> 15u;
h = h * 0x2C1B3C6Du;
h ^= h >> 12u;
return h % NAME_SLOTS;
}
// ── Per-label helpers ───────────────────────────────────────────────────────────
// Is any cell this label covers already occupied by a label that DREW? A pending label
// never occupies anything, so "occupied at all" means "occupied by somebody else".
fn blocked(f: Footprint) -> bool {
for (var cy = f.cy0; cy <= f.cy1; cy++) {
for (var cx = f.cx0; cx <= f.cx1; cx++) {
if atomicLoad(&state[OCC_BASE + cy * GRID_W + cx]) != 0u {
return true;
}
}
}
return false;
}
// Does this label still own every claim cell it bid for?
//
// Equality, not `>=`: `atomicMax` left exactly one value per cell and priorities are
// unique by construction, so `==` identifies THE winner. `>=` would let a runner-up
// through wherever it tied.
fn owns_claim(f: Footprint, priority: u32) -> bool {
for (var cy = f.cy0; cy <= f.cy1; cy++) {
for (var cx = f.cx0; cx <= f.cx1; cx++) {
if atomicLoad(&state[CLAIM_BASE + cy * GRID_W + cx]) != priority {
return false;
}
}
}
return true;
}
// ── Pass 1a: clear the whole frame ──────────────────────────────────────────────
//
// Every grid region plus the per-label state (zero == LS_PENDING). Forgetting this does
// not blank the map on frame 1 — wgpu zero-initialises the buffer — it leaves LAST
// frame's priorities standing, so a later frame's labels lose to ghosts. The proof
// resolves the same collider at two different LODs to catch it; replaying an identical
// frame cannot, because it re-claims the same cells with the same priorities.
@compute @workgroup_size(64)
fn cs_clear_frame(@builtin(global_invocation_id) id: vec3<u32>) {
let i = GRID_BASE + id.x;
if i >= arrayLength(&state) {
return;
}
atomicStore(&state[i], 0u);
}
// ── Pass 1b: clear the per-round regions ────────────────────────────────────────
//
// CLAIM and NAME_CLAIM only. Occupancy and the won-name marks must survive the round,
// or a label would stop being blocked by ink already on the paper.
@compute @workgroup_size(64)
fn cs_clear_round(@builtin(global_invocation_id) id: vec3<u32>) {
let n = id.x;
if n < GRID_CELLS {
atomicStore(&state[CLAIM_BASE + n], 0u);
} else if n < GRID_CELLS + NAME_SLOTS {
atomicStore(&state[NAME_CLAIM_BASE + (n - GRID_CELLS)], 0u);
}
}
// ── Pass 2: settle, then claim ──────────────────────────────────────────────────
@compute @workgroup_size(64)
fn cs_claim(@builtin(global_invocation_id) id: vec3<u32>) {
let i = id.x;
if i >= arrayLength(&cands) {
return;
}
if atomicLoad(&state[STATE_BASE + i]) != LS_PENDING {
return;
}
let c = cands[i];
let f = footprint_of(c);
if !f.ok {
return;
}
// Whatever drew last round decides this label permanently.
if blocked(f) {
atomicStore(&state[STATE_BASE + i], LS_BLOCKED);
return;
}
for (var cy = f.cy0; cy <= f.cy1; cy++) {
for (var cx = f.cx0; cx <= f.cx1; cx++) {
atomicMax(&state[CLAIM_BASE + cy * GRID_W + cx], c.priority);
}
}
}
// ── Pass 3: name bid — only labels that won their PIXELS compete on their NAME ──
//
// This ordering is the repeat filter. Bidding during the claim pass instead let a
// spatially-contested label hold its name slot and suppress the one instance of that
// street which could have drawn. The CPU pass never had the problem because its
// `placed_names` only ever held names it had actually PLACED.
@compute @workgroup_size(64)
fn cs_name_bid(@builtin(global_invocation_id) id: vec3<u32>) {
let i = id.x;
if i >= arrayLength(&cands) {
return;
}
if atomicLoad(&state[STATE_BASE + i]) != LS_PENDING {
return;
}
let c = cands[i];
let f = footprint_of(c);
if !f.ok || !owns_claim(f, c.priority) {
return;
}
atomicMax(&state[NAME_CLAIM_BASE + repeat_slot(c.name_hash, f.ncx, f.ncy)], c.priority);
}
// ── Pass 4: emit + indirect dispatch ────────────────────────────────────────────
@compute @workgroup_size(64)
fn cs_emit(@builtin(global_invocation_id) id: vec3<u32>) {
let i = id.x;
if i >= arrayLength(&cands) {
return;
}
if atomicLoad(&state[STATE_BASE + i]) != LS_PENDING {
return;
}
let c = cands[i];
let f = footprint_of(c);
if !f.ok || !owns_claim(f, c.priority) {
return;
}
// The repeat filter: this name must not have been taken at a HIGHER priority in this
// coarse cell or any of its 8 neighbours — as a better BID this round, or as ink
// already on the paper.
for (var dy = -1; dy <= 1; dy++) {
for (var dx = -1; dx <= 1; dx++) {
let slot = repeat_slot(c.name_hash, f.ncx + dx, f.ncy + dy);
if atomicLoad(&state[NAME_CLAIM_BASE + slot]) > c.priority
|| atomicLoad(&state[NAME_WON_BASE + slot]) > c.priority {
// RETIRE, do not merely skip. This label owns its pixels but will not
// draw; leaving it PENDING would have it win the claim again next round,
// be rejected again, and hold that patch of paper against every
// lower-priority label for the whole frame. That zombie made the repeat
// filter REDUCE the distinct-name count it exists to raise (19 -> 12 on
// real data), and no number of extra rounds could clear it.
atomicStore(&state[STATE_BASE + i], LS_RETIRED);
return;
}
}
}
// ── it draws ──────────────────────────────────────────────────────────────
atomicStore(&state[STATE_BASE + i], LS_WON);
for (var cy = f.cy0; cy <= f.cy1; cy++) {
for (var cx = f.cx0; cx <= f.cx1; cx++) {
atomicMax(&state[OCC_BASE + cy * GRID_W + cx], c.priority);
}
}
atomicMax(&state[NAME_WON_BASE + repeat_slot(c.name_hash, f.ncx, f.ncy)], c.priority);
// A blocker occupies its pixels and prints nothing — which is how e88969f's "the
// address pins reserve their boxes FIRST" survives onto a lane with no ordering. It
// must reach the occupancy write above, so this return is here and not at the top.
if (c.flags & FLAG_BLOCKER) != 0u || c.glyph_count == 0u {
return;
}
// The ink budget. `W_LABELS_WON` counts the labels that passed the rule AND doubles
// as the budget allocator, because those are the same set counted in the same order
// — a second counter beside it would only be a second place to drift. Above the
// budget the survivors are the ones that arrived first, which is nondeterministic;
// see `label_grid::MAX_VISIBLE_LABELS`.
let slot = atomicAdd(&state[W_LABELS_WON], 1u);
if slot >= MAX_VISIBLE_LABELS {
return;
}
let n = c.glyph_count * INSTANCES_PER_GLYPH;
// THE indirect dispatch: the draw's instance count is produced here and consumed by
// `draw_indirect` without the CPU ever learning it.
let base = atomicAdd(&state[W_INSTANCE_COUNT], n);
if base + n > arrayLength(&out_glyphs) {
// Leaving the count past capacity would make `draw_indirect` read off the end of
// the vertex buffer. Give the claim back and record it: unlike the cull's
// pre-item-2 leak, this counter drives a real draw.
atomicSub(&state[W_INSTANCE_COUNT], n);
atomicAdd(&state[W_OVERFLOW], 1u);
return;
}
// Counted HERE, past both gates, so `labels_drawn` is exactly the number of labels
// whose glyphs are in `out_glyphs`. It used to BE the budget allocator and so counted
// every label that REACHED the gate — a number that reads correct on every fixture
// under the budget and is silently wrong on one over it.
// `a_dense_field_with_no_overlap_culls_nothing` is the arm that can tell the
// difference: 300 labels win the rule, 48 reach the paper.
atomicAdd(&state[W_LABELS_DRAWN], 1u);
// `label_grid::HALO_OFFSETS`, then the ink pass at zero offset. The ink is the LAST
// instance of each glyph's block, so within one label it rasterises over its own
// halo. (A `var`, not a module `const`: a const array cannot be indexed by a runtime
// value — the same reason `msdf_vs` declares its `corners` locally.)
var taps = array<vec2<f32>, 5>(
vec2<f32>( 1.0, 0.0),
vec2<f32>(-1.0, 0.0),
vec2<f32>( 0.0, 1.0),
vec2<f32>( 0.0, -1.0),
vec2<f32>( 0.0, 0.0),
);
let ctr = label_center_px(c);
for (var g = 0u; g < c.glyph_count; g++) {
let src = glyph_src[c.glyph_start + g];
for (var t = 0u; t < INSTANCES_PER_GLYPH; t++) {
let is_ink = t == HALO_TAPS;
let o = ctr + taps[t];
var out: GlyphOut;
out.rect_min = o + src.off_min;
out.rect_max = o + src.off_max;
out.uv_min = src.uv_min;
out.uv_max = src.uv_max;
out.color = select(U.halo, U.ink, is_ink);
out_glyphs[base + g * INSTANCES_PER_GLYPH + t] = out;
}
}
}