facett-map/tests/cull_indirect.rs
//! **GFX_V2 item 2 — the GPU cull really drives the draw** (only with `--features wgpu`).
//!
//! Before this file existed, NOTHING in the tree proved the cull compute pass runs.
//! The survey that opened this work found the pass had been orphaned: it compacted
//! the 24 B hairline stream into `compact_buf` and bumped `indirect_buf` every frame
//! while `paint` drew a completely separate CPU-baked quad buffer over a CPU LOD
//! prefix, with no frustum test at all. The `draw_pipeline` built from `DRAW_WGSL`
//! was never bound. Deleting the whole compute pass would not have failed a single
//! test — the definition of dead GPU work behind a green suite.
//!
//! So the assertions here are on the number the DEVICE produced and the draw
//! consumes: `DrawIndirectArgs::vertex_count`, read back out of the very buffer
//! `draw_indirect` reads. Not `renderer.drew()`, not a chunk count, not a CPU-side
//! `Vec::len()`. If the cull stopped running, or culled the wrong ways, or wrote the
//! wrong stride, this number moves.
//!
//! Run: `cargo test -p facett-map --features wgpu --test cull_indirect -- --nocapture`
#![cfg(feature = "wgpu")]
use facett_map::Way;
use facett_map::gpu::{GpuGeometry, GpuMapRenderer, build_gpu_geometry, ways_origin};
/// The ONE GPU gate for this crate's device-backed tests (`tests/common/mod.rs`):
/// a missing adapter becomes a SKIP row in the matrix and a panic under CI, never a
/// silent `ok`.
mod common;
/// Four roads, each in its own well-separated 0.1°-wide lon band, so a viewport can
/// select any prefix of them by x alone and the answer is unambiguous.
///
/// Deliberately NOT all one LOD: the last is city-tier, so the LOD arm has something
/// to reject that the frustum arm would keep. A fixture where every way shares a tier
/// cannot tell the two tests apart.
fn banded_roads() -> Vec<Way> {
let c = egui::Color32::from_rgb(0xE8, 0xB0, 0x30);
let mut ways = vec![
Way::line(1, "road", vec![(9.00, 47.00), (9.05, 47.02), (9.10, 47.00)], c, 6.0),
Way::line(2, "road", vec![(10.00, 47.00), (10.05, 47.02), (10.10, 47.00)], c, 5.0),
Way::line(3, "road", vec![(11.00, 47.00), (11.10, 47.00)], c, 4.0),
Way::line(4, "road", vec![(12.00, 47.00), (12.04, 47.01), (12.07, 47.00), (12.10, 47.02)], c, 3.0),
];
ways[3].style.lod = facett_map::Lod::City;
ways
}
/// **Where each fixture way's span landed.** `build_gpu_geometry` emits the line
/// stream in `facett_map::way_paint_key` order (the declared paint order) rather than
/// in the vector's order, so `line_meta[i]` is no longer `ways[i]`. This is the same
/// sort the bake does, asked of the fixture, so the test indexes by WAY and never by
/// the accident of where a way sat in the input vector.
fn meta_slot(ways: &[Way]) -> Vec<usize> {
let mut order: Vec<usize> = (0..ways.len()).collect();
order.sort_by_key(|&i| facett_map::way_paint_key(&ways[i]));
let mut slot = vec![usize::MAX; ways.len()];
for (s, &i) in order.iter().enumerate() {
slot[i] = s;
}
slot
}
fn geometry(ways: &[Way]) -> (GpuGeometry, [f64; 2]) {
let origin = ways_origin(ways.iter()).expect("non-empty fixture has an origin");
(build_gpu_geometry(ways.iter().map(|w| (w, w.style.lod as u32)), origin), origin)
}
/// 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,
]
}
/// Drive the REAL cull dispatch (`GpuMapRenderer::run_cull`, the same call
/// `MapPaintCallback::prepare` makes) and read back what the draw would use.
fn culled_vertex_count(
renderer: &GpuMapRenderer,
device: &wgpu::Device,
queue: &wgpu::Queue,
bbox: [f32; 4],
lod: u32,
) -> u32 {
let mut enc = device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("cull_test") });
renderer.run_cull(queue, &mut enc, bbox, lod, 0.0);
queue.submit(Some(enc.finish()));
device.poll(wgpu::PollType::wait_indefinitely()).ok();
renderer.line_draw_counts(device, queue).iter().sum()
}
/// **The proof.** The vertex count `draw_indirect` consumes is computed on the GPU,
/// tracks the viewport and the LOD, and equals — exactly, not approximately — the
/// quad vertices of the ways that should have survived.
///
/// Exactness is the point. A "more ways visible ⇒ count went up" assertion would
/// pass on a cull that kept everything, and a count-based assertion alone would pass
/// on a cull that kept the WRONG ways of the same size. Both are pinned here by
/// comparing against `line_meta`, the per-way spans the CPU baked.
#[test]
fn gpu_cull_drives_the_indirect_draw_count() {
let ways = banded_roads();
let (geo, origin) = geometry(&ways);
// The per-way quad spans the cull is expected to select from.
assert_eq!(geo.line_meta.len(), 4, "one meta per way");
let total: u32 = geo.line_meta.iter().map(|m| m.vert_count).sum();
assert_eq!(
total as usize,
geo.line_vertices.len(),
"the per-way spans must tile the whole quad buffer with no gap or overlap"
);
// Non-vacuous fixture: the ways differ in size, so a cull that keeps the wrong
// subset cannot accidentally produce the right count.
let slot = meta_slot(&ways);
// Per-WAY quad counts (`sizes[k]` = fixture way k+1), not per-stream-position.
let sizes: Vec<u32> = slot.iter().map(|&s| geo.line_meta[s].vert_count).collect();
assert!(
sizes.iter().collect::<std::collections::HashSet<_>>().len() > 1,
"the ways must have DIFFERENT vertex counts, else the count proves nothing about WHICH survived: {sizes:?}"
);
let Some((device, queue)) = common::device_or_skip("gpu_cull_drives_the_indirect_draw_count", "the compute cull drives the indirect draw count")
else {
return;
};
let mut renderer = GpuMapRenderer::new(&device, wgpu::TextureFormat::Rgba8Unorm);
renderer.upload_for_test(&device, &queue, &geo);
assert!(renderer.line_chunk_count() >= 1, "the quad lane uploaded at least one cull chunk");
// The DEEPEST tier, so nothing is rejected by LOD and the frustum is the only filter.
// This was a literal `2` meaning "city"; when `Lod::District` was inserted at index 2
// the literal started meaning "district" and this test began measuring the frustum
// through a half-closed LOD gate. Derived, it means what it says.
let lod_all = facett_map::style::Lod::Detail as u32;
// ── 1. The whole world: every way survives ────────────────────────────────
let all = culled_vertex_count(&renderer, &device, &queue, viewport(0.0, 20.0, origin), lod_all);
assert_eq!(all, total, "a viewport containing everything must draw every quad vertex");
// ── 2. A viewport over the first two bands only ───────────────────────────
let want_two: u32 = sizes[0] + sizes[1];
let two = culled_vertex_count(&renderer, &device, &queue, viewport(8.9, 10.2, origin), lod_all);
assert_eq!(
two, want_two,
"a viewport over bands 1-2 must draw exactly those two ways' quads \
({want_two} = {} + {}), got {two} of a possible {total}",
sizes[0], sizes[1]
);
assert!(two < all, "culling must actually REMOVE something: {two} vs {all}");
// ── 3. A viewport over the third band only ────────────────────────────────
let one = culled_vertex_count(&renderer, &device, &queue, viewport(10.9, 11.2, origin), lod_all);
assert_eq!(one, sizes[2], "a viewport over band 3 alone must draw exactly way 3's quads");
// ── 4. THE NON-VACUOUS ANCHOR: nothing visible draws nothing ──────────────
let none = culled_vertex_count(&renderer, &device, &queue, viewport(30.0, 31.0, origin), lod_all);
assert_eq!(none, 0, "an empty viewport must draw NOTHING — got {none} vertices out of {total}");
// ── 5. LOD is applied per way, on the GPU ─────────────────────────────────
// Way 4 is city-tier. At country LOD it must drop even though it is inside the
// frustum — and the other three must remain, so this cannot pass by dropping all.
let city_way = sizes[3];
let country = culled_vertex_count(&renderer, &device, &queue, viewport(0.0, 20.0, origin), 0);
assert_eq!(
country,
total - city_way,
"at country LOD the city-tier way ({city_way} verts) must drop and the rest stay: \
expected {}, got {country}",
total - city_way
);
assert!(country > 0, "the LOD filter must not drop everything — that would pass arm 4 vacuously");
// ── 6. The compacted BYTES, not just the count ────────────────────────────
// `atomicAdd` accumulates `vert_count` whatever the stride is, so a cull
// specialised with the wrong `VPV` returns the RIGHT count over garbage. This
// reads what the vertex shader will actually pull.
//
// TWO TRAPS had to be designed out, and the first version of this check fell
// into both — it passed with the line pipeline's `VPV` override deleted:
//
// (a) IDENTITY COPY. When the survivors are a PREFIX of the buffer, `dst`
// accumulates in the same order as `vert_start`, so source and destination
// indices are equal and the "copy" is a no-op that lands correct bytes at
// correct offsets for ANY stride. So the selection below is the LAST two
// ways: `vert_start` 24/30 compacting down to `dst` 0/6, a genuine
// relocation that a wrong stride cannot fake.
// (b) RESIDUE. `compact_buf` is never cleared between frames, so an earlier
// whole-world cull leaves correct bytes lying in the range a later,
// narrower cull only partly rewrites. Hence a FRESH renderer here, whose
// compact buffer wgpu zero-initialises, so anything the cull fails to write
// reads as zero instead of as last frame's right answer.
let mut fresh = GpuMapRenderer::new(&device, wgpu::TextureFormat::Rgba8Unorm);
fresh.upload_for_test(&device, &queue, &geo);
let want_last: u32 = sizes[2] + sizes[3];
let last_two = culled_vertex_count(&fresh, &device, &queue, viewport(10.9, 12.2, origin), lod_all);
assert_eq!(last_two, want_last, "the last two bands survive this viewport");
let got = fresh.line_compact_readback(&device, &queue, 0, want_last);
assert_eq!(got.len() as u32, want_last, "read back exactly the vertices the draw would consume");
// In STREAM order, not fixture order: `build_gpu_geometry` emits by declared paint
// key, and the compaction now preserves that, so the expected sequence is the two
// survivors ordered by where their spans SIT — not by where they sat in the input
// vector (way 4 is `Lod::City` and therefore ranks below the country-tier ways).
let (m2, m3) = {
let (a, b) = (&geo.line_meta[slot[2]], &geo.line_meta[slot[3]]);
if a.vert_start <= b.vert_start { (a, b) } else { (b, a) }
};
assert!(
m2.vert_start > 0 || m3.vert_start > 0,
"the compaction must RELOCATE at least one survivor (srcs {} / {}), not copy in place",
m2.vert_start,
m3.vert_start
);
let want: Vec<_> = geo.line_vertices[m2.vert_start as usize..(m2.vert_start + m2.vert_count) as usize]
.iter()
.chain(&geo.line_vertices[m3.vert_start as usize..(m3.vert_start + m3.vert_count) as usize])
.copied()
.collect();
// Compared as a SEQUENCE. It used to be a multiset, with a comment saying that
// "ways claim their slots with `atomicAdd`, so which way lands first is genuinely
// nondeterministic and asserting an order would be asserting an accident". That
// was true and it was the bug: the GPU lane's paint order was the workgroup
// scheduler's. `cull.wgsl` compacts through a prefix sum now, so bake order is a
// guarantee, and `tests/cull_order.rs` is the guard that holds it. Sorting here
// would throw that guarantee away on the very stream that carries it.
// Reported as a mismatch COUNT plus the first offender: dumping two 24-element
// vertex vectors buries the diagnosis in two screens of float.
let bad = got.iter().zip(&want).filter(|(a, b)| a != b).count();
let zeros = got.iter().filter(|v| bytemuck::bytes_of(*v).iter().all(|b| *b == 0)).count();
assert!(
bad == 0,
"the compacted stream must be the surviving ways' quads byte-for-byte — {bad} of {} \
vertices differ ({zeros} are all-zero, i.e. never written). A wrong VPV stride yields \
the RIGHT count over shifted garbage, which is exactly what this catches.\n first bad \
got: {:?}\n first bad want: {:?}",
got.len(),
got.iter().zip(&want).find(|(a, b)| a != b).map(|(a, _)| a),
got.iter().zip(&want).find(|(a, b)| a != b).map(|(_, b)| b),
);
// Non-vacuous: a proper subset, and not all-zeros (which is what a stride that
// under-copies actually leaves behind on a fresh buffer).
assert!(
(want.len() as u32) < total,
"the compacted set must be a PROPER subset of the uploaded quads ({} of {total})",
want.len()
);
assert!(
got.iter().any(|v| v.pos != [0.0, 0.0] && v.half_px != 0.0),
"the compacted stream must contain real geometry, not a zeroed buffer"
);
eprintln!(
"[cull] GPU-computed draw counts — all={all} bands1-2={two} band3={one} empty={none} \
country-lod={country} (total {total}, per-way {sizes:?}); relocated+compacted {} verts match byte-for-byte",
got.len()
);
nornir_testmatrix::functional_status(
"facett-map",
"gpu_cull_drives_the_indirect_draw_count",
all == total && none == 0 && two == want_two && country == total - city_way,
&format!("indirect vertex_count: all={all} two={two} one={one} none={none} country={country} total={total}"),
);
}
/// Regrouping the quad expansion from one whole-tier call into one call per way (so
/// each way gets a `line_meta` span) must not change a single vertex.
///
/// This is the guard for the refactor that made per-way culling possible at all. It
/// needs no GPU, and it is the reason the visual output is unchanged: the bytes
/// uploaded are identical, only their bookkeeping is new.
#[test]
fn per_way_quad_expansion_concatenates_to_the_whole_tier_expansion() {
use facett_map::gpu::line::{LinePolyline, build_line_geometry};
let ways = banded_roads();
let (geo, origin) = geometry(&ways);
// Rebuild the country tier the OLD way: one `build_line_geometry` over all of it —
// in the DECLARED order the bake now emits, so this compares expansion against
// expansion and not expansion against a re-sort.
let slot = meta_slot(&ways);
let mut country: Vec<&Way> = ways.iter().filter(|w| w.style.lod as u32 == 0).collect();
country.sort_by_key(|w| facett_map::way_paint_key(w));
let polys: Vec<LinePolyline> = country
.iter()
.map(|w| LinePolyline {
pts: w
.lonlat
.iter()
.map(|&(lon, lat)| {
let (mx, my) = facett_map::mercator(lon, lat);
[(mx - origin[0]) as f32, (my - origin[1]) as f32]
})
.collect(),
width_px: w.style.width,
col: {
let [r, g, b, a] = w.style.color.to_array();
[f32::from(r) / 255.0, f32::from(g) / 255.0, f32::from(b) / 255.0, f32::from(a) / 255.0]
}, depth: facett_map::paint_order::depth_of(facett_map::way_rank(w), false), dash: w.style.dashed() })
.collect();
assert!(polys.len() >= 2, "need several polylines for concatenation to mean anything");
let whole_tier = build_line_geometry(&polys);
// The per-way path produces the country tier as each country way's own
// `line_meta` span, concatenated in the declared order. It is no longer a
// CONTIGUOUS prefix of the stream: `line_lod_end` is now a cumulative per-tier
// vertex COUNT (still monotone, still ending at the stream length) because the
// stream is ordered by paint key, not bucketed by LOD tier — the bucketing made a
// `Country` motorway draw UNDER every `City` cul-de-sac.
let per_way: Vec<_> = ways
.iter()
.enumerate()
.filter(|(_, w)| w.style.lod as u32 == 0)
.map(|(i, _)| geo.line_meta[slot[i]])
.collect::<Vec<_>>()
.into_iter()
.flat_map(|m| geo.line_vertices[m.vert_start as usize..(m.vert_start + m.vert_count) as usize].to_vec())
.collect();
assert_eq!(per_way.len(), whole_tier.len(), "per-way expansion emits the same vertex COUNT");
assert_eq!(per_way.as_slice(), whole_tier.as_slice(), "per-way expansion is byte-identical to the whole-tier expansion");
assert!(!whole_tier.is_empty(), "and it actually expanded something");
// `line_lod_end` still tells the truth about tiers, and still ends at the stream.
assert!(geo.line_lod_end.windows(2).all(|w| w[0] <= w[1]), "line_lod_end stays monotone");
assert_eq!(
*geo.line_lod_end.last().unwrap() as usize,
geo.line_vertices.len(),
"line_lod_end still accounts for every vertex"
);
}