facett-core/tests/gpu_picking.rs
//! **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
);
}