facett-core 0.1.19

//! **GFX_V2 item 8 — the colour-encoded ID pass really resolves a click** (needs
//! `--features wgpu`).
//!
//! Every assertion here is on **the id the DEVICE wrote, read back out of the
//! `R32Uint` target at a named pixel**. Not `pass.recorded()`, not a vertex count,
//! not "the target is non-blank". Item 2's lesson on this track was that a pixel
//! assertion stayed green through **0 device frames** because a CPU painter was
//! drawing; the equivalent lie here would be a picking test that passes while the
//! shader writes a constant. So:
//!
//! * **six distinct objects** are on screen at once, never one. A single-object
//!   probe cannot tell a working id encode from `return 0x01000000;`.
//! * their ids are chosen to sit on the **channel-carry boundaries** of the 4-byte
//!   little-endian wire encoding, including `0x0100_0000` (whose low three bytes are
//!   all zero — a pass that dropped a channel would still get this one right, which
//!   is exactly why the others are here) and `0xFFFF_FFFF` (the ceiling).
//! * a probe on **empty space** must return the reserved miss sentinel, and the
//!   sentinel must be distinguishable from a real feature 0.
//! * two **adjacent** quads are probed either side of their shared edge, with ids
//!   whose arithmetic mean is *itself a plausible id* — so any filtering, blending
//!   or MSAA resolve produces a visibly wrong answer rather than a rounding error.
//!
//! Run: `cargo test -p facett-core --features wgpu --test gpu_picking -- --nocapture`
#![cfg(feature = "wgpu")]

use facett_core::engine::pick::{PickId, MAX_FEATURE};
use facett_core::render::gpu::picking::{PickBatch, PickPass, PickTarget, PICK_DEPTH_FORMAT};
use facett_core::render::gpu::request_best_adapter;
use facett_core::render::gpu::preferred_backends;
use egui::{pos2, Rect};

/// The component every row and every SKIP in this file is filed under.
const GATE_COMPONENT: &str = "facett-core";
/// What a skipped arm of this file would have proven, so a SKIPPED matrix row is
/// readable without opening the test.
const GATE_DETAIL: &str = "the GPU id-buffer pick lane on a real device — one id per object, no bleed, and parity with the CPU lane";

/// A real device via the **sanctioned** selector (`request_best_adapter`), not an
/// eleventh hand-rolled `request_adapter(&Default::default())` — that one asks for
/// `LowPower` and picks the wrong GPU on a hybrid box.
///
/// `None` (⇒ the test skips) only when the box has no WebGPU-class adapter at all.
fn device() -> Option<(wgpu::Device, wgpu::Queue)> {
    let instance = wgpu::Instance::new(wgpu::InstanceDescriptor {
        backends: preferred_backends(),
        flags: wgpu::InstanceFlags::from_build_config().with_env(),
        backend_options: wgpu::BackendOptions::from_env_or_default(),
        memory_budget_thresholds: wgpu::MemoryBudgetThresholds::default(),
        display: None,
    });
    let adapter = request_best_adapter(&instance, preferred_backends())?;
    // ATTRIBUTION, not decoration. Item 2's key red on this track was a pixel
    // assertion passing through ZERO device frames because a CPU painter drew. Here
    // the equivalent lie is the graceful `None` skip above turning the whole file
    // green on a box that never ran a shader. Print the device that answered, so
    // `-- --nocapture` distinguishes "6 ids proven" from "6 ids skipped".
    let info = adapter.get_info();
    eprintln!("[gpu_picking] device: {} ({:?}, {:?})", info.name, info.backend, info.device_type);
    pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor {
        label: Some("gpu_picking_test"),
        required_features: wgpu::Features::empty(),
        required_limits: wgpu::Limits::downlevel_defaults(),
        ..Default::default()
    }))
    .ok()
}

const W: u32 = 200;
const H: u32 = 120;

