facett-pipeline/tests/flowsim_gpu_lane.rs
//! **THE GPU ADVECT PATH, DRIVEN BY A REAL HOST — proven on real pixels.**
//! (`--features flowsim-gpu`)
//!
//! `GFX_V2` / `flowsim-spike.md` §9.6 recorded the defect this file exists to close:
//!
//! > *"The GPU advect path is never driven by a real host. `FlowGpu` compiles and
//! > `GpuParticles::advect` is proven correct against `advect_cpu` on a real device,
//! > but no host wires it into an `egui_wgpu` paint callback. Every rendered frame is
//! > the CPU painter path."*
//!
//! A pixel diff alone cannot close it, because **the CPU painter draws a perfectly good
//! overlay too** — any "the particles reached the framebuffer" assertion is green
//! whichever lane painted them, which is exactly how the gap survived. So every claim
//! here is a PAIR: a pixel measurement (did real content land?) alongside lane
//! attribution (`flow_frames_drawn()`, the count of advect+draw frames the device
//! actually encoded). Neither half is sufficient:
//!
//! | claim | the lie it catches |
//! |---|---|
//! | GPU lane on: overlay pixels **and** the frame counter moved | the counter moved but nothing drew / something drew but the CPU drew it |
//! | CPU lane on: overlay pixels **and** the counter did NOT move | the "GPU" lane was silently the CPU painter all along |
//! | two clocks on ONE renderer differ by a third of the footprint | `advect` is never dispatched and the dots are frozen on the device |
//! | a dead DAG: 0 overlay px **and** 0 counter movement | a stopped pipeline glows, or claims the lane |
//! | every CPU dot has a GPU dot within 4 px, and the extents match | **the dots are painted, animated, GPU-drawn — and 32 px off the edges** |
//!
//! That last row is not hypothetical: it is what the first version of this lane did, and
//! the first version of this file was green for it. See `FlowFrame`'s note.
//!
//! Run: `cargo test -p facett-pipeline --features flowsim-gpu --test flowsim_gpu_lane`
//! Skips loudly with no wgpu adapter; never a false pass.
#![cfg(feature = "flowsim-gpu")]
use std::sync::atomic::{AtomicU64, Ordering};
use std::sync::Arc;
use egui_kittest::wgpu::{create_render_state, default_wgpu_setup, WgpuTestRenderer};
use egui_kittest::Harness;
use facett_core::Facet;
use facett_flowsim::gpu::host::{flow_frames_drawn, flow_renderer_installed, install_flow_renderer};
use facett_flowsim::{FlowMode, FlowMsg};
use facett_pipeline::{PipelineNode, PipelineView, Status};
const EDGES: [(usize, usize); 3] = [(1, 0), (2, 1), (3, 2)];
/// The same 4-stage ETL fixture `robot_flowsim.rs` uses: `live` decides only whether
/// the stages report a throughput, so every drawn label is identical across a
/// live/dead pair and a diff is the overlay alone.
fn dag(live: bool) -> Vec<PipelineNode> {
let tp = |n: &str, st: Status, rows: u64, r: f32| {
let node = PipelineNode::new(n, st).with_rows(rows);
if live {
node.with_throughput_rows_s(r)
} else {
node
}
};
vec![
tp("raw", Status::Done, 120_000, 500.0),
tp("clean", Status::Running, 48_000, 9_000.0),
tp("agg", Status::Running, 12_000, 2_000.0),
tp("out", Status::Running, 4_000, 300.0),
]
}
fn view(live: bool, mode: FlowMode, t: f64, gpu: bool) -> PipelineView {
let mut v = PipelineView::from_nodes("ETL", dag(live), EDGES.to_vec())
// Kill the marching-ants + glow pulse: the overlay is then the only motion.
.with_animation(false)
.with_clock(t)
.with_flowsim(true);
let fs = v.flowsim_mut().unwrap();
fs.update(FlowMsg::SetMode(mode));
fs.set_gpu_active(gpu);
v
}
/// Pixels differing by more than a hair of dithering noise (same measure as
/// `robot_flowsim.rs`, so the two files' numbers are comparable).
fn diff_px(a: &image::RgbaImage, b: &image::RgbaImage) -> usize {
assert_eq!(a.dimensions(), b.dimensions());
a.pixels()
.zip(b.pixels())
.filter(|(p, q)| (0..3).map(|i| (p.0[i] as i32 - q.0[i] as i32).abs()).sum::<i32>() > 12)
.count()
}
/// A harness on its **own fresh** `RenderState` (egui_kittest rejects a reused one), with
/// the flow renderer installed into it, reading its clock from the returned cell.
///
/// The cell is what makes an on-device `advect` step observable: a single harness can be
/// rendered twice, and because both renders go through the SAME `RenderState` they share
/// one persistent `FlowRenderer` — so the second frame advects the buffers the first one
/// seeded instead of re-uploading them. Comparisons that only need one frame each use
/// separate harnesses, exactly as `robot_flowsim.rs` does.
fn harness(live: bool, mode: FlowMode, gpu: bool, t0: f64) -> (Harness<'static>, Arc<AtomicU64>) {
let rs = create_render_state(default_wgpu_setup(), Default::default());
install_flow_renderer(&rs);
let clock = Arc::new(AtomicU64::new(t0.to_bits()));
let c = clock.clone();
let h = Harness::builder()
.with_size(egui::vec2(900.0, 620.0))
.renderer(WgpuTestRenderer::from_render_state(rs))
.build_ui(move |ui| {
facett_core::set_theme(ui.ctx(), facett_core::Theme::sci_fi());
let t = f64::from_bits(c.load(Ordering::Relaxed));
view(live, mode, t, gpu).ui(ui);
});
(h, clock)
}
/// One frame of `(live, mode, gpu)` at clock `t`, or `None` with a loud note when there
/// is no adapter.
fn frame(live: bool, mode: FlowMode, gpu: bool, t: f64) -> Option<image::RgbaImage> {
let (mut h, _c) = harness(live, mode, gpu, t);
h.run_steps(3);
match h.render() {
Ok(img) => Some(img),
Err(e) => {
eprintln!("[flowsim-gpu] no wgpu backend ({e}) — GPU lane pixel proof skipped");
None
}
}
}
/// Where a frame's overlay actually IS: the centroid and bounding box of the pixels it
/// added over the same frame with the overlay Off. The applied-output answer to "did the
/// dots land on the edges", which no "is it non-blank" measure can give.
struct Region {
cx: f64,
cy: f64,
n: usize,
/// `(min_x, max_x, min_y, max_y)`.
bbox: (u32, u32, u32, u32),
/// Per-pixel "the overlay lit this", for [`covered_fraction`].
mask: Vec<bool>,
w: u32,
h: u32,
}
fn overlay_region(with: &image::RgbaImage, without: &image::RgbaImage) -> Region {
assert_eq!(with.dimensions(), without.dimensions());
let (w, h) = with.dimensions();
let mut mask = vec![false; (w * h) as usize];
let (mut sx, mut sy, mut n) = (0f64, 0f64, 0usize);
let (mut x0, mut x1, mut y0, mut y1) = (u32::MAX, 0u32, u32::MAX, 0u32);
for (x, y, p) in with.enumerate_pixels() {
let q = without.get_pixel(x, y);
if (0..3).map(|i| (p.0[i] as i32 - q.0[i] as i32).abs()).sum::<i32>() > 12 {
mask[(y * w + x) as usize] = true;
sx += f64::from(x);
sy += f64::from(y);
n += 1;
x0 = x0.min(x);
x1 = x1.max(x);
y0 = y0.min(y);
y1 = y1.max(y);
}
}
assert!(n > 0, "the overlay lit no pixels at all — nothing to locate");
Region { cx: sx / n as f64, cy: sy / n as f64, n, bbox: (x0, x1, y0, y1), mask, w, h }
}
/// Fraction of `a`'s overlay pixels that have one of `b`'s within `r` px — "is `a`'s
/// overlay drawn where `b`'s is".
///
/// This, not the centroid, is the alignment measure. Two lanes drawing the same dots with
/// different sprite sizes and intensities weight their pixel mass differently along an
/// edge (measured: a 22 px x-centroid gap with the bounding boxes 3 px apart), so a
/// centroid test would be both a false red here and blind to a shift *along* the edges.
/// Coverage tolerates the size difference and collapses to ~0 for any real displacement.
fn covered_fraction(a: &Region, b: &Region, r: i64) -> f64 {
assert_eq!((a.w, a.h), (b.w, b.h));
let mut hit = 0usize;
for y in 0..a.h as i64 {
for x in 0..a.w as i64 {
if !a.mask[(y * a.w as i64 + x) as usize] {
continue;
}
let near = (-r..=r).any(|dy| {
(-r..=r).any(|dx| {
let (nx, ny) = (x + dx, y + dy);
nx >= 0
&& ny >= 0
&& nx < a.w as i64
&& ny < a.h as i64
&& b.mask[(ny * b.w as i64 + nx) as usize]
})
});
if near {
hit += 1;
}
}
}
hit as f64 / a.n as f64
}
/// How many particles the plan seeds — the DERIVED floor for "the overlay reached the
/// framebuffer". A magic constant would drift silently as the plan changes.
fn planned_particles() -> usize {
view(true, FlowMode::Particles, 0.0, false).state_json()["flowsim"]["total_particles"]
.as_u64()
.unwrap() as usize
}
/// **THE PROOF.** Every paired claim in one wgpu session, against one device.
#[test]
fn the_gpu_lane_paints_the_overlay_and_the_cpu_painter_stands_down() {
let n = planned_particles();
assert!(n > 0, "the fixture must plan particles at all");
let floor = n * 15; // each dot is a several-px sprite; see robot_flowsim.rs
let total = 900 * 620;
// ── The baseline: the overlay switched Off. Both lanes decline, so this is the
// bare DAG (base curve + arrowheads + node boxes + canvas). ────────────────
let Some(off) = frame(true, FlowMode::Off, true, 0.0) else { return };
assert!(flow_renderer_installed(), "installing must arm the gate that unlocks the lane");
// ── 1. THE GPU LANE. Pixels AND attribution, in one render. ─────────────────
let before_gpu = flow_frames_drawn();
let gpu_t0 = frame(true, FlowMode::Particles, true, 0.0).expect("gpu frame");
let gpu_frames = flow_frames_drawn() - before_gpu;
let gpu_px = diff_px(&gpu_t0, &off);
eprintln!(
"[flowsim-gpu] GPU lane: {gpu_px}/{total} px differ from the Off frame \
({n} particles planned, derived floor {floor}); device encoded {gpu_frames} advect+draw frames"
);
assert!(
gpu_px > floor,
"the GPU overlay must visibly reach the framebuffer — only {gpu_px} px differ from the \
Off frame, for {n} planned particles"
);
assert!(
gpu_frames > 0,
"ATTRIBUTION: the pixels above must have come from the DEVICE. The advect+draw counter \
did not move ({gpu_frames}), so the CPU painter drew them and the GPU lane is still the \
dead code GFX_V2 §9.6 described."
);
// ── 1b. THE SAME STORY THROUGH THE SHARED REGISTRY. ─────────────────────────
// `flow_frames_drawn()` is this lane's private counter. `facett_core::render::lane`
// is the one an **About box row** is rendered from ("GPU lane: IN USE — N frame(s)"),
// and that row is only honest if it moves for the same real frames this test just
// measured in pixels. Asserting it HERE is what makes the row a measurement rather
// than a claim: this is the only place in the workspace where a GPU lane, real
// pixels and the registry are all in one process at once.
let shared = facett_core::render::lane::lane_use();
assert!(
matches!(shared, facett_core::render::lane::GpuLaneUse::InUse { frames } if frames >= gpu_frames),
"the SHARED lane registry must report IN USE with at least the {gpu_frames} frame(s) the \
device encoded, or an About box reading it understates a lane that is plainly painting: \
got {shared:?}"
);
assert!(
facett_core::render::lane::lanes()
.iter()
.any(|l| l.short_name() == "FlowRenderer" && l.frames > 0),
"the registry must NAME the lane that painted (FlowRenderer), so the row can say who: {:?}",
facett_core::render::lane::lanes()
);
// ── 2. THE CPU LANE, unchanged. The pairing that makes claim 1 mean something:
// the CPU painter still draws a full overlay, and does it WITHOUT the device.
// Without this half, claim 1 is satisfied by a lane that broke the fallback.
let before_cpu = flow_frames_drawn();
let cpu_t0 = frame(true, FlowMode::Particles, false, 0.0).expect("cpu frame");
let cpu_frames = flow_frames_drawn() - before_cpu;
let cpu_px = diff_px(&cpu_t0, &off);
eprintln!("[flowsim-gpu] CPU lane: {cpu_px}/{total} px differ from Off; device frames {cpu_frames}");
assert!(
cpu_px > floor,
"the CPU fallback is the golden path and must still paint the overlay ({cpu_px} px)"
);
assert_eq!(
cpu_frames, 0,
"a host that did NOT opt in must not reach the device — the counter moved by {cpu_frames}"
);
// …and the shared registry must not have counted that CPU frame either, or an
// About box would claim the GPU lane painted a frame the device never saw.
let cpu_shared = facett_core::render::lane::lanes()
.iter()
.filter(|l| l.short_name() == "FlowRenderer")
.map(|l| l.frames)
.sum::<u64>();
assert_eq!(
cpu_shared,
flow_frames_drawn(),
"the shared registry's count must track the device counter exactly — a drift here is an \
About row inventing frames"
);
// ── 3. THE ALIGNMENT CLAIM — the GPU dots land ON the edges. ────────────────
//
// This is the guard the first version of this file did not have, and it is the one
// that matters most. Claims 1, 2 and 4 were ALL green while the GPU overlay sat
// 32.6 px below the edges it was supposed to ride: an offset overlay is present,
// is animated, and differs from the CPU frame. The lane was wired to the wrong
// space (screen instead of the paint callback's viewport) and only the PNG showed
// it. So the claim has to be about WHERE, measured against the lane that is
// already known correct.
//
// Both lanes are asked for the same frame, so every CPU dot must have a GPU dot
// essentially on top of it — and vice versa.
let g = overlay_region(&gpu_t0, &off);
let c = overlay_region(&cpu_t0, &off);
eprintln!(
"[flowsim-gpu] GPU overlay: {} px, centroid ({:.1}, {:.1}), bbox {:?}",
g.n, g.cx, g.cy, g.bbox
);
eprintln!(
"[flowsim-gpu] CPU overlay: {} px, centroid ({:.1}, {:.1}), bbox {:?}",
c.n, c.cx, c.cy, c.bbox
);
let cpu_covered = covered_fraction(&c, &g, 4);
let gpu_covered = covered_fraction(&g, &c, 4);
eprintln!(
"[flowsim-gpu] alignment: {:.1}% of CPU dots have a GPU dot within 4 px; \
{:.1}% the other way",
100.0 * cpu_covered,
100.0 * gpu_covered
);
assert!(
cpu_covered > 0.90,
"the GPU overlay must ride the SAME edges as the CPU overlay — only {:.1}% of the CPU \
dots have a GPU dot within 4 px. That is the 32 px space bug: an overlay that is \
present, animated and GPU-drawn, and still in the wrong place.",
100.0 * cpu_covered
);
assert!(
gpu_covered > 0.80,
"…and the GPU overlay must not be drawing anywhere the CPU overlay is not — only \
{:.1}% of GPU dots sit within 4 px of a CPU dot, so the lane is painting somewhere \
of its own.",
100.0 * gpu_covered
);
// The extent too, so a lane scaled about a coincident centre cannot pass.
let (gw, gh) = (g.bbox.1 - g.bbox.0, g.bbox.3 - g.bbox.2);
let (cw, ch) = (c.bbox.1 - c.bbox.0, c.bbox.3 - c.bbox.2);
eprintln!("[flowsim-gpu] lane overlay extent: GPU {gw}x{gh} px · CPU {cw}x{ch} px");
assert!(
(gw as i64 - cw as i64).abs() < 12 && (gh as i64 - ch as i64).abs() < 12,
"the GPU overlay must span the same edges as the CPU one — GPU {gw}x{gh} vs CPU \
{cw}x{ch}. A rescaled overlay is the same space bug seen through its extent."
);
// ── 4. `advect` RUNS ON THE DEVICE: one renderer, two clocks. ───────────────
// The clock is deliberately NOT in the reseed key, so the second frame advects the
// buffers the first one seeded rather than re-uploading them. If nothing dispatched
// `cs_advect` the dots would be frozen and this is exactly 0.
let (mut h_move, clock) = harness(true, FlowMode::Particles, true, 0.0);
h_move.run_steps(3);
let a = h_move.render().expect("advect frame t0");
clock.store(0.42f64.to_bits(), Ordering::Relaxed);
h_move.run_steps(3);
let b = h_move.render().expect("advect frame t1");
let motion_px = diff_px(&a, &b);
eprintln!(
"[flowsim-gpu] advect on device: {motion_px} px moved between t=0.0 and t=0.42 \
({}% of the {gpu_px} px overlay footprint)",
100 * motion_px / gpu_px.max(1)
);
assert!(
motion_px * 3 > gpu_px,
"the device must ADVECT the particles between two clocks — only {motion_px} px moved \
against a {gpu_px} px overlay footprint. Frozen dots mean `advect` is never dispatched \
and the GPU lane is a static upload."
);
// ── 5. THE ANCHOR: a dead DAG neither glows nor claims the lane. ────────────
// Without this, every claim above would also pass for a lane that painted glow
// over a pipeline that has stopped — the difference between a live picture and a
// decorative one.
let before_dead = flow_frames_drawn();
let dead_on = frame(false, FlowMode::Particles, true, 0.0).expect("dead frame");
let dead_frames = flow_frames_drawn() - before_dead;
let dead_off = frame(false, FlowMode::Off, true, 0.0).expect("dead off frame");
let dead_px = diff_px(&dead_on, &dead_off);
eprintln!("[flowsim-gpu] dead DAG: overlay changed {dead_px} px (must be 0); device frames {dead_frames} (must be 0)");
assert_eq!(
dead_px, 0,
"a pipeline with zero throughput must not glow — the GPU overlay lit {dead_px} px over \
stages moving no rows"
);
assert_eq!(
dead_frames, 0,
"a dead DAG must not claim the GPU lane either — the counter moved by {dead_frames}, so \
`draw_flow` said yes for a frame with nothing in it and the CPU painter was skipped for \
no reason"
);
#[cfg(feature = "testmatrix")]
facett_core::testmatrix::emit(
"facett-pipeline.flowsim_gpu_lane",
"gpu_advect_path_is_driven_by_a_host",
gpu_px > floor && gpu_frames > 0 && cpu_frames == 0 && cpu_covered > 0.90 && gpu_covered > 0.80 && motion_px * 3 > gpu_px && dead_px == 0 && dead_frames == 0,
&format!(
"gpu={gpu_px}px/{gpu_frames}frames cpu={cpu_px}px/{cpu_frames}frames \
align={:.0}/{:.0}% motion={motion_px}px dead={dead_px}px/{dead_frames}frames",
100.0 * cpu_covered,
100.0 * gpu_covered
),
);
}
/// **See the frame.** The GPU-advected overlay next to the CPU-painted one, so the two
/// renderers are something a human can compare.
#[test]
fn flowsim_gpu_png_blobs() {
let Some(gpu) = frame(true, FlowMode::Particles, true, 0.31) else { return };
let cpu = frame(true, FlowMode::Particles, false, 0.31).expect("cpu frame");
let dir = std::path::Path::new(env!("CARGO_MANIFEST_DIR")).join("tests/snapshots");
std::fs::create_dir_all(&dir).unwrap();
let (gp, cp) = (dir.join("flowsim_gpu_lane.png"), dir.join("flowsim_cpu_lane.png"));
gpu.save(&gp).unwrap();
cpu.save(&cp).unwrap();
for p in [&gp, &cp] {
let n = std::fs::metadata(p).unwrap().len();
assert!(n > 1000, "{} is a real frame ({n} bytes)", p.display());
eprintln!("[flowsim-gpu] wrote {} ({n} B)", p.display());
}
}