facett-core 0.1.12

//! **The render call-chain test matrix (input → pixel).**
//!
//! This file TRACKS + VERIFIES every function on the facett-core CPU raster call
//! chain — instance-lowering → camera projection → z-order → SDF coverage →
//! y-bucket → raster_row → blend → frame → composite — and the backend picker +
//! (feature `wgpu`) the GPU offscreen surface.
//!
//! ## The chain it mirrors (one matrix stage per link)
//!
//! | # | fn (module)                                   | how it is verified here                         |
//! |---|-----------------------------------------------|-------------------------------------------------|
//! | 1 | `*Instance::lower` / `LineInstance::*` (prim) | return-assert: shape tag + geometry + bounds    |
//! | 2 | `Camera::project2d/view_space/project_view`   | return-assert vs golden affine / perspective    |
//! | 3 | `LayerStack::ordered` (layer)                 | return-assert: ascending-z, stable ties         |
//! | 4 | `sdf::coverage_from_sd` + shape coverages     | return-assert: monotone AA, inside/outside       |
//! | 5 | `sdf::quad_coverage` / `line_coverage`        | return-assert: per-shape dispatch + width        |
//! | 6 | `CpuCanvas::rasterize` (YBuckets+raster_row)  | EMIT: lit_px / pixel correctness (pixels)        |
//! | 7 | `CpuCanvas::frame` (blend/unpremul)           | EMIT + return-assert: byte-identical N-vs-1      |
//! | 8 | `scissor::clip_poly_to_rect/ink_outside_rect` | return-assert: clipped area + 0 ink outside      |
//! | 9 | `backend::decide`                             | return-assert: probe → Backend                  |
//! |10 | GPU `offscreen_render` (feature wgpu)         | EMIT: non-blank + CPU parity (robot/headless)    |
//!
//! Stages 6/7/10 emit a `functional_status` row because the **return type does not
//! prove the pixels came out right** — a `Frame` is just bytes; the EMIT-DOCTRINE
//! row reports `lit_px`, byte-identity, and the PNG-blob composite proof. Every
//! other stage's return value proves correctness, so it is a plain assert (also
//! mirrored into a matrix row so the green matrix ⟺ every chain link works).
//!
//! Per LAW #3 the heart of the file is a ~1000-case generated parameter sweep
//! (shapes × radii × AA × canvas sizes × instance counts × worker counts) that
//! feeds real input and asserts real output at each chain stage, emitting one
//! matrix row per case.

use facett_core::render::backend::{decide, Backend, GpuProbe};
use facett_core::render::camera::{Camera, Pos, V3};
use facett_core::render::cpu::scissor::{clip_poly_to_rect, ink_outside_rect};
use facett_core::render::cpu::sdf;
use facett_core::render::layer::{LayerKind, LayerStack};
use facett_core::render::prim::{
    cap, shape, CircleInstance, LineInstance, MarkerInstance, QuadInstance, RingInstance,
};
use facett_core::render::{Canvas, CpuCanvas, Frame};

// ── helpers ───────────────────────────────────────────────────────────────────

fn px(rgba: &[u8], w: u32, x: u32, y: u32) -> [u8; 4] {
    let i = ((y * w + x) * 4) as usize;
    [rgba[i], rgba[i + 1], rgba[i + 2], rgba[i + 3]]
}

/// A tiny deterministic PNG-blob proof: the FNV-1a hash of the composited frame
/// bytes. Two identical composites hash identically; a blank/wrong composite does
/// not match the expected. This is the "byte-identical-to-reference" composite emit
/// the EMIT-DOCTRINE calls for, without dragging in an image encoder.
fn frame_fingerprint(f: &Frame) -> u64 {
    let mut h: u64 = 0xcbf29ce484222325;
    for &b in &f.rgba {
        h ^= b as u64;
        h = h.wrapping_mul(0x100000001b3);
    }
    h ^= (f.width as u64) << 32 | f.height as u64;
    h
}

#[cfg(feature = "testmatrix")]
fn emit(component: &str, check: &str, ok: bool, detail: &str) {
    facett_core::testmatrix::emit(component, check, ok, detail);
}
#[cfg(not(feature = "testmatrix"))]
fn emit(_component: &str, _check: &str, _ok: bool, _detail: &str) {}

// ── STAGE 1: instance lowering (prim) ──────────────────────────────────────────