/// The six-object fixture. Each entry is `(name, rect, id, probe)` where `probe` is a
/// pixel **strictly inside** the rect, and every id is a different byte pattern
/// exercising a different carry boundary of the little-endian wire encoding.
fn fixture() -> Vec<(&'static str, Rect, PickId, (u32, u32))> {
    vec![
        // low three bytes ALL ZERO — the floor of the real id space. A pass that
        // wrote only the alpha/high byte would still pass on this one alone.
        ("floor 0x01000000", Rect::from_min_max(pos2(10.0, 10.0), pos2(50.0, 50.0)), PickId::new(1, 0), (30, 30)),
        // byte 0 saturated, one below the R→G carry.
        ("pre-carry 0x010000FF", Rect::from_min_max(pos2(60.0, 10.0), pos2(100.0, 50.0)), PickId::new(1, 0xFF), (80, 30)),
        // the R→G carry itself.
        ("R→G 0x02000100", Rect::from_min_max(pos2(110.0, 10.0), pos2(150.0, 50.0)), PickId::new(2, 0x100), (130, 30)),
        // the G→B carry.
        ("G→B 0x01010000", Rect::from_min_max(pos2(10.0, 60.0), pos2(50.0, 100.0)), PickId::new(1, 0x1_0000), (30, 80)),
        // feature ceiling: all 24 feature bits set, layer still small.
        ("feature ceiling 0x03FFFFFF", Rect::from_min_max(pos2(60.0, 60.0), pos2(100.0, 100.0)), PickId::new(3, MAX_FEATURE), (80, 80)),
        // the absolute ceiling of the id space.
        ("ceiling 0xFFFFFFFF", Rect::from_min_max(pos2(110.0, 60.0), pos2(150.0, 100.0)), PickId::new(255, MAX_FEATURE), (130, 80)),
    ]
}

/// Run one id pass over `batch` into a fresh target and hand back the target.
///
/// A **FRESH** target every time, deliberately (item 2's RESIDUE finding): a reused
/// target holds the previous pass's correct answers, so a pass that wrote nothing at
/// all would read back last frame's right result. wgpu zero-initialises a new
/// texture, so an unwritten texel here reads `0` = the miss sentinel.
fn run_pass(device: &wgpu::Device, queue: &wgpu::Queue, batch: &PickBatch, depth: bool) -> (PickTarget, u32) {
    let mut target = if depth { PickTarget::with_depth() } else { PickTarget::new() };
    target.ensure(device, W, H);
    let mut pass = PickPass::new(device, depth.then_some(PICK_DEPTH_FORMAT));
    pass.set_viewport(queue, W, H);
    pass.upload(device, queue, batch);

    let mut enc = device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("pick_test") });
    let drawn = {
        let mut rp = target.begin_pass(&mut enc).expect("target is allocated");
        pass.record(&mut rp)
    };
    queue.submit(Some(enc.finish()));
    device.poll(wgpu::PollType::wait_indefinitely()).ok();
    (target, drawn)
}

