facett-map/tests/cull_order.rs
//! **The GPU cull must compact IN ORDER** — `cull.wgsl`'s three-pass prefix sum.
//!
//! # What was wrong
//!
//! The cull claimed its destination with `let dst = atomicAdd(&draw_count, n)`. That is
//! a correct *compaction* and a broken *painter*: the offset a way lands at is whatever
//! order the workgroups happened to reach the atomic in. So the draw order of every way
//! on the GPU lane — which building covers which, whether the road is over the lake —
//! was decided per frame by the GPU's scheduler, and `facett_map::paint_order`'s sort
//! could not reach the screen through it.
//!
//! MEASURED on oden / RTX 4090, 2026-08-22, on the code as shipped, 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 %)**
//! * two dispatches of the same buffers — **21 888 of 24 000 differ (91.2 %)**
//!
//! After the ordered compaction: **0** and **0**.
//!
//! # Why the assertions are what they are
//!
//! The whole-world case alone would be a weak guard: with nothing culled, a correct
//! ordered compaction IS the identity copy, and so is a compaction that never moved
//! anything. The partial-viewport case carries the proof — the survivors are a
//! scattered subset, they must land as a contiguous prefix, and that prefix must equal
//! the survivors' spans concatenated **in bake order, as a SEQUENCE**. A multiset
//! comparison passes on the scrambled shipped code; that is exactly the assertion
//! `cull_indirect.rs` used to make, and the comment beside it said in so many words
//! that asserting an order "would be asserting an accident". It is not an accident any
//! more, and this file is why.
//!
//! Run under the box's shared GPU lock:
//! `flock ~/scratch/locks/gpu.lock cargo test -p facett-map --features wgpu --test cull_order`
#![cfg(feature = "wgpu")]
use facett_map::Way;
use facett_map::gpu::{GpuMapRenderer, build_gpu_geometry, ways_origin};
mod common;
/// `n` short roads laid out west to east, one every 0.004°, three vertices each. The
/// lengths are uniform on purpose: a cull that keeps the WRONG ways of the same size
/// produces the right vertex count, so the byte comparison — not the count — is the
/// oracle here.
fn many_roads(n: i64) -> Vec<Way> {
let c = egui::Color32::from_rgb(0xE8, 0xB0, 0x30);
(0..n)
.map(|i| {
let lon = 9.0 + (i as f64) * 0.004;
Way::line(i + 1, "road", vec![(lon, 47.0), (lon + 0.002, 47.01), (lon + 0.004, 47.0)], c, 3.0)
})
.collect()
}
/// An origin-local viewport bbox spanning `[lon0, lon1]` at latitude ~47, in exactly
/// the space `build_gpu_geometry` rebased the per-way bboxes into.
fn viewport(lon0: f64, lon1: f64, origin: [f64; 2]) -> [f32; 4] {
let (x0, y0) = facett_map::mercator(lon0, 46.0);
let (x1, y1) = facett_map::mercator(lon1, 48.0);
[
(x0 - origin[0]) as f32,
(y0.min(y1) - origin[1]) as f32,
(x1 - origin[0]) as f32,
(y0.max(y1) - origin[1]) as f32,
]
}
#[test]
fn the_cull_compacts_in_bake_order_and_does_it_the_same_way_every_time() {
let Some((device, queue)) = common::device_or_skip(
"the_cull_compacts_in_bake_order_and_does_it_the_same_way_every_time",
"the GPU cull's stream compaction preserves the declared paint order",
) else {
return;
};
let ways = many_roads(2000);
let origin = ways_origin(ways.iter()).expect("non-empty fixture has an origin");
let geo = build_gpu_geometry(ways.iter().map(|w| (w, w.style.lod as u32)), origin);
let total: u32 = geo.line_meta.iter().map(|m| m.vert_count).sum();
assert_eq!(total as usize, geo.line_vertices.len(), "the spans tile the whole quad buffer");
assert!(geo.line_meta.len() >= 500, "the fixture must spread over many workgroups (64 ways each)");
let mut r = GpuMapRenderer::new(&device, wgpu::TextureFormat::Rgba8Unorm);
r.upload_for_test(&device, &queue, &geo);
assert!(r.line_chunk_count() >= 1, "the quad lane uploaded at least one cull chunk");
let cull = |bbox: [f32; 4]| {
let mut enc = device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("cull_order") });
r.run_cull(&queue, &mut enc, bbox, 3, 0.0);
queue.submit(Some(enc.finish()));
device.poll(wgpu::PollType::wait_indefinitely()).ok();
let n: u32 = r.line_draw_counts(&device, &queue).iter().sum();
(n, r.line_compact_readback(&device, &queue, 0, n.min(total)))
};
// ── 1. NOTHING CULLED: the compacted stream is the baked stream, exactly ──
let whole = viewport(0.0, 20.0, origin);
let (n_all, a) = cull(whole);
assert_eq!(n_all, total, "a viewport containing everything must draw every quad vertex");
let scrambled = geo.line_vertices.iter().zip(&a).filter(|(x, y)| x != y).count();
assert_eq!(
scrambled, 0,
"the compacted stream must BE the baked stream — {scrambled} of {} vertices are out of \
place. The bake sorts by `paint_order`; if the cull re-orders behind it, nothing the \
CPU decided about the stack reaches the screen.",
a.len()
);
// ── 2. RUN TO RUN: the same input must give the same bytes ────────────────
let (_, b) = cull(whole);
let (_, c) = cull(whole);
let ab = a.iter().zip(&b).filter(|(x, y)| x != y).count();
let ac = a.iter().zip(&c).filter(|(x, y)| x != y).count();
assert_eq!(
(ab, ac),
(0, 0),
"two dispatches of the SAME buffers must compact identically — {ab} and {ac} of {} \
vertices moved between runs. A stack that changes every frame is a flicker, not a map.",
a.len()
);
// ── 3. THE REAL PROOF: a partial viewport, compared as a SEQUENCE ─────────
let part = viewport(11.0, 13.0, origin);
let (n_part, got) = cull(part);
assert!(
n_part > 0 && n_part < total,
"the partial viewport must cull SOMETHING and keep something: {n_part} of {total}"
);
// Which ways survive, asked of the same bbox test the shader applies, in BAKE order.
let [vx0, vy0, vx1, vy1] = part;
let survivors: Vec<usize> = geo
.line_meta
.iter()
.enumerate()
.filter(|(_, m)| {
!(m.bbox_max[0] < vx0 || m.bbox_min[0] > vx1 || m.bbox_max[1] < vy0 || m.bbox_min[1] > vy1)
})
.map(|(i, _)| i)
.collect();
let want: Vec<_> = survivors
.iter()
.flat_map(|&i| {
let m = geo.line_meta[i];
geo.line_vertices[m.vert_start as usize..(m.vert_start + m.vert_count) as usize].to_vec()
})
.collect();
assert_eq!(want.len() as u32, n_part, "the CPU's survivor set and the GPU's count must agree");
// NON-VACUOUS: the first survivor must start past 0, or this is arm 1 again and a
// shader that never relocated anything would pass.
let first = geo.line_meta[survivors[0]];
assert!(
first.vert_start > 0,
"the first survivor must start past 0 (src {}), else the compaction is an identity copy",
first.vert_start
);
let bad = got.iter().zip(&want).filter(|(x, y)| x != y).count();
assert_eq!(
bad, 0,
"the compacted survivors must appear in BAKE order — {bad} of {} vertices are wrong. \
Compared as a sequence, deliberately: as a multiset this passes on a shader that \
scrambles every way, which is what the `atomicAdd` claim did.",
want.len()
);
// ── 4. And the partial case is stable too ─────────────────────────────────
let (_, got2) = cull(part);
let drift = got.iter().zip(&got2).filter(|(x, y)| x != y).count();
assert_eq!(drift, 0, "the partial cull must be stable run to run — {drift} vertices moved");
eprintln!(
"[cull-order] whole-world {n_all}/{total} verts, 0 out of bake order, 0 run-to-run drift; \
partial {n_part} verts over {} survivors starting at src {}, 0 out of order",
survivors.len(),
first.vert_start
);
nornir_testmatrix::functional_status(
"facett-map",
"cull_compacts_in_bake_order",
scrambled == 0 && ab == 0 && bad == 0 && drift == 0,
&format!("whole {n_all}/{total} in order; partial {n_part} over {} survivors in order", survivors.len()),
);
}