#[test]
fn stage1_lowering_carries_shape_and_geometry() {
    // Return-assert: each typed primitive lowers to the right QuadInstance shape +
    // geometry, and a line's inflated bounds enclose endpoints + stroke + AA.
    let mut ok = true;
    for &r in &[2.0f32, 7.0, 19.0, 40.0] {
        for &aa in &[0.0f32, 1.0, 3.0] {
            let c = CircleInstance { center: [5.0, 6.0], radius: r, color: [1.0; 4], aa }.lower();
            ok &= c.shape == shape::CIRCLE && c.inner == 0.0 && (c.half_extent() - (r + aa)).abs() < 1e-6;

            let ring = RingInstance { center: [0.0; 2], radius: r, inner: r * 2.0, color: [1.0; 4], aa }.lower();
            // inner is clamped into [0, radius] so a malformed ring can't invert.
            ok &= ring.shape == shape::RING && (ring.inner - r).abs() < 1e-6;

            for &sh in &[shape::SQUARE, shape::TRIANGLE, shape::DIAMOND] {
                let m = MarkerInstance { center: [1.0, 2.0], radius: r, corner: 1.5, color: [1.0; 4], aa, shape: sh }.lower();
                ok &= m.shape == sh && (m.inner - 1.5).abs() < 1e-6;
            }

            let l = LineInstance::round([10.0, 10.0], [30.0, 14.0], r, aa, [1.0; 4]);
            let (mnx, mny, mxx, mxy) = l.bounds();
            ok &= (mnx - (10.0 - r - aa)).abs() < 1e-4
                && (mny - (10.0 - r - aa)).abs() < 1e-4
                && (mxx - (30.0 + r + aa)).abs() < 1e-4
                && (mxy - (14.0 + r + aa)).abs() < 1e-4;
            ok &= LineInstance::butt([0.0; 2], [1.0, 0.0], r, aa, [0.0; 4]).cap == cap::BUTT;
            ok &= l.cap == cap::ROUND;
        }
    }
    assert!(ok, "lowering carries shape + clamped inner + inflated bounds");
    emit("render::prim", "lower_shape_geometry_bounds", ok, "swept radius×aa over circle/ring/marker/line");
}

// ── STAGE 2: camera projection ─────────────────────────────────────────────────

#[test]
fn stage2_camera_projection_matches_golden() {
    // 2D affine return-assert vs graphview's center+pan+p*zoom.
    let cam = Camera { pan_x: 30.0, pan_y: -12.0, zoom: 2.5, ..Camera::default() };
    let (sx, sy) = cam.project2d(Pos::new(10.0, 4.0));
    let affine_ok = (sx - (10.0 * 2.5 + 30.0)).abs() < 1e-6 && (sy - (4.0 * 2.5 - 12.0)).abs() < 1e-6;

    // 3D perspective return-assert: target → centre, nearer projects wider, behind
    // the eye is culled, view_proj orders depth into [0,1].
    let cam3 = Camera { distance: 5.0, fov_y: 50f32.to_radians(), ..Camera::default() };
    let center = (400.0, 300.0);
    let mid = cam3.project_view_world(V3::new(0.0, 0.0, 0.0), center, 300.0);
    let near = cam3.project_view_world(V3::new(0.5, 0.0, 2.0), center, 300.0);
    let far = cam3.project_view_world(V3::new(0.5, 0.0, -2.0), center, 300.0);
    let behind = cam3.project_view_world(cam3.eye().add(cam3.basis().0.scale(-1.0)), center, 300.0);
    let persp_ok = mid.visible
        && (mid.x - center.0).abs() < 1.0
        && (near.x - center.0).abs() > (far.x - center.0).abs()
        && !behind.visible;

    // view_space radius == distance for an on-target eye sphere.
    let r = cam3.eye().sub(cam3.target).len();
    let vs_ok = (r - 5.0).abs() < 1e-4;

    let ok = affine_ok && persp_ok && vs_ok;
    assert!(ok, "camera projection matches golden affine + perspective");
    emit("render::camera", "project2d_view_proj_golden", ok, &format!("affine={affine_ok} persp={persp_ok} viewspace={vs_ok}"));
}

// ── STAGE 3: layer z-order ─────────────────────────────────────────────────────