/// **The headline.** Six distinct objects, six known pixels, six exact ids — plus
/// empty space resolving to the miss sentinel.
#[test]
fn six_distinct_objects_each_resolve_to_their_own_id_at_a_known_pixel() {
    let Some((device, queue)) = device() else {
        facett_core::testmatrix::gpu_skip(GATE_COMPONENT, "six_distinct_objects_each_resolve_to_their_own_id_at_a_known_pixel", "no WebGPU-class adapter", GATE_DETAIL);
        return;
    };
    let fx = fixture();
    let mut batch = PickBatch::new();
    for (_, rect, id, _) in &fx {
        batch.push_quad(*rect, *id);
    }
    let (target, drawn) = run_pass(&device, &queue, &batch, false);
    assert_eq!(drawn, 36, "six quads is 36 vertices — the pass must actually record them");

    // Every id distinct, so a constant-returning shader cannot satisfy more than one.
    let ids: Vec<u32> = fx.iter().map(|(_, _, id, _)| id.0).collect();
    let mut uniq = ids.clone();
    uniq.sort_unstable();
    uniq.dedup();
    assert_eq!(uniq.len(), fx.len(), "the fixture must have six DISTINCT ids or it proves nothing");

    for (name, _, want, (px, py)) in &fx {
        let got = target.read_id(&device, &queue, *px, *py);
        assert_eq!(
            got, *want,
            "{name}: pixel ({px},{py}) read {:#010x} (layer {} feature {}), expected {:#010x} (layer {} feature {})",
            got.0,
            got.layer(),
            got.feature(),
            want.0,
            want.layer(),
            want.feature()
        );
        assert!(!got.is_nothing(), "{name}: a drawn object must never read as a miss");
    }

    // A MISS: the gutter right of every quad. Must be the sentinel — not id 0 meaning
    // "feature 0", not a neighbour's id, not garbage.
    for (px, py) in [(180, 110), (0, 0), (155, 55), (105, 55)] {
        let got = target.read_id(&device, &queue, px, py);
        assert!(
            got.is_nothing(),
            "empty pixel ({px},{py}) read {:#010x} — a miss must be PickId::NOTHING",
            got.0
        );
    }
    // ...and the sentinel is NOT what a real feature 0 reads as, which is the whole
    // point of reserving layer 0.
    let feature_zero = target.read_id(&device, &queue, 30, 30);
    assert_eq!(feature_zero, PickId::new(1, 0));
    assert_eq!(feature_zero.feature(), 0, "this object really is feature 0");
    assert!(!feature_zero.is_nothing(), "feature 0 must be distinguishable from a miss");
}

/// **No bleed across a shared edge.** Two quads meet at exactly x = 100. The ids are
/// feature 2 and feature 4 of the same layer, so their arithmetic mean is feature 3 —
/// a perfectly plausible object that no pixel is allowed to report. Any filtering,
/// blending or MSAA resolve on the id target produces it.
#[test]
fn adjacent_objects_do_not_bleed_across_their_shared_edge() {
    let Some((device, queue)) = device() else {
        facett_core::testmatrix::gpu_skip(GATE_COMPONENT, "adjacent_objects_do_not_bleed_across_their_shared_edge", "no WebGPU-class adapter", GATE_DETAIL);
        return;
    };
    let left = PickId::new(1, 2); // 0x01000002
    let right = PickId::new(1, 4); // 0x01000004
    let blend = PickId::new(1, 3); // 0x01000003 — the forbidden average
    assert_eq!((left.0 + right.0) / 2, blend.0, "the fixture's mean really is a valid id");

    let mut batch = PickBatch::new();
    batch.push_quad(Rect::from_min_max(pos2(0.0, 0.0), pos2(100.0, 60.0)), left);
    batch.push_quad(Rect::from_min_max(pos2(100.0, 0.0), pos2(200.0, 60.0)), right);
    let (target, _) = run_pass(&device, &queue, &batch, false);

    let a = target.read_id(&device, &queue, 99, 30);
    let b = target.read_id(&device, &queue, 100, 30);
    assert_eq!(a, left, "1 px left of the edge must be the left object, got {:#010x}", a.0);
    assert_eq!(b, right, "1 px right of the edge must be the right object, got {:#010x}", b.0);
    assert_ne!(a, b, "the edge separates two DIFFERENT ids — otherwise this proves nothing");

    // Sweep the whole 20 px band across the seam: not one texel may carry the mean,
    // and every texel must be exactly one of the two real ids.
    let band = target.read_region(&device, &queue, 90, 20, 20, 20);
    assert_eq!(band.len(), 400, "the readback must return the region it was asked for");
    for (i, got) in band.iter().enumerate() {
        let x = 90 + (i % 20) as u32;
        assert_ne!(*got, blend, "texel x={x} carries the BLEND of two ids — the target is filtering");
        assert!(
            *got == left || *got == right,
            "texel x={x} read {:#010x}, which is neither neighbour — the id space is being interpolated",
            got.0
        );
        let want = if x < 100 { left } else { right };
        assert_eq!(*got, want, "texel x={x} is on the wrong side of the seam");
    }
}

