facett-map/tests/gpu_lane_share.rs
//! **WHAT SHARE OF THE FRAME DOES THE GPU LANE ACTUALLY TAKE?**
//!
//! The owner, 2026-08-30, after using the client: *"drawing is not fresh… it feels like my
//! first Turbo Pascal 1996 map, it's not smooth double buffered or anything, it kind of
//! fast then slow"* — and then, precisely: *"the lines must be PUT in the GPU and THEN
//! not manipulated from egui."*
//!
//! The evidence offered for "the lane owns nothing" was korp's `FIRST-PAINT` line:
//!
//! ```text
//! FIRST-PAINT lod=5 after 91 ms — … upload 25 us (1280 bytes) · gpu 5589 us
//! ```
//!
//! **1 280 bytes**, set beside a "1 264 ways" from elsewhere in the same log. That
//! reading is wrong twice over, and this file is what proves it.
//!
//! * `LaneClockStats::last_upload_bytes` is the bytes of ONE upload EVENT, and the tile
//! lane calls `note_upload` once per hand-over GROUP — hundreds per viewport. The
//! printed number is whichever group uploaded last; it was never a per-frame figure.
//! MEASURED here: a 640-way scene warms in **480 upload events / 744 960 bytes**, i.e.
//! ~1 552 bytes per event — the same order as the 1 280 that was read as a whole frame.
//! * The two numbers are not even the same moment. `1280 bytes` is `oden-ui-fresh.log`
//! line 33, `lod=5`, 91 ms in; `1264 ways` is line 451, a `lod=15` street-level serve
//! four hundred lines later.
//!
//! Bytes cannot answer "is the lane carrying its weight" at all, because the *right*
//! answer to a pan over ground already in VRAM is **zero bytes**. **Ways on each side**
//! can, and `MapFacet::gpu_lane_share` counts them off the very draw list the frame
//! paints from.
//!
//! # What this file measures, on a real device
//!
//! All figures below are from an RTX 4090 (Vulkan, `DiscreteGpu`) — the adapter each
//! test prints and refuses to substitute a software rasteriser for.
//!
//! 1. `the_lane_takes_the_ways_and_the_cpu_keeps_only_what_has_no_way` — the RATIO, at
//! three cameras, over a korp-shaped scene (1 200 ways, 150 producer-stamped tiles,
//! 8 feature classes — the shape of the Stockholm viewport in the owner's log).
//! MEASURED: **1 200 of 1 200 ways on the lane at all three cameras, 0 on the CPU
//! painter** (RED arm: the adapter-less plan puts 840 on it).
//! 2. `a_pan_over_cached_ground_re_bakes_nothing` — re-bakes counted across a pan, with
//! the RED arm right beside it (perturb one vertex per frame ⇒ the content key moves
//! ⇒ the count climbs, and it must come back down when the perturbation stops).
//! MEASURED: **480 upload events warming, then 0 across 12 panned frames**; RED
//! **5 760**; settled back to **0**.
//! 3. `a_panning_frame_does_not_re_group_the_ways` — the grouped `O(ways)` walk, counted
//! at two different frame counts so a per-frame walk cannot pass. MEASURED: **2 over
//! 7 panned frames and 2 over 23**; RED (a layer toggled every frame) **9 and 25**.
//! 4. `the_group_census_at_a_korp_shaped_viewport` — the number nobody was bounding:
//! groups are `tiles × bands`, and every group is a VRAM entry, a cull dispatch, up
//! to three `draw_indirect`s **and five per-frame uniform writes**. MEASURED: **900
//! groups over 150 tiles (6.00/tile) for 1 200 ways ⇒ ≥5 400 `queue.write_buffer`
//! calls per frame**, panning or still. That is the cost that is left.
//! 5. `a_stamped_frame_does_not_fold_the_whole_lane_origin` — the THIRD `O(ways)` walk,
//! in the caller: `add_gpu_ways_callback` folded `ways_origin` over the whole lane
//! before it asked whether the frame was even taking the grouped path, which discards
//! that origin. MEASURED: **0 folds across 11 panned frames** on korp's stamped shape;
//! RED (unstamped single band, the legacy path that needs it) **13**.
//!
//! **THE ADAPTER IS PRINTED.** A lane figure measured on `llvmpipe` is meaningless and
//! this host can silently select it when the 4090 is busy, so every test here says which
//! adapter it got and refuses to report a lane number from a CPU one.
//!
//! Run (needs the real device, opens NO window):
//! ```text
//! env -u WAYLAND_DISPLAY -u DISPLAY FACETT_TEST_GPU=discrete \
//! cargo test -p facett-map --features wgpu --test gpu_lane_share -- --nocapture --test-threads=1
//! ```
#![cfg(feature = "wgpu")]
use facett_core::render::lane;
use facett_map::layer::{LatLon, MapFacet, MapTheme, MapView, WaysLayer};
use facett_map::{MapPalette, Way};
/// **The device opt-in, shared** — `common/gpu_gate.rs`, lane GG #5.
///
/// This file's own copy was the LAXEST of the three: it took any value at all and filed
/// **no matrix row** for a skip, while the two siblings demanded `discrete|software` and
/// filed a failing row. So one `FACETT_TEST_GPU=integrated` ran this file and skipped
/// those — the same env var meaning two different things three files apart, which is the
/// twin LAW 5 forbids arriving as a *semantic* difference rather than a duplicated body.
#[path = "common/gpu_gate.rs"]
mod gpu_gate;
/// The owner's own window, near enough: the log's viewport reports 104–156 covering tiles
/// on a screen this size.
const CANVAS: egui::Vec2 = egui::vec2(1600.0, 1000.0);
#[path = "common/one_gpu.rs"]
mod one_gpu;
// ── the korp-shaped scene ───────────────────────────────────────────────────
/// The eight feature classes a real Stockholm tile carries, closed-ring or not. Taken
/// from what korp's classifier actually emits, not invented: each one lands in its own
/// `paint_order::way_draw_band`, which is what makes the group census real.
const CLASSES: &[(&[(&str, &str)], bool)] = &[
(&[("natural", "water")], true),
(&[("landuse", "residential")], true),
(&[("leisure", "park")], true),
(&[("building", "yes")], true),
(&[("highway", "motorway")], false),
(&[("highway", "residential")], false),
(&[("highway", "footway")], false),
(&[("waterway", "river")], false),
];
/// **A korp-shaped way set**: `tiles` producer-stamped tiles × [`CLASSES`] × `per_class`
/// ways, laid out over a Stockholm-sized patch. 150 × 8 × 1 = 1 200 ways over 150 tiles
/// is the shape of the viewport in `oden-ui-fresh.log` (1 264 ways, 104–156 tiles).
///
/// `jitter` perturbs every coordinate — the RED arm's lever: a scene built with a
/// different jitter is a different content key for every group.
fn korp_shaped(tiles: u64, per_class: usize, jitter: f64) -> Vec<Way> {
let pal = MapPalette::cartographic(false);
let mut ways = Vec::with_capacity(tiles as usize * CLASSES.len() * per_class);
let mut id = 1i64;
// CLASS-MAJOR, because korp's `push_way_layers_shared` cuts the slab into one
// contiguous range per feature class. Emitting tile-major would make every "class"
// layer a mixture and the fixture would not be the shape under test.
for (tags, closed) in CLASSES {
for t in 1..=tiles {
// A 15×N grid of tiles over central Stockholm.
let (tx, ty) = (((t - 1) % 15) as f64, ((t - 1) / 15) as f64);
let (lon0, lat0) = (18.0 + tx * 0.02, 59.30 + ty * 0.01);
for k in 0..per_class {
let d = 0.0015 * k as f64;
let (lon, lat) = (lon0 + d + jitter, lat0 + d + jitter);
let tagv: Vec<(String, String)> =
tags.iter().map(|(k, v)| ((*k).to_string(), (*v).to_string())).collect();
let coords: Vec<(f64, f64)> = if *closed {
let w = 0.0008;
vec![
(lon - w, lat - w),
(lon + w, lat - w),
(lon + w, lat + w),
(lon - w, lat + w),
(lon - w, lat - w),
]
} else {
vec![(lon, lat), (lon + 0.006, lat + 0.001), (lon + 0.012, lat)]
};
let Some(mut w) = Way::from_lonlat_with(id, &coords, &tagv, &pal) else { continue };
// The stamp that flips facett's hand-over to the content-keyed VRAM
// cache — korp writes it in `classify_rows` from the served row's tile.
w.tile_key = t_key(tx, ty);
ways.push(w);
id += 1;
}
}
}
ways
}
/// A producer's tile stamp. `| 1` keeps it out of the `0` REST bucket; the shift keeps
/// the two axes from folding into each other (they did: `ty | 1` collapsed rows 0/1,
/// 2/3, … and a "150-tile" fixture measured 76 tiles).
fn t_key(tx: f64, ty: f64) -> u64 {
((tx as u64) << 32) | ((ty as u64) << 1) | 1
}
/// The facet korp builds: ONE `WaysLayer::shared` per feature class over a single slab,
/// which is exactly `map_paint::push_way_layers_shared`.
fn korp_facet(view: MapView, ways: Vec<Way>) -> MapFacet {
let mut f = MapFacet::new("lane-share", view, MapTheme::default());
let slab: std::sync::Arc<[Way]> = ways.into();
// korp splits the slab by class into contiguous ranges (`push_way_layers_shared`);
// the fixture is emitted class-major, so stripe `c` IS class `c`.
let n = slab.len();
let per = n.div_ceil(CLASSES.len());
for (c, _) in CLASSES.iter().enumerate() {
let range = (c * per).min(n)..((c + 1) * per).min(n);
if range.is_empty() {
continue;
}
f.push_layer(Box::new(WaysLayer::shared(format!("osm-class-{c}"), std::sync::Arc::clone(&slab), range)));
}
f
}
// ── the device ──────────────────────────────────────────────────────────────
/// A fresh real `RenderState` with `GpuMapRenderer` installed — korp's own eframe
/// creation closure, and deliberately NOT `egui_kittest`'s default setup, whose selector
/// prefers software rasterisers.
fn fresh_render_state() -> egui_wgpu::RenderState {
let setup = facett_core::render::gpu::adapter_wgpu::facett_wgpu_options().wgpu_setup;
let rs = egui_kittest::wgpu::create_render_state(setup, Default::default());
facett_map::gpu::install_renderer(&rs);
rs
}
/// **SAY WHICH ADAPTER, AND REFUSE A CPU ONE.** This box selects `llvmpipe` without
/// erroring when the 4090 is saturated (lane SPIKE, 2026-08-30), and every number in
/// this file is meaningless on a software rasteriser. Returns the adapter's name so the
/// report can carry it.
fn adapter_name_or_bail(rs: &egui_wgpu::RenderState) -> String {
let info = rs.adapter.get_info();
let name = format!("{} ({:?}, {:?})", info.name, info.backend, info.device_type);
eprintln!("LOOK: adapter = {name}");
assert_ne!(
info.device_type,
wgpu::DeviceType::Cpu,
"this host handed the test a SOFTWARE rasteriser ({name}) — every lane figure \
measured on it is meaningless, so the measurement is refused rather than \
reported"
);
name
}
/// Drive a facet through the real device for `steps` frames, running a hook before each
/// one (the pan). Every step is rendered, because the upload rides on the frame whose
/// callback carries the payload — a `run_steps(n); render()` shape submits only the last
/// one and the lane sits legitimately at 0.
fn drive_with(mut f: MapFacet, steps: u32, mut before_each: impl FnMut(u32, &mut MapFacet)) -> u64 {
use egui_kittest::wgpu::WgpuTestRenderer;
let _turn = one_gpu::hold();
let rs = fresh_render_state();
adapter_name_or_bail(&rs);
let before = lane::lane_use().frames();
let mut step = 0u32;
let mut h = egui_kittest::Harness::builder()
.with_size(CANVAS)
.renderer(WgpuTestRenderer::from_render_state(rs))
.build_ui(move |ui| {
before_each(step, &mut f);
step += 1;
f.ui(ui);
});
for _ in 0..steps {
h.render().expect("the wgpu backend produced a frame");
h.step();
}
lane::lane_use().frames() - before
}
fn gated(name: &str) -> bool {
gpu_gate::gated("facett-map", name)
}
/// Cumulative bytes this lane has handed the device — the field that can answer a
/// per-frame question, unlike `last_upload_bytes`.
fn bytes_total() -> u64 {
lane::clock_stats_of(facett_map::gpu::lane_name()).upload_bytes_total
}
fn uploads_total() -> u64 {
lane::clock_stats_of(facett_map::gpu::lane_name()).uploads_total
}
// ── 1. THE RATIO ────────────────────────────────────────────────────────────
/// **THE LANE'S SHARE, AT THREE REAL CAMERAS.**
///
/// The claim under test is not "bytes moved" — it is *who draws the ways*. At every
/// camera the lane must own **every way**, and the CPU painter must be left with only
/// what has no `Way` at all (labels, points) plus a landmass silhouette if one is
/// pushed.
///
/// **Seen RED** by asking the same body for the pure-CPU plan
/// (`features_for_cpu_for_test(&[], false, false)`, which is what a host with no adapter
/// gets): the CPU way count jumps from 0 to all 1 200. A guard that could not move
/// between those two numbers would be comparing a value to a copy of itself.
#[test]
fn the_lane_takes_the_ways_and_the_cpu_keeps_only_what_has_no_way() {
if !gated("the_lane_takes_the_ways_and_the_cpu_keeps_only_what_has_no_way") {
return;
}
let _turn = one_gpu::hold();
let rs = fresh_render_state();
let adapter = adapter_name_or_bail(&rs);
let ways = korp_shaped(150, 1, 0.0);
let total = ways.len();
let cameras = [
("sweden z6", LatLon::new(62.0, 17.0), 6.0),
("stockholm z13", LatLon::new(59.33, 18.06), 13.0),
("street z16", LatLon::new(59.33, 18.06), 16.0),
];
let mut rows = Vec::new();
for (name, at, zoom) in cameras {
let view = MapView::new(at, zoom, (CANVAS.x, CANVAS.y));
let f = korp_facet(view, ways.clone());
let s = f.gpu_lane_share();
let on_gpu = s.lane_lines + s.lane_fills;
let share = on_gpu as f64 / (on_gpu + s.cpu_ways).max(1) as f64;
eprintln!(
"LOOK: {name:<14} lane={:?} owns={:?} floor={:?} — GPU {on_gpu} ways \
({} line + {} fill) · CPU {} ways of {} features · {} groups / {} tiles · share {:.1}%",
s.lane, s.owns, s.floor, s.lane_lines, s.lane_fills, s.cpu_ways, s.cpu_feats, s.groups, s.tiles,
share * 100.0
);
assert!(s.lane.is_gpu(), "{name}: the plan declined the lane — {:?}", s.lane);
assert!(s.owns.lines && s.owns.fills, "{name}: the lane claimed only {:?}", s.owns);
assert_eq!(on_gpu, total, "{name}: the lane must bake every way of the scene");
assert_eq!(s.cpu_ways, 0, "{name}: {} ways were left on the CPU painter", s.cpu_ways);
rows.push((name, share));
}
// ── THE RED ARM, in the same body: the pure-CPU plan every adapter-less host
// takes. If this does not move, the assertions above are watching nothing.
let view = MapView::new(LatLon::new(59.33, 18.06), 13.0, (CANVAS.x, CANVAS.y));
let f = korp_facet(view, ways);
let cpu_only = f.features_for_cpu_for_test(&[], false, false);
let cpu_only_ways = cpu_only
.iter()
.filter(|d| {
matches!(d.geom, facett_map::layer::DrawGeom::Line(_) | facett_map::layer::DrawGeom::Polygon { .. })
})
.count();
// The CPU painter culls to the viewport and gates on LOD before it materialises a
// feature; the GPU bake does neither (it bakes the layer's whole slice and culls on
// the device). So the honest comparison is against the ways the CPU would DRAW, not
// against the scene size — and that number is what the lane took to zero.
eprintln!(
"LOOK: RED arm — the pure-CPU plan puts {cpu_only_ways} ways on the painter \
(of {total} in the scene); with the lane it is 0"
);
assert!(
cpu_only_ways > 0,
"the CPU-only arm drew no way at all, so `cpu_ways == 0` above proves nothing"
);
nornir_testmatrix::functional_status(
"facett-map",
"gpu_lane_way_share",
rows.iter().all(|(_, s)| *s > 0.99),
&format!(
"{total} ways, 3 cameras, all 100% on the lane (CPU-only arm draws \
{cpu_only_ways}) — {adapter}"
),
);
}
// ── 2. RE-BAKES PER PAN ─────────────────────────────────────────────────────
/// **A PAN OVER GROUND ALREADY IN VRAM MUST RE-BAKE NOTHING.**
///
/// This is the half of "fast then slow" the complaint points at, and the answer here is
/// a measured number rather than a reading of the code: 24 frames of real panning across
/// a korp-shaped scene, counting upload EVENTS and BYTES on each side of the pan.
///
/// The tile lane's content key is a function of the group's ways and the group's OWN
/// bake origin (`ways_origin` over that group) — **the camera is not in it**. So a pan
/// cannot invalidate a bake, and the number must be 0.
///
/// **THE RED ARM IS IN THIS TEST.** A second run rebuilds the scene with a fresh jitter
/// every frame, which is exactly "force the signature to change every frame": the count
/// must climb, and then a third run with the jitter switched back off must return to 0.
/// Without that third leg a climbing count could be the harness warming up.
#[test]
fn a_pan_over_cached_ground_re_bakes_nothing() {
if !gated("a_pan_over_cached_ground_re_bakes_nothing") {
return;
}
const FRAMES: u32 = 24;
let base = korp_shaped(80, 1, 0.0);
let view = || MapView::new(LatLon::new(59.33, 18.06), 13.0, (CANVAS.x, CANVAS.y));
// ── GREEN ARM: one scene, panned. The first frames bake it; the rest must not.
let warm_uploads;
let pan_uploads;
let pan_bytes;
{
let ways = base.clone();
let f = korp_facet(view(), ways);
let split = std::sync::Arc::new(std::sync::atomic::AtomicU64::new(0));
let (u0, b0) = (uploads_total(), bytes_total());
let sp = std::sync::Arc::clone(&split);
let frames = drive_with(f, FRAMES, move |step, f| {
if step == FRAMES / 2 {
// Bank the counters at the halfway mark: everything after this is pure
// pan over ground the device already holds.
sp.store(uploads_total(), std::sync::atomic::Ordering::Relaxed);
}
if step >= FRAMES / 2 {
// A REAL pan, through the same `pan_by_px` a drag on the canvas drives.
f.view_mut().pan_by_px(-40.0, -25.0);
}
});
assert!(frames > 0, "the lane encoded no frame — nothing below is about the GPU lane");
let mid = split.load(std::sync::atomic::Ordering::Relaxed);
warm_uploads = mid - u0;
pan_uploads = uploads_total() - mid;
pan_bytes = bytes_total() - b0;
eprintln!(
"LOOK: pan — {warm_uploads} upload events warming the scene, then \
{pan_uploads} across {} panned frames ({} total bytes over the whole run)",
FRAMES / 2,
pan_bytes
);
}
assert!(warm_uploads > 0, "the scene never uploaded at all — the measurement has no subject");
assert_eq!(pan_uploads, 0, "a pan over cached ground re-baked {pan_uploads} groups");
// ── RED ARM: the same pan with the geometry moving under it every frame.
let churn_uploads = {
let u0 = uploads_total();
let f = korp_facet(view(), korp_shaped(80, 1, 0.0));
drive_with(f, FRAMES / 2, move |step, f| {
// Rebuild the way set with a fresh jitter: every group's content key moves,
// which is what "the signature changed" means to the tile lane.
let ways = korp_shaped(80, 1, 1e-5 * f64::from(step + 1));
let mut fresh = korp_facet(*f.view(), ways);
std::mem::swap(f, &mut fresh);
});
uploads_total() - u0
};
eprintln!("LOOK: RED arm — geometry moving every frame re-baked {churn_uploads} groups");
assert!(
churn_uploads > pan_uploads,
"forcing the content key to change every frame did NOT raise the re-bake count \
({churn_uploads} vs {pan_uploads}) — the green arm above is watching nothing"
);
// ── AND BACK: the number must return once the churn stops.
let settled = {
let f = korp_facet(view(), base);
let u_mid = std::sync::Arc::new(std::sync::atomic::AtomicU64::new(0));
let um = std::sync::Arc::clone(&u_mid);
drive_with(f, FRAMES, move |step, f| {
if step == FRAMES / 2 {
um.store(uploads_total(), std::sync::atomic::Ordering::Relaxed);
}
if step >= FRAMES / 2 {
f.view_mut().pan_by_px(-40.0, -25.0);
}
});
uploads_total() - u_mid.load(std::sync::atomic::Ordering::Relaxed)
};
eprintln!("LOOK: settled — {settled} re-bakes across {} panned frames after the churn", FRAMES / 2);
assert_eq!(settled, 0, "the re-bake count did not come back down after the churn stopped");
nornir_testmatrix::functional_status(
"facett-map",
"gpu_lane_rebakes_per_pan",
pan_uploads == 0 && churn_uploads > 0 && settled == 0,
&format!("pan {pan_uploads} re-bakes, forced-churn {churn_uploads}, settled {settled}"),
);
}
// ── 3. THE PER-FRAME WORK ───────────────────────────────────────────────────
/// **A STILL OR PANNING FRAME MUST NOT RE-GROUP THE WAYS.**
///
/// The owner: *"the lines must be PUT in the GPU and THEN not manipulated from egui."*
/// The payload already obeyed that — a pan over cached ground uploads nothing (test 2).
/// The GROUPING around it did not: `add_gpu_tiled_callback` walked every way of every
/// taken layer twice into a fresh `HashMap`, collected the keys and sorted them, on
/// **every frame**, only to hand over `geometry: None` for all of them.
///
/// Nothing in that grouping depends on the camera. The count below is the applied fact —
/// grouped `O(ways)` walks over N frames of real panning — and both arms are here:
///
/// * **GREEN**: one apply, panned. The walk runs once.
/// * **RED**: the layers replaced every frame (a fresh `geometry_generation` each time,
/// which is what korp does on a real apply). The walk must run every frame, or the
/// green arm is a guard that cannot go red.
#[test]
fn a_panning_frame_does_not_re_group_the_ways() {
if !gated("a_panning_frame_does_not_re_group_the_ways") {
return;
}
let base = korp_shaped(80, 1, 0.0);
let view = || MapView::new(LatLon::new(59.33, 18.06), 13.0, (CANVAS.x, CANVAS.y));
// The count as of the START of the last frame, so it covers `frames - 1` rendered
// frames — the same arithmetic in both arms, which is what makes them comparable.
let walks = |frames: u32, toggle: bool, trace: bool| -> usize {
let seen = std::sync::Arc::new(std::sync::atomic::AtomicUsize::new(0));
let log = std::sync::Arc::new(std::sync::Mutex::new(Vec::<(u32, usize, u64)>::new()));
let (sv, lg) = (std::sync::Arc::clone(&seen), std::sync::Arc::clone(&log));
let f = korp_facet(view(), base.clone());
drive_with(f, frames, move |step, f| {
sv.store(f.tile_group_walk_count(), std::sync::atomic::Ordering::Relaxed);
if trace {
lg.lock().unwrap().push((step, f.tile_group_walk_count(), uploads_total()));
}
if toggle {
// The SAME facet, its layer set changed every frame — a checkbox
// toggled, which is what korp's layer panel does and what the memo key
// is keyed on. The plan cannot be reused, so the walk must run.
f.set_layer_enabled("osm-class-7", step % 2 == 0);
}
if step > 0 {
f.view_mut().pan_by_px(-40.0, -25.0);
}
});
if trace {
// ONE lock. `&log.lock().unwrap()[..6.min(log.lock().unwrap().len())]` is a
// self-DEADLOCK — `std::sync::Mutex` is not reentrant, so the inner `.len()`
// lock blocks forever on the guard the outer index expression still holds.
// It hung this test on the real device with the adapter line already
// printed, which reads exactly like a wedged GPU and is not one.
let head: Vec<(u32, usize, u64)> = {
let t = log.lock().unwrap();
t[..6.min(t.len())].to_vec()
};
eprintln!("LOOK: per-frame (step, walks, uploads) = {head:?}");
}
seen.load(std::sync::atomic::Ordering::Relaxed)
};
// ── GREEN: the same apply, panned for 8 frames and for 24. If the walk were still
// per-frame these two could not agree; that they DO is the whole claim, and it
// is a claim no single number can make.
let (g8, g24) = (walks(8, false, true), walks(24, false, false));
eprintln!("LOOK: grouped way walks — {g8} over 7 panned frames, {g24} over 23, on one apply");
// ── RED: the layer set changed every frame. The plan key moves, so the walk must
// run again — and the count must GROW with the frame count, which is exactly
// what the green arm says it no longer does.
let (r8, r24) = (walks(8, true, false), walks(24, true, false));
eprintln!("LOOK: RED arm — a layer toggled every frame walks {r8} times over 7 frames, {r24} over 23");
assert_eq!(
g8, g24,
"the grouped way walk still tracks the FRAME count on one apply ({g8} over 7 frames, \
{g24} over 23) — the plan memo is not holding"
);
assert!(
g24 <= 2,
"one apply cost {g24} grouped way walks; at most two are explainable (the plan \
build, and the one later frame that still owed the device bytes)"
);
assert!(
r24 > r8 && r24 > g24,
"changing the layer set every frame did NOT make the grouped-walk count grow with \
the frames ({r8} over 7, {r24} over 23, against {g24} green) — the green arm above \
is a guard that cannot go red"
);
nornir_testmatrix::functional_status(
"facett-map",
"gpu_lane_frame_regrouping",
g8 == g24 && g24 <= 2 && r24 > r8,
&format!("{g8} walks over 7 panned frames and {g24} over 23 on one apply (RED: {r8} / {r24})"),
);
}
// ── 4. THE GROUP CENSUS ─────────────────────────────────────────────────────
/// **HOW MANY HAND-OVER GROUPS A REAL VIEWPORT PRODUCES, AND WHAT EACH ONE COSTS.**
///
/// `gpu_lane_paint_order::the_band_major_group_census_stays_bounded` bounds the RATIO
/// (groups per tile ≤ classes). Nothing bounded the ABSOLUTE number, and the absolute
/// number is what a frame pays: `MapPaintCallback::prepare` writes **five** uniform
/// buffers per group per frame (draw, fill, and the pretty-line casing/fill/outline
/// trio) and `run_cull_tiles_clocked` then writes one indirect reset per chunk plus one
/// viewport per group — every frame, panning or still.
///
/// This test does not fail a threshold it invented; it PRINTS the number for a
/// korp-shaped viewport so the cost is on the record, and asserts only the arithmetic
/// that must hold (`groups == tiles × distinct bands`), which is what makes the printed
/// number trustworthy.
#[test]
fn the_group_census_at_a_korp_shaped_viewport() {
if !gated("the_group_census_at_a_korp_shaped_viewport") {
return;
}
let _turn = one_gpu::hold();
let rs = fresh_render_state();
let adapter = adapter_name_or_bail(&rs);
let ways = korp_shaped(150, 1, 0.0);
let f = korp_facet(MapView::new(LatLon::new(59.33, 18.06), 13.0, (CANVAS.x, CANVAS.y)), ways);
let s = f.gpu_lane_share();
let per_tile = s.groups as f64 / s.tiles.max(1) as f64;
// Five uniform writes per group in `prepare`, plus one viewport write per group in
// the cull — the floor of the per-frame `queue.write_buffer` traffic, before the
// per-chunk indirect resets.
let writes = s.groups * 6;
eprintln!(
"LOOK: census — {} groups over {} tiles ({per_tile:.2}/tile, {} classes) for \
{} ways ⇒ ≥{writes} queue.write_buffer calls PER FRAME · {adapter}",
s.groups,
s.tiles,
CLASSES.len(),
s.lane_lines + s.lane_fills
);
assert_eq!(s.tiles, 150, "the fixture must really span 150 producer-stamped tiles");
assert!(per_tile > 1.0, "the band is not in the group key and this census measures the wrong thing");
assert!(
per_tile <= CLASSES.len() as f64,
"{per_tile:.2} groups per tile exceeds the {} classes in the fixture — something \
finer than the draw band is splitting groups",
CLASSES.len()
);
nornir_testmatrix::functional_status(
"facett-map",
"gpu_lane_group_census",
per_tile > 1.0 && per_tile <= CLASSES.len() as f64,
&format!("{} groups / {} tiles = {per_tile:.2} per tile, ≥{writes} writes/frame", s.groups, s.tiles),
);
}
// ── 5. THE THIRD WALK ───────────────────────────────────────────────────────
/// **A single class, unstamped** — the LEGACY whole-payload arm of
/// [`MapFacet::add_gpu_ways_callback`]: one `way_draw_band`, `tile_key == 0` on every
/// way, so `gpu_lane_needs_grouping` is false and the frame takes the un-grouped path
/// that genuinely needs a whole-lane `ways_origin`.
fn one_band_unstamped(n: usize) -> Vec<Way> {
let pal = MapPalette::cartographic(false);
let mut ways = Vec::with_capacity(n);
for k in 0..n {
let d = 0.0004 * k as f64;
let (lon, lat) = (18.0 + d, 59.30 + d);
let coords = [(lon, lat), (lon + 0.006, lat + 0.001), (lon + 0.012, lat)];
let tags = [("highway", "residential")];
if let Some(w) = Way::from_lonlat_with(k as i64 + 1, &coords, &tags, &pal) {
// NO `tile_key` — this is the unstamped bucket the legacy path serves.
ways.push(w);
}
}
ways
}
/// A facet with ONE ways layer — one class, so one draw band.
fn one_layer_facet(view: MapView, ways: Vec<Way>) -> MapFacet {
let mut f = MapFacet::new("lane-share-legacy", view, MapTheme::default());
let slab: std::sync::Arc<[Way]> = ways.into();
let n = slab.len();
f.push_layer(Box::new(WaysLayer::shared("osm-one-band".to_string(), slab, 0..n)));
f
}
/// **THE THIRD `O(ways)` WALK A FRAME PAID, AND WHY korp PAID IT FOR NOTHING.**
///
/// Two walks were named when this lane started — the plan build and the bake gather —
/// and they are now one (`a_panning_frame_does_not_re_group_the_ways`). There was a
/// THIRD, in the caller: `add_gpu_ways_callback` ran `crate::gpu::ways_origin` over
/// every way of every taken layer **before** it asked whether the frame was even going
/// to take the grouped path — and the grouped path throws that origin away, because the
/// tile lane bakes every group against the group's OWN origin (`TileGroupPlan::origin`).
///
/// korp is always the grouped path: every way it serves carries a producer stamp. So on
/// the machine the owner is complaining about, that fold ran over the whole way set, sixty
/// times a second, and its result was dropped on the next line. `gpu_lane_needs_grouping`
/// short-circuits on the first stamped way, so asking it first costs `O(1)` there.
///
/// This is the applied fact, not a reading of the code:
///
/// * **GREEN** — a korp-shaped (stamped, many-band) scene, panned: **0** whole-lane
/// origin folds, however many frames are drawn.
/// * **RED** — one class, unstamped: the legacy path, which genuinely needs the origin,
/// still folds once per frame. Without this arm the green arm would be indistinguishable
/// from a counter nobody increments.
#[test]
fn a_stamped_frame_does_not_fold_the_whole_lane_origin() {
if !gated("a_stamped_frame_does_not_fold_the_whole_lane_origin") {
return;
}
const FRAMES: u32 = 12;
let view = || MapView::new(LatLon::new(59.33, 18.06), 13.0, (CANVAS.x, CANVAS.y));
// ── GREEN: korp's own shape — stamped ways, the grouped path.
let (grouped_walks, grouped_frames) = {
let seen = std::sync::Arc::new(std::sync::atomic::AtomicUsize::new(0));
let sv = std::sync::Arc::clone(&seen);
let f = korp_facet(view(), korp_shaped(40, 1, 0.0));
let frames = drive_with(f, FRAMES, move |_, f| {
sv.store(f.lane_origin_walk_count(), std::sync::atomic::Ordering::Relaxed);
f.view_mut().pan_by_px(-40.0, -25.0);
});
(seen.load(std::sync::atomic::Ordering::Relaxed), frames)
};
eprintln!(
"LOOK: stamped/grouped — {grouped_walks} whole-lane origin folds across {} frames \
({grouped_frames} lane frames encoded)",
FRAMES - 1
);
// ── RED: one band, no stamps — the legacy path, which needs the origin every frame.
let legacy_walks = {
let seen = std::sync::Arc::new(std::sync::atomic::AtomicUsize::new(0));
let sv = std::sync::Arc::clone(&seen);
let f = one_layer_facet(view(), one_band_unstamped(600));
drive_with(f, FRAMES, move |_, f| {
sv.store(f.lane_origin_walk_count(), std::sync::atomic::Ordering::Relaxed);
f.view_mut().pan_by_px(-40.0, -25.0);
});
seen.load(std::sync::atomic::Ordering::Relaxed)
};
eprintln!("LOOK: RED arm — unstamped single band folds the origin {legacy_walks} times over the same frames");
assert!(grouped_frames > 0, "the lane encoded no frame — nothing here is about the GPU lane");
assert_eq!(
grouped_walks, 0,
"a stamped frame folded the whole-lane origin {grouped_walks} times for a value the \
grouped path discards"
);
assert!(
legacy_walks >= usize::try_from(FRAMES - 1).expect("frame count fits"),
"the legacy arm folded the origin only {legacy_walks} times over {} frames — it is \
not exercising the path that needs it, so the 0 above proves nothing",
FRAMES - 1
);
nornir_testmatrix::functional_status(
"facett-map",
"gpu_lane_origin_fold",
grouped_walks == 0 && legacy_walks > 0,
&format!("stamped {grouped_walks} folds, unstamped-legacy {legacy_walks} over {} frames", FRAMES - 1),
);
}