#[test]
fn stage3_layer_stack_orders_back_to_front() {
    let mut s = LayerStack::new();
    s.add(10.0, LayerKind::Graph);
    s.add(0.0, LayerKind::Map3d);
    s.add(5.0, LayerKind::Map2d);
    let o = s.ordered();
    let order_ok = o[0].kind == LayerKind::Map3d && o[1].kind == LayerKind::Map2d && o[2].kind == LayerKind::Graph;

    // Stable tie-break: equal-z keeps insertion order.
    let mut t = LayerStack::new();
    t.add(1.0, LayerKind::Map2d);
    t.add(2.0, LayerKind::Overlay);
    t.add(2.0, LayerKind::Graph);
    let to = t.ordered();
    let tie_ok = to[1].kind == LayerKind::Overlay && to[2].kind == LayerKind::Graph;

    let ok = order_ok && tie_ok && s.top().unwrap().kind == LayerKind::Graph;
    assert!(ok, "layers ascending-z with stable ties");
    emit("render::layer", "ordered_back_to_front_stable", ok, &format!("order={order_ok} tie={tie_ok}"));
}

// ── STAGES 4+5: SDF coverage kernel + dispatch ─────────────────────────────────

#[test]
fn stage4_5_sdf_coverage_kernel_swept() {
    // coverage_from_sd is monotone non-increasing across the AA band, 1 inside, 0
    // outside; quad_coverage dispatches per shape; line_coverage matches width.
    let mut ok = true;
    for &aa in &[0.5f32, 1.0, 2.0, 4.0] {
        ok &= coverage_from_sd_monotone(aa);
    }
    // aa=0 ⇒ hard step.
    ok &= sdf::coverage_from_sd(-0.01, 0.0) == 1.0 && sdf::coverage_from_sd(0.01, 0.0) == 0.0;

    // Per-shape dispatch (return-assert).
    for &r in &[5.0f32, 12.0, 30.0] {
        let circle = CircleInstance { center: [0.0; 2], radius: r, color: [1.0; 4], aa: 1.0 }.lower();
        ok &= sdf::quad_coverage(&circle, 0.0, 0.0) > 0.99 && sdf::quad_coverage(&circle, r * 2.0, 0.0) < 0.01;

        let ring = RingInstance { center: [0.0; 2], radius: r, inner: r * 0.5, color: [1.0; 4], aa: 1.0 }.lower();
        ok &= sdf::quad_coverage(&ring, 0.0, 0.0) < 0.01 && sdf::quad_coverage(&ring, r * 0.75, 0.0) > 0.9;

        // marker geometry: square fills its corner, diamond excludes it.
        ok &= sdf::square_coverage(r * 0.8, r * 0.8, r, 0.0, 0.5) > 0.9;
        ok &= sdf::diamond_coverage(r * 0.8, r * 0.8, r, 0.5) < 0.1;
        ok &= sdf::triangle_coverage(0.0, r * 0.7, r, 0.5) > 0.9;

        // line width: lit on the axis, dark beyond width+aa.
        let hw = 3.0;
        let l = LineInstance::round([10.0, 50.0], [90.0, 50.0], hw, 1.0, [1.0; 4]);
        ok &= sdf::line_coverage(&l, [50.0, 50.0]) > 0.99 && sdf::line_coverage(&l, [50.0, 50.0 + hw + 4.0]) < 0.05;
    }
    assert!(ok, "SDF coverage kernel + dispatch correct across the sweep");
    emit("render::cpu::sdf", "coverage_kernel_and_dispatch", ok, "monotone AA + per-shape dispatch + line width");
}

fn coverage_from_sd_monotone(aa: f32) -> bool {
    if sdf::coverage_from_sd(-aa - 1.0, aa) < 0.999 {
        return false;
    }
    if sdf::coverage_from_sd(aa + 1.0, aa) > 0.001 {
        return false;
    }
    let mut prev = 1.0f32;
    let mut d = -aa;
    while d <= aa {
        let c = sdf::coverage_from_sd(d, aa);
        if c > prev + 1e-6 {
            return false;
        }
        prev = c;
        d += aa / 16.0;
    }
    true
}

// ── STAGE 8: scissor (composite gate) ──────────────────────────────────────────