/// With a depth attachment the **nearest** id wins, and it wins in **both submission
/// orders**. Without a depth test the answer is "whichever drew last", so the two
/// orders disagree — which is what makes this pair of assertions non-vacuous.
#[test]
fn the_nearest_id_wins_regardless_of_draw_order() {
    let Some((device, queue)) = device() else {
        facett_core::testmatrix::gpu_skip(GATE_COMPONENT, "the_nearest_id_wins_regardless_of_draw_order", "no WebGPU-class adapter", GATE_DETAIL);
        return;
    };
    let near = PickId::new(4, 11);
    let far = PickId::new(4, 22);
    let overlap = Rect::from_min_max(pos2(40.0, 30.0), pos2(120.0, 90.0));

    let mut answers = Vec::new();
    for near_first in [true, false] {
        let mut batch = PickBatch::new();
        let push_near = |b: &mut PickBatch| b.push_quad_at_depth(overlap, near, 0.25);
        let push_far = |b: &mut PickBatch| b.push_quad_at_depth(overlap, far, 0.75);
        if near_first {
            push_near(&mut batch);
            push_far(&mut batch);
        } else {
            push_far(&mut batch);
            push_near(&mut batch);
        }
        let (target, drawn) = run_pass(&device, &queue, &batch, true);
        assert_eq!(drawn, 12, "two quads, 12 vertices");
        assert!(target.has_depth(), "this lane must have a depth attachment");
        let got = target.read_id(&device, &queue, 80, 60);
        assert_eq!(
            got, near,
            "near_first={near_first}: the visible (nearest) object is {:#010x}, read {:#010x}",
            near.0, got.0
        );
        answers.push(got);
    }
    assert_eq!(answers[0], answers[1], "the answer must not depend on submission order");
}

/// The pass **clears** to the miss sentinel — it does not merely start on a
/// zero-initialised texture.
///
/// This is the residue trap from item 2, inverted: draw a full-target id, prove it is
/// there, then run a second pass over the SAME target with nothing to draw. If the
/// clear were `LoadOp::Load` the first pass's ids would still be sitting there and a
/// "clicked empty space" probe would return a stale object.
#[test]
fn a_second_pass_clears_the_previous_frames_ids_to_the_miss_sentinel() {
    let Some((device, queue)) = device() else {
        facett_core::testmatrix::gpu_skip(GATE_COMPONENT, "a_second_pass_clears_the_previous_frames_ids_to_the_miss_sentinel", "no WebGPU-class adapter", GATE_DETAIL);
        return;
    };
    let id = PickId::new(9, 12345);
    let mut batch = PickBatch::new();
    batch.push_quad(Rect::from_min_max(pos2(0.0, 0.0), pos2(W as f32, H as f32)), id);
    let (target, drawn) = run_pass(&device, &queue, &batch, false);
    assert_eq!(drawn, 6);
    assert_eq!(target.read_id(&device, &queue, 100, 60), id, "frame 1 must have painted the whole target");

    // Frame 2 on the SAME target: an empty batch, so only the clear runs.
    let mut empty = PickPass::new(&device, None);
    empty.set_viewport(&queue, W, H);
    empty.upload(&device, &queue, &PickBatch::new());
    let mut enc = device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("pick_clear") });
    let drawn2 = {
        let mut rp = target.begin_pass(&mut enc).expect("still allocated");
        empty.record(&mut rp)
    };
    queue.submit(Some(enc.finish()));
    device.poll(wgpu::PollType::wait_indefinitely()).ok();
    assert_eq!(drawn2, 0, "an empty batch draws nothing — but the pass still ran");

    for (px, py) in [(100, 60), (0, 0), (W - 1, H - 1)] {
        let got = target.read_id(&device, &queue, px, py);
        assert!(
            got.is_nothing(),
            "({px},{py}) still reads {:#010x} from the PREVIOUS frame — the pass is not clearing",
            got.0
        );
    }
}

/// The logical-coordinate entry point lands on the same texel as the physical one, at
/// a non-1.0 `pixels_per_point` — where an unscaled click silently probes the wrong
/// object. A hi-dpi host is the common case, so `ppp = 1.0` is the identity value
/// this test must not sit on.
#[test]
fn a_logical_click_at_hidpi_resolves_the_object_under_the_cursor() {
    let Some((device, queue)) = device() else {
        facett_core::testmatrix::gpu_skip(GATE_COMPONENT, "a_logical_click_at_hidpi_resolves_the_object_under_the_cursor", "no WebGPU-class adapter", GATE_DETAIL);
        return;
    };
    let a = PickId::new(1, 100);
    let b = PickId::new(1, 200);
    let mut batch = PickBatch::new();
    // In PHYSICAL px: a on the left half, b on the right half.
    batch.push_quad(Rect::from_min_max(pos2(0.0, 0.0), pos2(100.0, 120.0)), a);
    batch.push_quad(Rect::from_min_max(pos2(100.0, 0.0), pos2(200.0, 120.0)), b);
    let (target, _) = run_pass(&device, &queue, &batch, false);

    // A cursor at logical x = 60 with ppp = 2.0 is physical x = 120 → object b.
    // Ignoring ppp would probe physical 60 → object a. The two answers differ, which
    // is the only reason this assertion means anything.
    assert_eq!(
        target.read_id_logical(&device, &queue, pos2(60.0, 30.0), 2.0),
        b,
        "a hi-dpi click must scale into physical texels"
    );
    assert_eq!(target.read_id(&device, &queue, 60, 60), a, "physical 60 really is the OTHER object");
    // And logical 30 @ ppp 2 = physical 60 = a, closing the loop.
    assert_eq!(target.read_id_logical(&device, &queue, pos2(30.0, 30.0), 2.0), a);

    // A cursor outside the widget is a miss, not a clamp onto the edge object.
    assert!(target.read_id_logical(&device, &queue, pos2(-4.0, 30.0), 2.0).is_nothing());
}