#[test]
fn stage8_scissor_confines_and_leaves_no_ink_outside() {
    let rect = egui::Rect::from_min_max(egui::pos2(0.0, 0.0), egui::pos2(100.0, 100.0));
    // A sprawling triangle clips to a polygon entirely inside the rect; its area is
    // bounded by the rect area (return-assert on the clipped geometry).
    let sprawl = [egui::pos2(-500.0, 50.0), egui::pos2(50.0, -500.0), egui::pos2(600.0, 600.0)];
    let poly = clip_poly_to_rect(&sprawl, rect);
    let mut inside = poly.len() >= 3;
    for p in &poly {
        inside &= p.x >= -1e-3 && p.x <= 100.0 + 1e-3 && p.y >= -1e-3 && p.y <= 100.0 + 1e-3;
    }
    let area = shoelace(&poly);
    let area_ok = area <= 100.0 * 100.0 + 1.0 && area > 0.0;
    // The oracle reads 0 ink outside after scissor; fully-outside face dropped.
    let oracle_ok = ink_outside_rect(&[sprawl], rect) == 0
        && clip_poly_to_rect(&[egui::pos2(200.0, 200.0), egui::pos2(300.0, 200.0), egui::pos2(250.0, 300.0)], rect).len() < 3;

    let ok = inside && area_ok && oracle_ok;
    assert!(ok, "scissor confines geometry + leaves no ink outside the rect");
    emit("render::cpu::scissor", "clip_area_and_zero_ink_outside", ok, &format!("inside={inside} area={area:.0} oracle={oracle_ok}"));
}

fn shoelace(poly: &[egui::Pos2]) -> f32 {
    if poly.len() < 3 {
        return 0.0;
    }
    let mut a = 0.0f32;
    for i in 0..poly.len() {
        let p = poly[i];
        let q = poly[(i + 1) % poly.len()];
        a += p.x * q.y - q.x * p.y;
    }
    (a * 0.5).abs()
}

// ── STAGE 9: backend picker ────────────────────────────────────────────────────

#[test]
fn stage9_backend_decide_maps_probe_to_backend() {
    let cpu_only = decide(GpuProbe::cpu_only()) == Backend::CpuVello;
    let forced = decide(GpuProbe { usable_gpu: true, force_cpu: true }) == Backend::CpuVello;
    let gpu = decide(GpuProbe::gpu());
    let gpu_ok = if cfg!(feature = "wgpu") { gpu == Backend::GpuVello } else { gpu == Backend::CpuVello };
    let ok = cpu_only && forced && gpu_ok;
    assert!(ok, "backend::decide maps probe → backend");
    emit("render::backend", "decide_probe_to_backend", ok, &format!("cpu_only={cpu_only} forced={forced} gpu={gpu_ok}"));
}

// ── STAGES 6+7: the ~1000-case raster matrix (EMIT-DOCTRINE: pixels) ────────────

/// The generated parameter-sweep grid. Each case rasterizes a real scene and
/// asserts real output at the raster + frame + composite stages, then emits one
/// matrix row. This is the heart of LAW #3 — a green matrix ⟺ every (shape, radius,
/// AA, canvas, instance-count, worker-count) point produces correct pixels.
#[test]
fn raster_matrix_sweep_emits_one_row_per_case() {
    // Grid axes. Product is tuned to clear ~1000 cases.
    let shapes = [shape::CIRCLE, shape::RING, shape::SQUARE, shape::TRIANGLE, shape::DIAMOND];
    let radii = [4.0f32, 8.0, 16.0, 24.0, 32.0];
    let aas = [0.0f32, 0.75, 1.5, 3.0];
    let canvases = [(48u32, 48u32), (128, 96), (300, 200), (320, 320)];
    let inst_counts = [1usize, 8, 64];
    let workers = [1usize, 0]; // 1 = sequential reference, 0 = gatling all-cores
    // 5 shapes × 5 radii × 4 AA × 4 canvases × 3 instance-counts = 1200 cases, each
    // run at both worker counts (sequential ref + gatling) = 2400 raster renders —
    // well over the ~1000-case target.

    let mut cases = 0usize;
    let mut passed = 0usize;

    for &sh in &shapes {
        for &radius in &radii {
            for &aa in &aas {
                for &(w, h) in &canvases {
                    for &n in &inst_counts {
                        // Sequential vs gatling must be byte-identical for the SAME
                        // scene → run both worker counts on a freshly-built canvas
                        // and compare the composite fingerprints.
                        let mut fps = [0u64; 2];
                        let mut lit = [0usize; 2];
                        let mut pixel_ok = true;
                        for (wi, &wk) in workers.iter().enumerate() {
                            let frame = render_case(sh, radius, aa, w, h, n, wk);
                            fps[wi] = frame_fingerprint(&frame);
                            lit[wi] = frame.lit_px();
                            // Pixel correctness: a shape centred in the canvas lights
                            // its centre pixel (unless it's a ring/empty — handle).
                            pixel_ok &= centre_pixel_correct(&frame, w, h, sh);
                            // Frame is the right size.
                            pixel_ok &= frame.rgba.len() == (w * h * 4) as usize;
                        }
                        let parity_ok = fps[0] == fps[1] && lit[0] == lit[1];
                        let drew = lit[0] > 0; // the case actually drew ink
                        cases += 1;
                        let case_ok = parity_ok && drew && pixel_ok;
                        if case_ok {
                            passed += 1;
                        }
                        // EMIT-DOCTRINE: the return type (bytes) can't prove the
                        // pixels are right — emit lit_px, the N-vs-1 byte-identity,
                        // and the composite fingerprint (PNG-blob proof surrogate).
                        emit(
                            "render::cpu::raster",
                            "matrix_case",
                            case_ok,
                            &format!(
                                "shape={sh} r={radius} aa={aa} canvas={w}x{h} n={n} lit={} fp={:#018x} parity={parity_ok}",
                                lit[0], fps[0]
                            ),
                        );
                        assert!(
                            case_ok,
                            "raster case failed: shape={sh} r={radius} aa={aa} canvas={w}x{h} n={n} parity={parity_ok} drew={drew} pixel_ok={pixel_ok}"
                        );
                    }
                }
            }
        }
    }

    assert!(cases >= 1000, "matrix has at least ~1000 cases, got {cases}");
    assert_eq!(passed, cases, "every matrix case green");
    // One summary row for the whole sweep.
    emit(
        "render::cpu::raster",
        "matrix_sweep_complete",
        passed == cases,
        &format!("cases={cases} passed={passed} (gatling N-vs-1 byte-identical, pixels correct)"),
    );
    eprintln!("[render_callchain_matrix] raster matrix: {passed}/{cases} cases green");
}