/// **The two lanes agree, texel by texel.** This is the proof `engine::pick` was
/// designed for: *"it is the same 32-bit id the CPU lane returns, so a test can assert
/// the two lanes agree."*
///
/// The GPU pass and `PickIndex` share nothing but the id encoding — one rasterises
/// quads on a 4090, the other bins anchors into a `ScreenGrid` on the CPU. Two
/// independent implementations returning the same `PickId` at 500 probe pixels is a
/// far stronger statement than either lane checked against its own expectations, and
/// it is the only assertion here that could catch a bug living in the *fixture*.
///
/// **The tie-break is the whole difficulty, and it is a real bug this caught.**
/// `PickShape::Box` documents "the LOWEST feature index that contains the probe wins —
/// paint order"; a colour-encoded pass with no depth answers "whichever quad was
/// submitted last". For overlapping chips those are DIFFERENT objects. So the GPU lane
/// goes through `push_quad_in_painter_order`, which maps paint index to depth
/// (lower = nearer) so the `Less` test reproduces the CPU rule by construction.
#[test]
fn the_gpu_lane_and_the_cpu_lane_return_the_same_id_at_every_probe() {
    let Some((device, queue)) = device() else {
        facett_core::testmatrix::gpu_skip(GATE_COMPONENT, "the_gpu_lane_and_the_cpu_lane_return_the_same_id_at_every_probe", "no WebGPU-class adapter", GATE_DETAIL);
        return;
    };
    use facett_core::engine::pick::{PickIndex, PickShape};
    use facett_core::render::gpu::picking::painter_depth;

    // Chips that OVERLAP — 60x28 px boxes on a 34x16 px pitch. Overlap is mandatory:
    // over disjoint chips every tie-break rule agrees, so the test would pass with
    // the painter-order depth deleted. Item 2's identity-copy trap, in another shape.
    const LAYER: u8 = 5;
    let half = egui::vec2(30.0, 14.0);
    let anchors: Vec<egui::Pos2> =
        (0..24).map(|i| pos2(34.0 + (i % 6) as f32 * 26.0, 30.0 + (i / 6) as f32 * 22.0)).collect();
    let n = anchors.len();

    let mut cpu = PickIndex::default();
    cpu.begin(1);
    cpu.push_layer(LAYER, &anchors, (0..n as u32).collect(), PickShape::Box(half));

    // SUBMITTED IN REVERSE PAINT ORDER — and this is the single most important line
    // in the file. The first version of this test pushed the chips in paint order,
    // and it stayed GREEN with `push_quad_in_painter_order` replaced by a plain
    // `push_quad`: with every quad at depth 0.0, a strict `Less` test keeps the
    // FIRST-submitted fragment, which in paint order happens to be the lowest index —
    // the CPU rule, reproduced by accident. The depth ramp was a no-op and the guard
    // could not see it.
    //
    // That is item 2's IDENTITY-COPY trap in a new shape: when the submission order
    // already equals the intended order, the mapping between them is unobservable.
    // Reversing the submission order is the same remedy item 2 used (select a genuine
    // relocation, not a prefix), and it makes the depth mapping load-bearing: paint
    // order now has to come from the DEPTH VALUE, because submission order actively
    // contradicts it.
    let mut batch = PickBatch::new();
    for (i, a) in anchors.iter().enumerate().rev() {
        batch.push_quad_in_painter_order(Rect::from_center_size(*a, half * 2.0), PickId::new(LAYER, i as u32), i, n);
    }
    // Painter order really is encoded: index 0 must be strictly nearer than index n-1.
    assert!(painter_depth(0, n) < painter_depth(n - 1, n), "the depth ramp must be monotonic");
    let (target, drawn) = run_pass(&device, &queue, &batch, true);
    assert_eq!(drawn as usize, n * 6, "every chip must reach the device");

    // Probe a grid covering the chip field and the empty margin around it.
    let mut agreed_hits = 0usize;
    let mut agreed_misses = 0usize;
    let mut overlap_probes = 0usize;
    for gx in 0..25u32 {
        for gy in 0..20u32 {
            let (px, py) = (gx * 8, gy * 6);
            if px >= W || py >= H {
                continue;
            }
            // The CPU lane probes the PIXEL CENTRE, which is the point the rasteriser
            // tested when it decided texel (px,py)'s coverage. Probing the corner
            // instead disagrees on every boundary texel — a half-pixel convention
            // mismatch, and the reason a naive version of this test is flaky.
            let probe = pos2(px as f32 + 0.5, py as f32 + 0.5);
            let want = cpu.pick(probe).id;
            let got = target.read_id(&device, &queue, px, py);
            assert_eq!(
                got, want,
                "pixel ({px},{py}): GPU says {:#010x} (feature {}), CPU says {:#010x} (feature {})",
                got.0,
                got.feature(),
                want.0,
                want.feature()
            );
            if want.is_nothing() {
                agreed_misses += 1;
            } else {
                agreed_hits += 1;
                // How many chips contain this probe? >1 means the tie-break decided it.
                let covering = anchors
                    .iter()
                    .filter(|c| (probe.x - c.x).abs() <= half.x && (probe.y - c.y).abs() <= half.y)
                    .count();
                if covering > 1 {
                    overlap_probes += 1;
                }
            }
        }
    }
    // NON-VACUITY. A sweep that only ever landed on empty space, or never on an
    // overlap, would agree trivially and prove nothing.
    assert!(agreed_hits > 100, "only {agreed_hits} probes hit a chip — the sweep proves nothing");
    assert!(agreed_misses > 10, "only {agreed_misses} probes missed — the sentinel path is untested");
    assert!(
        overlap_probes > 50,
        "only {overlap_probes} probes landed where chips OVERLAP — the tie-break rule is untested, \
         and it is the one place the two lanes can legitimately disagree"
    );
    eprintln!(
        "[gpu_picking] two lanes agreed on {} hits / {} misses, {} of them contested by >1 chip",
        agreed_hits, agreed_misses, overlap_probes
    );
}