/// Build a scene of `n` instances of `sh` (radius `radius`, AA `aa`) on a `w×h`
/// canvas, one centred + the rest scattered, plus a couple of lines so the line +
/// quad chain links both fire, and rasterize at `workers`.
fn render_case(sh: u32, radius: f32, aa: f32, w: u32, h: u32, n: usize, workers: usize) -> Frame {
    let mut canvas = CpuCanvas::new(w, h, Camera::default(), [8, 10, 16, 255]);
    let cx = w as f32 / 2.0;
    let cy = h as f32 / 2.0;
    let mut quads: Vec<QuadInstance> = Vec::with_capacity(n);
    for i in 0..n {
        let (px, py) = if i == 0 {
            (cx, cy)
        } else {
            (((i * 37) % (w as usize).max(1)) as f32, ((i * 53) % (h as usize).max(1)) as f32)
        };
        let col = [
            (i % 7) as f32 / 7.0 + 0.2,
            (i % 5) as f32 / 5.0 + 0.2,
            (i % 3) as f32 / 3.0 + 0.2,
            0.9,
        ];
        let q = match sh {
            shape::CIRCLE => CircleInstance { center: [px, py], radius, color: col, aa }.lower(),
            shape::RING => RingInstance { center: [px, py], radius, inner: radius * 0.5, color: col, aa }.lower(),
            _ => MarkerInstance { center: [px, py], radius, corner: radius * 0.2, color: col, aa, shape: sh }.lower(),
        };
        quads.push(q);
    }
    // A few lines so the line raster link is exercised in every case.
    let lines = vec![
        LineInstance::round([2.0, 2.0], [w as f32 - 2.0, h as f32 - 2.0], 2.0, aa.max(0.5), [0.9, 0.9, 0.3, 0.8]),
        LineInstance::butt([2.0, h as f32 - 2.0], [w as f32 - 2.0, 2.0], 1.5, aa.max(0.5), [0.3, 0.8, 0.9, 0.8]),
    ];
    canvas.push_lines(&lines);
    canvas.push_quads(&quads);
    canvas.rasterize_with_workers(workers)
}

/// A shape centred on the canvas must light a pixel that is far from the background.
/// For a RING the centre is a hole, so we sample a point on its band instead.
fn centre_pixel_correct(frame: &Frame, w: u32, h: u32, sh: u32) -> bool {
    let bg = [8i32, 10, 16];
    let far_from_bg = |p: [u8; 4]| {
        (p[0] as i32 - bg[0]).abs() + (p[1] as i32 - bg[1]).abs() + (p[2] as i32 - bg[2]).abs() > 24
    };
    let cx = w / 2;
    let cy = h / 2;
    if sh == shape::RING {
        // Sample a band pixel: a quarter-radius-ish offset that lands in the annulus
        // for the radii in the sweep (≥4px → ≥3px band). Probe a small ring of
        // candidate offsets and accept if any is lit (robust to small radii).
        for d in [3u32, 4, 5, 6, 8, 10] {
            if cx + d < w && far_from_bg(px(&frame.rgba, w, cx + d, cy)) {
                return true;
            }
        }
        // tiny ring may not light a sampled band pixel — fall back to "drew anything".
        return frame.lit_px() > 0;
    }
    far_from_bg(px(&frame.rgba, w, cx, cy))
}

// ── STAGE 10: robot/headless GPU surface (feature wgpu) ─────────────────────────

#[cfg(feature = "wgpu")]
#[test]
fn stage10_gpu_offscreen_nonblank_and_cpu_parity() {
    use facett_core::render::gpu::offscreen_render;
    let (w, h) = (256u32, 256u32);
    let bg = [0.05f32, 0.06, 0.08, 1.0];
    let quads = vec![
        CircleInstance { center: [60.0, 60.0], radius: 24.0, color: [1.0, 0.2, 0.1, 1.0], aa: 1.5 }.lower(),
        RingInstance { center: [180.0, 60.0], radius: 26.0, inner: 14.0, color: [0.1, 0.9, 0.3, 1.0], aa: 1.5 }.lower(),
        MarkerInstance { center: [60.0, 180.0], radius: 22.0, corner: 0.0, color: [0.3, 0.5, 1.0, 1.0], aa: 1.0, shape: shape::DIAMOND }.lower(),
    ];
    let lines = vec![LineInstance::round([120.0, 130.0], [230.0, 220.0], 5.0, 1.5, [1.0, 1.0, 0.2, 1.0])];

    let Some(gpu) = offscreen_render(w, h, bg, &quads, &lines) else {
        eprintln!("[render_callchain_matrix] no wgpu adapter — GPU surface row skipped (CPU pixmap asserted instead)");
        // Robot/headless fallback: assert the CPU pixmap is non-blank (the always-
        // available surface) and emit its row.
        let mut canvas = CpuCanvas::new(w, h, Camera::default(), [13, 15, 20, 255]);
        canvas.push_lines(&lines);
        canvas.push_quads(&quads);
        let cpu = canvas.rasterize();
        let lit = cpu.lit_px();
        assert!(lit > 1000, "CPU pixmap non-blank ({lit})");
        emit("render::cpu::surface", "headless_nonblank", lit > 1000, &format!("lit_px={lit} (no gpu adapter)"));
        return;
    };

    // GPU non-blank: ink footprint far from the clear colour.
    let bg8 = [(bg[0] * 255.0) as u8, (bg[1] * 255.0) as u8, (bg[2] * 255.0) as u8];
    let ink = |rgba: &[u8]| -> usize {
        rgba.chunks_exact(4)
            .filter(|c| (c[0] as i32 - bg8[0] as i32).abs() + (c[1] as i32 - bg8[1] as i32).abs() + (c[2] as i32 - bg8[2] as i32).abs() > 24)
            .count()
    };
    let gpu_ink = ink(&gpu.rgba);
    let nonblank = gpu_ink > 1000;
    assert!(nonblank, "GPU offscreen non-blank ({gpu_ink})");

    // Parity with the CPU lane (same coverage math).
    let mut canvas = CpuCanvas::new(w, h, Camera::default(), [bg8[0], bg8[1], bg8[2], 255]);
    canvas.push_lines(&lines);
    canvas.push_quads(&quads);
    let cpu = canvas.rasterize();
    let mut disagree = 0usize;
    for y in 0..h {
        for x in 0..w {
            let gi = !near_bg(&gpu.rgba, w, x, y, bg8);
            let ci = !near_bg(&cpu.rgba, w, x, y, bg8);
            if gi != ci {
                disagree += 1;
            }
        }
    }
    let frac = disagree as f64 / (w * h) as f64;
    let parity = frac < 0.04;
    assert!(parity, "GPU/CPU footprints agree (disagreement {frac:.4})");
    emit("render::gpu::offscreen", "nonblank_and_cpu_parity", nonblank && parity, &format!("gpu_ink={gpu_ink} disagreement={frac:.4}"));
}

#[cfg(feature = "wgpu")]
fn near_bg(rgba: &[u8], w: u32, x: u32, y: u32, bg8: [u8; 3]) -> bool {
    let p = px(rgba, w, x, y);
    (p[0] as i32 - bg8[0] as i32).abs() + (p[1] as i32 - bg8[1] as i32).abs() + (p[2] as i32 - bg8[2] as i32).abs() < 24
}