facett-core/src/render/gpu/label_collide.rs
//! **GPU label collision + the indirect label draw** (feature `wgpu`) — GFX_V2 §3.B,
//! item 4.
//!
//! Three compute passes decide which labels survive, and the third one *produces the
//! draw*: it appends the winners' glyph quads and bumps `instance_count` on a
//! `DrawIndirectArgs` by `atomicAdd`. The render pass then reads that count straight
//! out of the same buffer at the same offset. **The CPU never learns how many labels
//! were drawn** — the same contract item 2 established for the line lane, for the same
//! reason: a readback is a stall, and a CPU-side count is a second opinion that can
//! disagree with the device.
//!
//! The collision rule itself is not here. It lives in [`crate::label_grid`], in
//! ordinary Rust, and `label_collide.wgsl` is its transcription;
//! [`cpu_and_gpu_lanes_resolve_the_same_labels`](tests) runs both over one fixture and
//! asserts the surviving sets are equal, which is what stops the fail-safe painter
//! drifting from the device.
//!
//! ## What draws the glyphs
//!
//! `label_draw.wgsl`, sampling **epaint's own font atlas** — the texture egui's text
//! pass already uses. That choice is what makes a GPU-drawn street name identical to
//! the CPU-painted one, and it keeps a font rasteriser, an offline `msdf-atlas-gen`
//! step and a second glyph cache out of the tree. [`super::MsdfText`] stays the right
//! pipeline for a real MSDF atlas; see `label_draw.wgsl` for why its median-of-three
//! reconstruction is wrong arithmetic over a coverage field.
use wgpu::util::DeviceExt as _;
use crate::label_grid::{
GlyphSrc, LabelCandidate, LabelFrame, LabelGlyphInstance, GRID_CELLS, GRID_WORDS,
INSTANCES_PER_GLYPH, LABEL_ROUNDS, MAX_VISIBLE_LABELS, NAME_SLOTS,
};
use super::DrawIndirectArgs;
/// The collision compute passes (`cs_clear_frame` / `cs_clear_round` / `cs_claim` /
/// `cs_name_bid` / `cs_emit`).
pub use crate::render::wgsl::LABEL_COLLIDE_WGSL;
/// The instanced glyph draw whose instance count comes from the emit pass.
pub use crate::render::wgsl::LABEL_DRAW_WGSL;
/// Draw instances the output buffer holds. `MAX_VISIBLE_LABELS` labels of up to 48
/// glyphs, each expanded to [`INSTANCES_PER_GLYPH`] — so the emit pass's capacity
/// guard is reachable only by a genuinely absurd label, not by a normal frame.
pub const LABEL_GLYPH_CAPACITY: u32 = MAX_VISIBLE_LABELS * 48 * INSTANCES_PER_GLYPH;
/// The compute uniform. 80 bytes; mirrors `LabelUniforms` in `label_collide.wgsl`
/// (`ink`/`halo` are `vec4`, so they sit at 16-aligned offsets 48 and 64).
#[repr(C)]
#[derive(Clone, Copy, Debug, PartialEq, bytemuck::Pod, bytemuck::Zeroable)]
struct LabelUniforms {
zoom: [f32; 2],
ref_pos: [f32; 2],
screen_center: [f32; 2],
viewport: [f32; 2],
inv_cell: [f32; 2],
inv_repeat: f32,
lod: u32,
ink: [f32; 4],
halo: [f32; 4],
}
impl LabelUniforms {
fn from_frame(f: &LabelFrame) -> Self {
let p = f.grid_params();
Self {
zoom: f.zoom,
ref_pos: f.ref_pos,
screen_center: f.screen_center,
viewport: p.viewport,
inv_cell: p.inv_cell,
inv_repeat: p.inv_repeat,
lod: p.lod,
ink: f.ink,
halo: f.halo,
}
}
}
/// What the emit pass counted this frame. **Diagnostics, not the draw**: the number
/// the draw uses is `instance_count` in the same buffer, which no CPU code reads on the
/// live path.
#[repr(C)]
#[derive(Clone, Copy, Debug, Default, PartialEq, bytemuck::Pod, bytemuck::Zeroable)]
pub struct LabelStats {
/// Labels that passed the collision rule — **before** the ink budget, so this is
/// the rule's own answer and can exceed [`MAX_VISIBLE_LABELS`]. It doubles as the
/// budget allocator on the device (see `cs_emit`).
pub labels_won: u32,
/// Labels whose glyphs are really in `out_glyphs` — past the ink budget *and* past
/// the capacity guard, so `labels_drawn <= MAX_VISIBLE_LABELS` always.
pub labels_drawn: u32,
/// Labels dropped for want of output capacity — should be 0 in any sane frame.
pub glyph_overflow: u32,
pub _pad: u32,
}
/// The first 8 words of the state buffer: the `DrawIndirectArgs` `draw_indirect` reads
/// at offset 0, immediately followed by the counters. The grid occupies every word
/// after these. Mirrors the `W_*` / `GRID_BASE` constants in `label_collide.wgsl`.
///
/// One buffer, not three, for two reasons that happen to agree. It keeps the compute
/// pass inside `max_storage_buffers_per_shader_stage = 4` — the floor
/// `Limits::downlevel_defaults()` still enforces, and the limit the first draft of this
/// module blew with six separate bindings. And the counters are reset and read in
/// lockstep with the count they gate, so separating them only created a second place to
/// forget.
#[repr(C)]
#[derive(Clone, Copy, Debug, PartialEq, bytemuck::Pod, bytemuck::Zeroable)]
pub struct LabelDrawState {
/// `vertex_count` is the generated quad's 6; `instance_count` is what the emit
/// pass accumulates and the draw consumes.
pub args: DrawIndirectArgs,
pub stats: LabelStats,
}
impl LabelDrawState {
/// The per-frame reset: a 6-vertex quad, zero instances, zero counters.
#[must_use]
pub fn reset() -> Self {
Self {
args: DrawIndirectArgs { vertex_count: 6, instance_count: 0, first_vertex: 0, first_instance: 0 },
stats: LabelStats::default(),
}
}
}
/// Words the draw args + counters occupy before the grid begins (`GRID_BASE`).
pub const STATE_HEADER_WORDS: u32 = 8;
/// The label lane: three compute pipelines over a shared grid, plus the instanced
/// glyph draw they feed through `draw_indirect`.
pub struct LabelCollider {
clear_frame_pipeline: wgpu::ComputePipeline,
clear_round_pipeline: wgpu::ComputePipeline,
claim_pipeline: wgpu::ComputePipeline,
name_bid_pipeline: wgpu::ComputePipeline,
emit_pipeline: wgpu::ComputePipeline,
compute_bgl: wgpu::BindGroupLayout,
draw_pipeline: wgpu::RenderPipeline,
// NO `draw_uniform_bgl` here, and that asymmetry with `compute_bgl`/`atlas_bgl`
// is the point. Those two ARE kept, because their bind groups are built LATE —
// `compute_bind` at `upload` (it needs the candidate count) and the atlas bind
// when an atlas first arrives — so the layout must outlive the constructor. The
// draw uniform has no such moment: `draw_uniform` is a fixed 16 bytes created
// here, `draw_uniform_bind` is created here with it, and the draw binds THAT
// group (`set_bind_group(0, &self.draw_uniform_bind, ..)`). Nothing ever needs
// the layout again. Keeping the field made `dead_code` fire on a struct where
// every other layout is live, which reads as a dropped binding rather than as
// what it is.
atlas_bgl: wgpu::BindGroupLayout,
sampler: wgpu::Sampler,
uniform: wgpu::Buffer,
draw_uniform: wgpu::Buffer,
draw_uniform_bind: wgpu::BindGroup,
/// Draw args + counters + the grid, in one `array<atomic<u32>>` (see
/// [`LabelDrawState`]). `INDIRECT` because `draw_indirect` reads its first 16 bytes.
state: wgpu::Buffer,
out_glyphs: wgpu::Buffer,
/// `None` until [`upload`](Self::upload) has been handed a non-empty candidate
/// set. WGSL's `arrayLength` needs a non-zero binding, and more to the point an
/// empty label set has nothing to collide.
cand_count: u32,
compute_bind: Option<wgpu::BindGroup>,
/// Kept alive for `compute_bind`.
_cand_buf: Option<wgpu::Buffer>,
_glyph_buf: Option<wgpu::Buffer>,
/// `None` until [`set_atlas`](Self::set_atlas). Without it the draw is skipped —
/// the fail-safe: no atlas must mean no labels, never garbage labels.
atlas: Option<(wgpu::Texture, wgpu::BindGroup)>,
atlas_size: [u32; 2],
target_format: wgpu::TextureFormat,
}
impl LabelCollider {
/// Build the pipelines. `target_format` is the colour format the labels composite
/// onto (premultiplied alpha over it).
pub fn new(device: &wgpu::Device, target_format: wgpu::TextureFormat) -> Self {
// ── The three compute passes: ONE module, ONE layout, three entry points ──
// Not three shaders. Clear/claim/emit share every constant, the grid indexing
// and `footprint_of`; splitting them into separate files is the twin LAW #5
// forbids, and worse, it would let the claim pass and the emit pass compute a
// label's cell range two subtly different ways — the exact bug class the whole
// `grid[cell] == priority` scheme cannot survive.
let collide = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("label_collide"),
source: wgpu::ShaderSource::Wgsl(LABEL_COLLIDE_WGSL.into()),
});
let compute_bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: Some("label_collide_bgl"),
entries: &[
storage_entry(0, false), // cands
storage_entry(1, false), // glyph_src
storage_entry(2, true), // state: draw args + counters + grid
storage_entry(3, true), // out_glyphs
uniform_entry(4),
],
});
let compute_pll = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("label_collide_pll"),
bind_group_layouts: &[Some(&compute_bgl)],
immediate_size: 0,
});
let pipe = |entry: &str, label: &str| {
device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: Some(label),
layout: Some(&compute_pll),
module: &collide,
entry_point: Some(entry),
compilation_options: wgpu::PipelineCompilationOptions::default(),
cache: None,
})
};
let clear_frame_pipeline = pipe("cs_clear_frame", "label_clear_frame");
let clear_round_pipeline = pipe("cs_clear_round", "label_clear_round");
let claim_pipeline = pipe("cs_claim", "label_claim");
let name_bid_pipeline = pipe("cs_name_bid", "label_name_bid");
let emit_pipeline = pipe("cs_emit", "label_emit");
// ── The glyph draw ────────────────────────────────────────────────────
let draw = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("label_draw"),
source: wgpu::ShaderSource::Wgsl(LABEL_DRAW_WGSL.into()),
});
let draw_uniform_bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: Some("label_draw_uniform_bgl"),
entries: &[wgpu::BindGroupLayoutEntry {
binding: 0,
visibility: wgpu::ShaderStages::VERTEX_FRAGMENT,
ty: wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Uniform,
has_dynamic_offset: false,
min_binding_size: None,
},
count: None,
}],
});
let atlas_bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: Some("label_atlas_bgl"),
entries: &[
wgpu::BindGroupLayoutEntry {
binding: 0,
visibility: wgpu::ShaderStages::FRAGMENT,
ty: wgpu::BindingType::Texture {
sample_type: wgpu::TextureSampleType::Float { filterable: true },
view_dimension: wgpu::TextureViewDimension::D2,
multisampled: false,
},
count: None,
},
wgpu::BindGroupLayoutEntry {
binding: 1,
visibility: wgpu::ShaderStages::FRAGMENT,
ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering),
count: None,
},
],
});
// Premultiplied source-over: the fragment already multiplied rgb by alpha.
let blend = Some(wgpu::BlendState {
color: wgpu::BlendComponent {
src_factor: wgpu::BlendFactor::One,
dst_factor: wgpu::BlendFactor::OneMinusSrcAlpha,
operation: wgpu::BlendOperation::Add,
},
alpha: wgpu::BlendComponent {
src_factor: wgpu::BlendFactor::One,
dst_factor: wgpu::BlendFactor::OneMinusSrcAlpha,
operation: wgpu::BlendOperation::Add,
},
});
// Offsets are derived from the type, then compile-time asserted, so a layout
// change cannot silently desync the shader (the `LINE_VERTEX_WORDS` idiom).
const _: () = assert!(std::mem::size_of::<LabelGlyphInstance>() == 48);
let attrs = [
wgpu::VertexAttribute { offset: 0, shader_location: 0, format: wgpu::VertexFormat::Float32x2 },
wgpu::VertexAttribute { offset: 8, shader_location: 1, format: wgpu::VertexFormat::Float32x2 },
wgpu::VertexAttribute { offset: 16, shader_location: 2, format: wgpu::VertexFormat::Float32x2 },
wgpu::VertexAttribute { offset: 24, shader_location: 3, format: wgpu::VertexFormat::Float32x2 },
wgpu::VertexAttribute { offset: 32, shader_location: 4, format: wgpu::VertexFormat::Float32x4 },
];
let draw_pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
label: Some("label_draw_pipeline"),
layout: Some(&device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("label_draw_pll"),
bind_group_layouts: &[Some(&draw_uniform_bgl), Some(&atlas_bgl)],
immediate_size: 0,
})),
vertex: wgpu::VertexState {
module: &draw,
entry_point: Some("label_vs"),
compilation_options: wgpu::PipelineCompilationOptions::default(),
buffers: &[Some(wgpu::VertexBufferLayout {
array_stride: std::mem::size_of::<LabelGlyphInstance>() as wgpu::BufferAddress,
step_mode: wgpu::VertexStepMode::Instance,
attributes: &attrs,
})],
},
fragment: Some(wgpu::FragmentState {
module: &draw,
entry_point: Some("label_fs"),
compilation_options: wgpu::PipelineCompilationOptions::default(),
targets: &[Some(wgpu::ColorTargetState { format: target_format, blend, write_mask: wgpu::ColorWrites::ALL })],
}),
primitive: wgpu::PrimitiveState { topology: wgpu::PrimitiveTopology::TriangleList, ..Default::default() },
depth_stencil: None,
multisample: crate::render::gpu::msaa_state(),
multiview_mask: None,
cache: None,
});
let sampler = device.create_sampler(&wgpu::SamplerDescriptor {
label: Some("label_atlas_sampler"),
mag_filter: wgpu::FilterMode::Linear,
min_filter: wgpu::FilterMode::Linear,
address_mode_u: wgpu::AddressMode::ClampToEdge,
address_mode_v: wgpu::AddressMode::ClampToEdge,
..Default::default()
});
let uniform = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("label_uniform"),
size: std::mem::size_of::<LabelUniforms>() as u64,
usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
mapped_at_creation: false,
});
let draw_uniform = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("label_draw_uniform"),
size: 16,
usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
mapped_at_creation: false,
});
let draw_uniform_bind = device.create_bind_group(&wgpu::BindGroupDescriptor {
label: Some("label_draw_uniform_bind"),
layout: &draw_uniform_bgl,
entries: &[wgpu::BindGroupEntry { binding: 0, resource: draw_uniform.as_entire_binding() }],
});
let state = Self::make_state(device, 0);
let out_glyphs = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("label_out_glyphs"),
size: u64::from(LABEL_GLYPH_CAPACITY) * std::mem::size_of::<LabelGlyphInstance>() as u64,
// COPY_SRC is for the proof only; the live draw reads this as VERTEX, on
// the device, with the count the compute pass wrote.
usage: wgpu::BufferUsages::VERTEX | wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC,
mapped_at_creation: false,
});
Self {
clear_frame_pipeline,
clear_round_pipeline,
claim_pipeline,
name_bid_pipeline,
emit_pipeline,
compute_bgl,
draw_pipeline,
atlas_bgl,
sampler,
uniform,
draw_uniform,
draw_uniform_bind,
state,
out_glyphs,
cand_count: 0,
compute_bind: None,
_cand_buf: None,
_glyph_buf: None,
atlas: None,
atlas_size: [0, 0],
target_format,
}
}
/// The state buffer: header + the four grid regions + one word per candidate.
///
/// Sized at upload rather than construction because the per-label region needs the
/// candidate count. `INDIRECT` because `draw_indirect` reads its first 16 bytes.
fn make_state(device: &wgpu::Device, cands: u32) -> wgpu::Buffer {
device.create_buffer(&wgpu::BufferDescriptor {
label: Some("label_state"),
size: u64::from(STATE_HEADER_WORDS + GRID_WORDS + cands.max(1)) * 4,
usage: wgpu::BufferUsages::INDIRECT
| wgpu::BufferUsages::STORAGE
| wgpu::BufferUsages::COPY_DST
| wgpu::BufferUsages::COPY_SRC,
mapped_at_creation: false,
})
}
/// The colour format the label draw targets.
#[must_use]
pub fn target_format(&self) -> wgpu::TextureFormat {
self.target_format
}
/// Candidates currently uploaded.
#[must_use]
pub fn candidate_count(&self) -> u32 {
self.cand_count
}
/// Is the lane able to draw? (candidates uploaded AND an atlas bound)
#[must_use]
pub fn ready(&self) -> bool {
self.cand_count > 0 && self.atlas.is_some()
}
/// **Upload the label set.** Called when the visible labels change — *not* every
/// frame: the candidates carry origin-local Mercator anchors, so pan and zoom are
/// a uniform update and the text layout is not redone.
///
/// An empty set clears the lane rather than leaving the previous one bound, so a
/// pane that stops offering labels stops drawing them.
pub fn upload(
&mut self,
device: &wgpu::Device,
_queue: &wgpu::Queue,
cands: &[LabelCandidate],
glyphs: &[GlyphSrc],
) {
if cands.is_empty() || glyphs.is_empty() {
self.cand_count = 0;
self.compute_bind = None;
self._cand_buf = None;
self._glyph_buf = None;
return;
}
let cand_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("label_cands"),
contents: bytemuck::cast_slice(cands),
usage: wgpu::BufferUsages::STORAGE,
});
let glyph_buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
label: Some("label_glyph_src"),
contents: bytemuck::cast_slice(glyphs),
usage: wgpu::BufferUsages::STORAGE,
});
// The per-label state region must fit this candidate set.
self.state = Self::make_state(device, cands.len() as u32);
self.compute_bind = Some(device.create_bind_group(&wgpu::BindGroupDescriptor {
label: Some("label_collide_bind"),
layout: &self.compute_bgl,
entries: &[
wgpu::BindGroupEntry { binding: 0, resource: cand_buf.as_entire_binding() },
wgpu::BindGroupEntry { binding: 1, resource: glyph_buf.as_entire_binding() },
wgpu::BindGroupEntry { binding: 2, resource: self.state.as_entire_binding() },
wgpu::BindGroupEntry { binding: 3, resource: self.out_glyphs.as_entire_binding() },
wgpu::BindGroupEntry { binding: 4, resource: self.uniform.as_entire_binding() },
],
}));
self._cand_buf = Some(cand_buf);
self._glyph_buf = Some(glyph_buf);
self.cand_count = cands.len() as u32;
}
/// Upload the glyph atlas — premultiplied `Rgba8Unorm` rows, `size[0] * size[1] *
/// 4` bytes. Feed it `egui::Context::fonts(|f| f.image())`: the same atlas the CPU
/// text pass samples, so the two lanes letter the map identically.
pub fn set_atlas(
&mut self,
device: &wgpu::Device,
queue: &wgpu::Queue,
size: [u32; 2],
rgba: &[u8],
) {
let (w, h) = (size[0].max(1), size[1].max(1));
let need = (w * h * 4) as usize;
if rgba.len() < need {
return; // refuse a short buffer rather than upload garbage
}
let tex = device.create_texture(&wgpu::TextureDescriptor {
label: Some("label_atlas"),
size: wgpu::Extent3d { width: w, height: h, depth_or_array_layers: 1 },
mip_level_count: 1,
sample_count: crate::render::gpu::NO_MSAA_SAMPLES,
dimension: wgpu::TextureDimension::D2,
format: wgpu::TextureFormat::Rgba8Unorm,
usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::COPY_DST,
view_formats: &[],
});
queue.write_texture(
wgpu::TexelCopyTextureInfo {
texture: &tex,
mip_level: 0,
origin: wgpu::Origin3d::ZERO,
aspect: wgpu::TextureAspect::All,
},
&rgba[..need],
wgpu::TexelCopyBufferLayout { offset: 0, bytes_per_row: Some(w * 4), rows_per_image: Some(h) },
wgpu::Extent3d { width: w, height: h, depth_or_array_layers: 1 },
);
let view = tex.create_view(&wgpu::TextureViewDescriptor::default());
let bind = device.create_bind_group(&wgpu::BindGroupDescriptor {
label: Some("label_atlas_bind"),
layout: &self.atlas_bgl,
entries: &[
wgpu::BindGroupEntry { binding: 0, resource: wgpu::BindingResource::TextureView(&view) },
wgpu::BindGroupEntry { binding: 1, resource: wgpu::BindingResource::Sampler(&self.sampler) },
],
});
self.atlas = Some((tex, bind));
self.atlas_size = [w, h];
}
/// The atlas dimensions currently bound (`[0, 0]` = none).
#[must_use]
pub fn atlas_size(&self) -> [u32; 2] {
self.atlas_size
}
/// **Run the three passes.** Publishes this frame's camera + palette, resets the
/// draw args and the counters, then dispatches clear → claim → emit.
///
/// One compute pass with the pipeline switched between dispatches, exactly as
/// `run_cull` does. The switch is the barrier that makes the rule
/// order-independent: every claim must land before any emit reads a cell.
///
/// Kept as its own public method (rather than inlined into a paint callback) for
/// the reason `run_cull` was lifted out of `prepare` — so the proof drives the
/// REAL dispatch instead of a re-spelled copy, which is how the cull went untested
/// for as long as it did.
pub fn run(&self, queue: &wgpu::Queue, encoder: &mut wgpu::CommandEncoder, frame: &LabelFrame) {
// `vertex_count: 6` is the generated quad; `instance_count: 0` is what the
// emit pass's `atomicAdd` accumulates onto and what `draw_indirect` reads. The
// GRID is not reset here — `cs_clear` owns that, and the proof runs three
// frames to catch it being skipped.
queue.write_buffer(&self.state, 0, bytemuck::bytes_of(&LabelDrawState::reset()));
queue.write_buffer(&self.uniform, 0, bytemuck::bytes_of(&LabelUniforms::from_frame(frame)));
let p = frame.grid_params();
queue.write_buffer(&self.draw_uniform, 0, bytemuck::cast_slice(&[p.viewport[0], p.viewport[1], 0.0, 0.0]));
let Some(bind) = &self.compute_bind else { return };
let mut cpass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
label: Some("label_collide"),
timestamp_writes: None,
});
cpass.set_bind_group(0, bind, &[]);
// One clear of everything (including the per-label state words), then the round
// loop. Each round is three dispatches, and the dispatch boundaries ARE the
// barriers: claim is the only pass that reads occupancy, emit the only one that
// writes it, so no invocation can observe a half-written region. A single kernel
// with `workgroupBarrier` could not give that — the grid spans far more cells
// than one workgroup.
let labels = self.cand_count.div_ceil(64);
cpass.set_pipeline(&self.clear_frame_pipeline);
cpass.dispatch_workgroups((GRID_WORDS + self.cand_count).div_ceil(64), 1, 1);
for round in 0..LABEL_ROUNDS {
if round > 0 {
// Round 0 needs no per-round clear: the frame clear just ran.
cpass.set_pipeline(&self.clear_round_pipeline);
cpass.dispatch_workgroups((GRID_CELLS + NAME_SLOTS).div_ceil(64), 1, 1);
}
cpass.set_pipeline(&self.claim_pipeline);
cpass.dispatch_workgroups(labels, 1, 1);
cpass.set_pipeline(&self.name_bid_pipeline);
cpass.dispatch_workgroups(labels, 1, 1);
cpass.set_pipeline(&self.emit_pipeline);
cpass.dispatch_workgroups(labels, 1, 1);
}
}
/// **The indirect label draw.** The instance count comes out of the buffer the
/// emit pass wrote; nothing here knows how many labels survived.
pub fn draw_indirect(&self, pass: &mut wgpu::RenderPass<'_>) {
let Some((_, atlas_bind)) = &self.atlas else { return };
if self.cand_count == 0 {
return;
}
pass.set_pipeline(&self.draw_pipeline);
pass.set_bind_group(0, &self.draw_uniform_bind, &[]);
pass.set_bind_group(1, atlas_bind, &[]);
pass.set_vertex_buffer(0, self.out_glyphs.slice(..));
pass.draw_indirect(&self.state, 0);
}
/// Record the label draw into an offscreen `target` — the self-contained form the
/// proof and a non-egui host use. `load` keeps what is already there.
#[allow(clippy::too_many_arguments)]
pub fn render_to(
&self,
encoder: &mut wgpu::CommandEncoder,
target: &wgpu::TextureView,
load: bool,
) {
let load_op =
if load { wgpu::LoadOp::Load } else { wgpu::LoadOp::Clear(wgpu::Color::TRANSPARENT) };
let mut rp = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
label: Some("label_draw_pass"),
color_attachments: &[Some(wgpu::RenderPassColorAttachment {
view: target,
resolve_target: None,
depth_slice: None,
ops: wgpu::Operations { load: load_op, store: wgpu::StoreOp::Store },
})],
depth_stencil_attachment: None,
timestamp_writes: None,
occlusion_query_set: None,
multiview_mask: None,
});
self.draw_indirect(&mut rp);
}
// ── Test hooks ────────────────────────────────────────────────────────────
//
// The live path never reads any of these. The whole point of item 4, like item 2,
// is that the count is produced and consumed on the device; copying it to the CPU
// here is how the proof observes what the draw saw.
/// Read back the draw args + counters — `args.instance_count` is the number the
/// draw uses, out of the very buffer and offset `draw_indirect` reads.
#[doc(hidden)]
pub fn state_readback(&self, device: &wgpu::Device, queue: &wgpu::Queue) -> LabelDrawState {
let bytes = read_buffer(device, queue, &self.state, std::mem::size_of::<LabelDrawState>() as u64);
*bytemuck::from_bytes(&bytes)
}
/// Read back the first `n` emitted glyph instances — the bytes the vertex shader
/// will actually pull.
#[doc(hidden)]
pub fn glyphs_readback(
&self,
device: &wgpu::Device,
queue: &wgpu::Queue,
n: u32,
) -> Vec<LabelGlyphInstance> {
let stride = std::mem::size_of::<LabelGlyphInstance>() as u64;
let n = n.min(LABEL_GLYPH_CAPACITY);
if n == 0 {
return Vec::new();
}
let bytes = read_buffer(device, queue, &self.out_glyphs, u64::from(n) * stride);
bytemuck::cast_slice::<u8, LabelGlyphInstance>(&bytes).to_vec()
}
}
/// Blocking buffer→CPU copy (the `line_draw_counts` idiom).
fn read_buffer(device: &wgpu::Device, queue: &wgpu::Queue, src: &wgpu::Buffer, size: u64) -> Vec<u8> {
let staging = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("label_readback"),
size,
usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ,
mapped_at_creation: false,
});
let mut enc = device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
enc.copy_buffer_to_buffer(src, 0, &staging, 0, size);
queue.submit(Some(enc.finish()));
let slice = staging.slice(..);
let (tx, rx) = std::sync::mpsc::channel();
slice.map_async(wgpu::MapMode::Read, move |r| {
let _ = tx.send(r);
});
device.poll(wgpu::PollType::wait_indefinitely()).ok();
let _ = rx.recv();
let out = slice.get_mapped_range().expect("wgpu: get_mapped_range on a mapped buffer").to_vec();
staging.unmap();
out
}
fn storage_entry(binding: u32, read_write: bool) -> wgpu::BindGroupLayoutEntry {
wgpu::BindGroupLayoutEntry {
binding,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Storage { read_only: !read_write },
has_dynamic_offset: false,
min_binding_size: None,
},
count: None,
}
}
fn uniform_entry(binding: u32) -> wgpu::BindGroupLayoutEntry {
wgpu::BindGroupLayoutEntry {
binding,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::Buffer {
ty: wgpu::BufferBindingType::Uniform,
has_dynamic_offset: false,
min_binding_size: None,
},
count: None,
}
}
#[cfg(test)]
mod tests {
use super::*;
use crate::label_grid::{
label_half_extent, name_hash, priority, repeat_slot, resolve, screen_labels, FLAG_BLOCKER,
FLAG_SCREEN_SPACE, GRID_H, GRID_W, REPEAT_CELL_PX,
};
/// The uniform block is 80 bytes with `ink`/`halo` at 16-aligned offsets, which is
/// what WGSL demands of a `vec4` in a uniform struct. A mismatch here does not
/// produce a validation error — it silently shifts every field the shader reads.
#[test]
fn uniform_layout_matches_the_wgsl_struct() {
assert_eq!(std::mem::size_of::<LabelUniforms>(), 80);
assert_eq!(std::mem::offset_of!(LabelUniforms, ink), 48);
assert_eq!(std::mem::offset_of!(LabelUniforms, halo), 64);
assert_eq!(std::mem::size_of::<LabelStats>(), 16);
// The state header must be exactly the 8 words `GRID_BASE` skips, with the
// draw args first — `draw_indirect` reads offset 0 of this same buffer.
assert_eq!(std::mem::size_of::<LabelDrawState>(), (STATE_HEADER_WORDS * 4) as usize);
assert_eq!(std::mem::offset_of!(LabelDrawState, args), 0);
assert_eq!(std::mem::offset_of!(LabelDrawState, stats), 16);
assert_eq!(LabelDrawState::reset().args.vertex_count, 6, "the quad is 6 vertices");
assert_eq!(LabelDrawState::reset().args.instance_count, 0, "the count starts at zero");
assert!(LABEL_COLLIDE_WGSL.contains("const GRID_BASE: u32 = 8u;"), "GRID_BASE must match STATE_HEADER_WORDS");
assert!(LABEL_COLLIDE_WGSL.contains("fn cs_clear"));
assert!(LABEL_COLLIDE_WGSL.contains("fn cs_claim"));
assert!(LABEL_COLLIDE_WGSL.contains("fn cs_emit"));
assert!(LABEL_DRAW_WGSL.contains("fn label_vs"));
assert!(LABEL_DRAW_WGSL.contains("fn label_fs"));
}
/// A headless device — **and the adapter it came from, printed**.
///
/// The adapter line is not decoration. Every test in this module opens with
/// `let Some(..) = headless_device() else { skip }`, and a graceful skip is exactly
/// how this file goes GREEN having run no shader at all: item 8's agent printed its
/// adapter for the same reason. If a FAILING run's replayed output (or a
/// `--nocapture` one) carries no `[label_collide_test] adapter — …` line, nothing
/// below it was proven on a device, whatever the test result says. libtest swallows
/// stderr otherwise, which is why the turn also REFUSES a software adapter rather
/// than trusting anyone to read the line.
fn headless_device(
) -> Option<(wgpu::Device, wgpu::Queue, String, crate::render::gputurn::probe::ProbeGuard)> {
// ONE WRITER for the bring-up — `crate::render::gputurn::probe::open`. This used to
// be twenty hand-rolled lines here and in seven other places; every copy asked
// `request_adapter(PowerPreference::default())` and so could be handed llvmpipe with
// no word said, and none of them took the GPU turn. `open` takes the turn, goes
// through facett's adapter policy (software ranked last), PRINTS the adapter it got,
// and refuses a software one. The turn travels out in the returned guard so it covers
// the whole test body — MEASURED 2026-08-31, serialising only the bring-up still left
// nine `label_collide` devices alive at once and deadlocked the binary for 13 minutes
// (20 threads in `futex_do_wait`, GPU at 0 %, 632 MiB held), in BOTH profiles; the
// same 74 `render::gpu::*` tests pass in 5.2 s serialised. Bind the guard, never `_`.
use crate::render::gputurn::probe;
// This file's tests print the adapter into their OWN assertion messages, so the
// name comes back out as well as going to stderr.
let p = probe::open("label_collide_test", probe::OnSoftware::Refuse, probe::downlevel)?;
let name = p.named.clone();
let (device, queue, guard) = p.split();
Some((device, queue, name, guard))
}
fn frame(viewport: [f32; 2], lod: u32) -> LabelFrame {
LabelFrame {
zoom: [1.0, 1.0],
ref_pos: [0.0, 0.0],
screen_center: [0.0, 0.0],
viewport,
lod,
ink: [0.9, 0.9, 0.95, 1.0],
halo: [0.02, 0.03, 0.05, 1.0],
}
}
/// A label at screen `pos`, `glyphs` glyphs wide, priority `rank`/`order`.
///
/// The glyph count VARIES per label on purpose. `instance_count` is
/// `5 * sum(glyph_count of the winners)`, so distinct counts make the number the
/// GPU produced identify *which* labels survived, not merely how many — the
/// `banded_roads` trick from item 2's proof, which is the difference between a
/// count assertion that means something and one that passes on the wrong set.
fn cand(pos: [f32; 2], glyphs: u32, glyph_start: u32, rank: u8, order: u32, name: &str) -> LabelCandidate {
LabelCandidate {
pos,
half_px: label_half_extent([glyphs as f32 * 7.0, 13.0]),
priority: priority(rank, order),
name_hash: name_hash(name),
glyph_start,
glyph_count: glyphs,
lod: 0,
flags: FLAG_SCREEN_SPACE,
_pad: [0; 2],
}
}
/// Glyph templates for a label of `n` glyphs, laid out left-to-right about the
/// centre. Non-degenerate UVs so a readback can tell a written instance from a
/// zeroed one — and `uv_min.y` encodes `label_id`, so the emitted instances say
/// WHICH labels drew, not merely how many.
///
/// That encoding is what lets the parity test compare the surviving SET. Comparing
/// `instance_count` alone only works while every label has a distinct glyph count,
/// which stops being true the moment the fixture is dense enough to exercise the
/// round loop.
fn glyph_run(n: u32, label_id: u32) -> Vec<GlyphSrc> {
let w = 7.0;
let span = n as f32 * w;
(0..n)
.map(|i| {
let x = -span * 0.5 + i as f32 * w;
GlyphSrc {
off_min: [x, -6.5],
off_max: [x + w, 6.5],
uv_min: [0.1 + 0.01 * i as f32, LABEL_UV_BASE + label_id as f32 * LABEL_UV_STEP],
uv_max: [0.1 + 0.01 * i as f32 + 0.008, 0.9],
}
})
.collect()
}
/// `uv_min.y` origin and per-label step used by [`glyph_run`] to tag a label.
const LABEL_UV_BASE: f32 = 0.01;
const LABEL_UV_STEP: f32 = 0.001;
/// Recover the set of label indices present in an emitted instance batch.
fn drawn_ids(glyphs: &[LabelGlyphInstance]) -> std::collections::BTreeSet<u32> {
glyphs
.iter()
.map(|g| ((g.uv_min[1] - LABEL_UV_BASE) / LABEL_UV_STEP).round() as u32)
.collect()
}
/// Build a candidate set + its concatenated glyph templates from `(pos, glyphs,
/// rank, order, name)` tuples, wiring `glyph_start` so no label reads another's
/// glyphs.
fn build(specs: &[([f32; 2], u32, u8, u32, &str)]) -> (Vec<LabelCandidate>, Vec<GlyphSrc>) {
let mut cands = Vec::new();
let mut glyphs = Vec::new();
for (pos, n, rank, order, name) in specs {
let start = glyphs.len() as u32;
glyphs.extend(glyph_run(*n, cands.len() as u32));
cands.push(cand(*pos, *n, start, *rank, *order, name));
}
(cands, glyphs)
}
/// Drive the REAL `run` and report what the draw would consume.
fn collide(
c: &LabelCollider,
device: &wgpu::Device,
queue: &wgpu::Queue,
f: &LabelFrame,
) -> (DrawIndirectArgs, LabelStats) {
let mut enc =
device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("label_test") });
c.run(queue, &mut enc, f);
queue.submit(Some(enc.finish()));
device.poll(wgpu::PollType::wait_indefinitely()).ok();
let st = c.state_readback(device, queue);
(st.args, st.stats)
}
fn collider(device: &wgpu::Device, queue: &wgpu::Queue, cands: &[LabelCandidate], glyphs: &[GlyphSrc]) -> LabelCollider {
let mut c = LabelCollider::new(device, wgpu::TextureFormat::Rgba8Unorm);
c.upload(device, queue, cands, glyphs);
// A tiny opaque atlas so `ready()` and the draw path are exercised.
c.set_atlas(device, queue, [4, 4], &[255u8; 4 * 4 * 4]);
c
}
/// **THE PROOF.** The instance count `draw_indirect` consumes is computed by the
/// collision pass, read back out of the very buffer the draw reads, and equals —
/// exactly — five instances per glyph of the labels that should have survived.
///
/// Every arm is on the number the DEVICE produced. Not `ready()`, not a CPU-side
/// prediction, not "some labels were drawn".
#[test]
fn gpu_collision_drives_the_indirect_label_instance_count() {
let Some((device, queue, _adapter, _gpu)) = headless_device() else {
eprintln!("[labels] no GPU adapter — skipping (CPU-only CI)");
return;
};
let f = frame([1280.0, 768.0], 2);
let per = INSTANCES_PER_GLYPH;
// ── 1. NON-VACUITY: nothing overlaps ⇒ every label draws ──────────────
// First, because a pass that culls everything would satisfy every "something
// was removed" assertion below.
let (cands, glyphs) =
build(&[([120.0, 100.0], 4, 5, 0, "a"), ([600.0, 100.0], 7, 5, 1, "b"), ([1100.0, 600.0], 11, 5, 2, "c")]);
let sizes: Vec<u32> = cands.iter().map(|c| c.glyph_count).collect();
assert!(
sizes.iter().collect::<std::collections::HashSet<_>>().len() == sizes.len(),
"the labels must have DISTINCT glyph counts, else the count cannot say WHICH survived: {sizes:?}"
);
let c = collider(&device, &queue, &cands, &glyphs);
let (args, stats) = collide(&c, &device, &queue, &f);
let all: u32 = sizes.iter().sum::<u32>() * per;
assert_eq!(args.instance_count, all, "no overlap ⇒ all {} labels draw ({all} instances)", sizes.len());
assert_eq!(args.vertex_count, 6, "the quad is 6 generated vertices");
assert_eq!(stats.labels_drawn, 3);
assert_eq!(stats.glyph_overflow, 0);
// ── 2. OVERLAP: N labels on one spot ⇒ exactly ONE survives, the best ──
// The ranks and the glyph counts are BOTH distinct, so the count identifies
// the survivor. A pass that kept "some" one would produce 4·5, 7·5 or 11·5;
// only the top-ranked one produces 9·5.
let (cands, glyphs) = build(&[
([400.0, 300.0], 4, 3, 0, "p"),
([403.0, 302.0], 9, 8, 1, "q"), // highest rank → the expected survivor
([398.0, 299.0], 6, 5, 2, "r"),
([401.0, 301.0], 13, 1, 3, "s"),
]);
let c2 = collider(&device, &queue, &cands, &glyphs);
let (args2, stats2) = collide(&c2, &device, &queue, &f);
assert_eq!(
args2.instance_count,
9 * per,
"one survivor and it is the RANK-8 label (9 glyphs → {} instances); got {} — \
4/6/13 glyphs would mean the wrong label won, 0 that all were culled",
9 * per,
args2.instance_count
);
assert_eq!(stats2.labels_won, 1, "exactly one label passed the rule");
assert!(args2.instance_count < all, "culling must actually REMOVE something");
// ── 3. THE EMPTY ANCHOR: nothing on screen ⇒ nothing drawn ────────────
let off = frame([1280.0, 768.0], 2);
let (cands, glyphs) = build(&[([-5000.0, -5000.0], 8, 5, 0, "gone")]);
let c3 = collider(&device, &queue, &cands, &glyphs);
let (args3, stats3) = collide(&c3, &device, &queue, &off);
assert_eq!(args3.instance_count, 0, "an off-screen label must draw NOTHING — got {}", args3.instance_count);
assert_eq!(stats3.labels_won, 0);
// ── 4. LOD is applied on the device ───────────────────────────────────
let (mut cands, glyphs) = build(&[([200.0, 200.0], 5, 5, 0, "country"), ([800.0, 200.0], 9, 5, 1, "city")]);
cands[1].lod = 2;
let c4 = collider(&device, &queue, &cands, &glyphs);
let (at_city, _) = collide(&c4, &device, &queue, &frame([1280.0, 768.0], 2));
let (at_country, _) = collide(&c4, &device, &queue, &frame([1280.0, 768.0], 0));
assert_eq!(at_city.instance_count, (5 + 9) * per, "at city LOD both draw");
assert_eq!(
at_country.instance_count,
5 * per,
"at country LOD the city-tier label drops and the OTHER stays — got {}",
at_country.instance_count
);
assert!(at_country.instance_count > 0, "the LOD filter must not drop everything (that would pass arm 3 vacuously)");
// ── 5. BLOCKERS reserve pixels and print nothing ──────────────────────
let (mut cands, glyphs) =
build(&[([500.0, 400.0], 3, 0xFF, 0, "#pin"), ([505.0, 402.0], 7, 8, 1, "over"), ([1000.0, 120.0], 11, 8, 2, "clear")]);
cands[0].flags |= FLAG_BLOCKER;
let c5 = collider(&device, &queue, &cands, &glyphs);
let (args5, _) = collide(&c5, &device, &queue, &f);
assert_eq!(
args5.instance_count,
11 * per,
"only the label away from the pin draws (11 glyphs); the blocker prints nothing \
and the label over it is suppressed — got {}",
args5.instance_count
);
// ── 6. THE RESIDUE TRAP ───────────────────────────────────────────────
// `cs_clear` is the only thing that empties the grid between frames; wgpu
// zero-initialises the buffer, so frame 1 is correct with or without it.
//
// THE FIRST VERSION OF THIS ARM WAS HOLLOW: it re-ran the SAME frame three
// times and asserted the count was stable. It passed with `cs_clear`'s body
// deleted, and had to — replaying an identical frame re-claims the same cells
// with the same priorities, so the stale values it reads back are exactly the
// ones it would have written. Residue is invisible unless the WINNERS CHANGE.
//
// So: two labels overlapping, the strong one city-tier and the weak one
// country-tier, resolved at one LOD and then the other. Frame 1 (city LOD) is
// won by the 11-glyph label; frame 2 (country LOD) drops it from the candidate
// set entirely, leaving the 4-glyph label — which can only win if frame 1's
// higher priority is GONE from those cells.
let (cands6, glyphs6) = build(&[([400.0, 300.0], 11, 9, 0, "strong"), ([402.0, 301.0], 4, 3, 1, "weak")]);
let mut cands6 = cands6;
cands6[0].lod = 2; // only a candidate at city LOD
cands6[1].lod = 0;
let c6 = collider(&device, &queue, &cands6, &glyphs6);
let (city, _) = collide(&c6, &device, &queue, &frame([1280.0, 768.0], 2));
assert_eq!(city.instance_count, 11 * per, "at city LOD the strong label wins");
let (country, _) = collide(&c6, &device, &queue, &frame([1280.0, 768.0], 0));
assert_eq!(
country.instance_count,
4 * per,
"at country LOD the weak label must win the cells the strong one held LAST \
frame — got {}; 0 means the grid was never cleared and it lost to a ghost",
country.instance_count
);
// …and back again, so the residue cannot be one-directional.
let (city2, _) = collide(&c6, &device, &queue, &frame([1280.0, 768.0], 2));
assert_eq!(city2.instance_count, 11 * per, "and back at city LOD the strong label wins again");
// ── 7. The emitted BYTES, not just the count ──────────────────────────
// `atomicAdd` accumulates whatever the emit pass claims, so a wrong stride or
// a wrong offset yields the RIGHT count over garbage. This reads what the
// vertex shader will pull.
let got = c.glyphs_readback(&device, &queue, all);
assert_eq!(got.len() as u32, all, "read back exactly the instances the draw consumes");
assert!(
got.iter().all(|g| g.rect_max[0] > g.rect_min[0] && g.rect_max[1] > g.rect_min[1]),
"every emitted quad has positive extent (a zeroed instance would not)"
);
assert!(
got.iter().all(|g| g.uv_max[0] > g.uv_min[0] && g.uv_min[0] >= 0.1),
"every instance carries its template's atlas UV, not zeros"
);
// The halo/ink split: exactly one instance in five is ink, and the halo taps
// sit ±1 px from it. A halo drawn at the same offset as the ink is invisible.
let ink = got.iter().filter(|g| g.color == f.ink).count() as u32;
let halo = got.iter().filter(|g| g.color == f.halo).count() as u32;
assert_eq!(ink, all / per, "one ink instance per glyph ({} of {all})", all / per);
assert_eq!(halo, all - all / per, "and {} halo taps", all - all / per);
// Every ink instance must have four halo taps at ±1 px around it.
let mut checked = 0;
for g in got.iter().filter(|g| g.color == f.ink) {
let neighbours = got
.iter()
.filter(|h| {
h.color == f.halo
&& h.uv_min == g.uv_min
&& (h.rect_min[0] - g.rect_min[0]).abs() + (h.rect_min[1] - g.rect_min[1]).abs() == 1.0
})
.count();
assert_eq!(neighbours, 4, "each ink glyph is haloed by 4 taps at ±1 px, found {neighbours}");
checked += 1;
}
assert!(checked >= 3, "the halo check actually ran over several glyphs, not zero");
eprintln!(
"[labels] GPU-computed instance_count — all={} overlap={} empty={} country-lod={} \
blocker={} (per-label glyphs {sizes:?}, {per} instances/glyph); {} instances match template UVs",
args.instance_count,
args2.instance_count,
args3.instance_count,
at_country.instance_count,
args5.instance_count,
got.len()
);
crate::testmatrix::emit(
"facett-core",
"gpu_collision_drives_the_indirect_label_instance_count",
args.instance_count == all
&& args2.instance_count == 9 * per
&& args3.instance_count == 0
&& at_country.instance_count == 5 * per,
&format!(
"indirect instance_count: all={} overlap={} empty={} country={} blocker={}",
args.instance_count, args2.instance_count, args3.instance_count, at_country.instance_count, args5.instance_count
),
);
}
/// **The parity guard — ONE rule, two executions.** The device and
/// [`resolve`](crate::label_grid::resolve) keep the SAME labels over a fixture
/// dense enough to exercise overlap, repeats, blockers and the LOD gate.
///
/// This is the guard that stops the fail-safe CPU painter drifting from the GPU
/// lane. It compares the surviving SET, reconstructed from `instance_count` via
/// distinct glyph counts, rather than a bare total — with 4/7/11/13/17 glyphs the
/// sum is unique to the subset, so an agreement on the number is an agreement on
/// the labels.
#[test]
fn cpu_and_gpu_lanes_resolve_the_same_labels() {
let Some((device, queue, _adapter, _gpu)) = headless_device() else {
eprintln!("[labels] no GPU adapter — skipping parity (CPU-only CI)");
return;
};
let f = frame([1280.0, 768.0], 1);
// Distinct powers-of-two-ish glyph counts ⇒ the sum identifies the subset.
let (mut cands, glyphs) = build(&[
([150.0, 90.0], 4, 6, 0, "Feldkircher Strasse"),
([420.0, 92.0], 7, 6, 1, "Feldkircher Strasse"), // repeat, lower priority
([430.0, 300.0], 11, 9, 2, "Planknerstrasse"),
([437.0, 305.0], 13, 4, 3, "Herrengasse"), // overlaps the above, loses
([1150.0, 700.0], 17, 7, 4, "Bahnhofstrasse"),
([700.0, 500.0], 5, 0xFF, 5, "#pin"), // blocker
([706.0, 503.0], 23, 8, 6, "Marktgasse"), // under the pin
([900.0, 220.0], 19, 8, 7, "Stadtle"),
]);
cands[5].flags |= FLAG_BLOCKER;
cands[7].lod = 2; // city tier, above the frame's LOD 1 → not a candidate
// CPU lane, through the SHARED rule.
let screen = screen_labels(&cands, &f);
let kept = resolve(&screen, &f.grid_params());
let cpu_instances: u32 =
kept.iter().map(|&i| cands[i as usize].glyph_count).sum::<u32>() * INSTANCES_PER_GLYPH;
// GPU lane.
let c = collider(&device, &queue, &cands, &glyphs);
let (args, stats) = collide(&c, &device, &queue, &f);
// Non-vacuity of the FIXTURE: it must genuinely cull, and genuinely keep.
assert!(!kept.is_empty(), "the fixture keeps something");
assert!(kept.len() < cands.len(), "the fixture culls something ({} of {} kept)", kept.len(), cands.len());
assert_eq!(
args.instance_count, cpu_instances,
"the device and the shared rule must keep the SAME labels — CPU kept {:?} \
({cpu_instances} instances), the GPU's draw args say {}",
kept, args.instance_count
);
assert_eq!(stats.labels_drawn as usize, kept.len(), "…and the same COUNT of labels");
eprintln!(
"[labels] parity — CPU kept {:?} = {cpu_instances} instances; GPU instance_count={} labels_drawn={}",
kept, args.instance_count, stats.labels_drawn
);
crate::testmatrix::emit(
"facett-core",
"cpu_and_gpu_lanes_resolve_the_same_labels",
args.instance_count == cpu_instances && stats.labels_drawn as usize == kept.len(),
&format!("cpu={cpu_instances} gpu={} labels={}", args.instance_count, stats.labels_drawn),
);
}
// ── The DENSE arms ─────────────────────────────────────────────────────────────
//
// Everything above runs on 3 to 8 labels, and the parity fixture — the densest of
// them — is 8. Eight labels on a 256x192 grid CANNOT CONTEND IT: `atomicMax` is
// handed one bidder per cell almost everywhere, so the arms above cannot tell a
// working priority race from a pass that simply draws whatever it is given. Nor can
// they tell a working winner from the WRONG winner once the fixture is big enough
// that several subsets sum to the same `instance_count`.
//
// So: 320 candidates in 40 tight knots of 8 mutually-overlapping rivals, and the
// readback says WHICH labels drew rather than only how many.
//
// The identity channel is `uv_min.y`, tagged per label by `glyph_run` and decoded by
// `drawn_ids` — read out of `out_glyphs`, which is the VERTEX buffer the draw pulls
// from, not a side table. `LABEL_ID_ROUND_TRIPS` proves that channel is injective
// over the whole fixture first, because an identity assertion over a lossy encoding
// is worse than a count.
//
// ⚠ THE HAZARD THIS FIXTURE IS SHAPED AROUND, and it is item 8's hazard exactly.
// Its cross-lane test passed with the depth mapping DELETED, because submission
// order already equalled intended order — when the order you are testing is the
// order you submitted, the mapping is unobservable. Label priority has the same
// shape: `priority(rank, order)` inverts `order` so a LOWER index wins, so a fixture
// that gives every rival the same rank makes "highest priority" and "submitted
// first" the same label, and `atomicMax` becomes unobservable. Every knot therefore
// draws its ranks from a SHUFFLED permutation, the test asserts the winning slot
// varies across knots, and it asserts outright that priority order is not submission
// order.
/// Knots across / down and rivals per knot: 8 x 5 x 8 = 320 candidates, 40 winners.
///
/// 40 stays under [`MAX_VISIBLE_LABELS`], deliberately: above the ink budget the
/// survivors are whoever `atomicAdd` served first, which is nondeterministic and
/// could not carry an exact identity assertion. The budget itself is proven by
/// [`a_dense_field_with_no_overlap_culls_nothing`], which goes over it on purpose.
const KNOT_COLS: usize = 8;
const KNOT_ROWS: usize = 5;
const KNOT_RIVALS: usize = 8;
const KNOTS: usize = KNOT_COLS * KNOT_ROWS;
/// xorshift32 — the shuffle source. Deterministic, so a red is reproducible.
struct Shuffle(u32);
impl Shuffle {
fn next(&mut self) -> u32 {
self.0 ^= self.0 << 13;
self.0 ^= self.0 >> 17;
self.0 ^= self.0 << 5;
self.0
}
/// Fisher-Yates permutation of `0..n`.
fn perm(&mut self, n: usize) -> Vec<usize> {
let mut v: Vec<usize> = (0..n).collect();
for i in (1..n).rev() {
let j = (self.next() % (i as u32 + 1)) as usize;
v.swap(i, j);
}
v
}
}
/// The inclusive spatial cell range a screen-space candidate covers.
///
/// Deliberately a SECOND spelling of `label_grid::footprint`, not a call to it. This
/// is the oracle that decides whether the fixture contends the grid at all, and an
/// oracle that calls the code under test cannot disagree with it — it would report
/// "the labels overlap" precisely when the shader thinks they do, which is the
/// question, not the answer.
fn cell_box(c: &LabelCandidate, p: &crate::label_grid::LabelGridParams) -> [u32; 4] {
let q = |v: f32, inv: f32, n: u32| (v * inv).floor().clamp(0.0, (n - 1) as f32) as u32;
[
q(c.pos[0] - c.half_px[0], p.inv_cell[0], GRID_W),
q(c.pos[1] - c.half_px[1], p.inv_cell[1], GRID_H),
q(c.pos[0] + c.half_px[0], p.inv_cell[0], GRID_W),
q(c.pos[1] + c.half_px[1], p.inv_cell[1], GRID_H),
]
}
/// Do two cell ranges share a cell — i.e. do these two labels contend?
fn cells_meet(a: [u32; 4], b: [u32; 4]) -> bool {
a[0] <= b[2] && b[0] <= a[2] && a[1] <= b[3] && b[1] <= a[3]
}
/// Names chosen so the **repeat filter cannot fire** on this fixture.
///
/// `repeat_slot` hashes `(name, coarse cell)` into [`NAME_SLOTS`], and a label is
/// suppressed if a HIGHER-priority label's slot lands anywhere in its 3x3
/// neighbourhood. Two *different* names can alias — `label_grid`'s own docs put that
/// at a low-single-digit percent per frame, and over 320 candidates the expected
/// number of aliased pairs is tens, not zero. Left to chance, several knots would
/// lose their rightful winner to a name collision and the closed-form expectation
/// below would simply be wrong, for a reason having nothing to do with `atomicMax`.
///
/// So the names are CONSTRUCTED to be alias-free rather than hoped to be: each
/// candidate takes the first suffix whose slot is outside every other candidate's
/// neighbourhood and whose neighbourhood excludes every other candidate's slot. The
/// repeat filter keeps its own dedicated arms (`same_name_repeats_are_thinned…` on
/// the CPU lane, the `Feldkircher Strasse` pair in the parity fixture); this arm is
/// about the spatial race, and mixing the two would only make a red ambiguous.
fn alias_free_names(cells: &[(i32, i32)]) -> Vec<String> {
let nbhd = |h: u32, cx: i32, cy: i32| -> [u32; 9] {
let mut o = [0u32; 9];
for (k, (dx, dy)) in (-1i32..=1).flat_map(|dy| (-1i32..=1).map(move |dx| (dx, dy))).enumerate() {
o[k] = repeat_slot(h, cx + dx, cy + dy);
}
o
};
let mut names: Vec<String> = Vec::new();
let mut chosen: Vec<u32> = Vec::new();
for (i, &(cx, cy)) in cells.iter().enumerate() {
let mut suffix = 0u32;
loop {
let name = format!("Strasse {i}/{suffix}");
let h = name_hash(&name);
let (my_slot, my_nbhd) = (repeat_slot(h, cx, cy), nbhd(h, cx, cy));
let clash = chosen.iter().enumerate().any(|(j, &hj)| {
let (jx, jy) = cells[j];
my_nbhd.contains(&repeat_slot(hj, jx, jy)) || nbhd(hj, jx, jy).contains(&my_slot)
});
if !clash {
names.push(name);
chosen.push(h);
break;
}
suffix += 1;
assert!(suffix < 100_000, "no alias-free name for candidate {i}");
}
}
names
}
/// The dense contended fixture. Returns the candidates, their glyph templates, the
/// **closed-form expectation** (one winner per knot: the highest priority in it) and
/// each knot's winning SLOT.
///
/// The expectation is closed-form on purpose — it is derived from the fixture's own
/// ranks, not from either lane. Every rival in a knot overlaps every other, so
/// whoever draws first occupies cells all seven others cover and settles them
/// permanently; the winner can only be the knot's `argmax priority`, and there can be
/// exactly one. Both assertions (which one, and only one) then have real content.
fn knot_fixture() -> (Vec<LabelCandidate>, Vec<GlyphSrc>, Vec<u32>, Vec<usize>) {
let mut sh = Shuffle(0x9E37_79B9);
let (mut pos, mut rank, mut glyphs_n) = (Vec::new(), Vec::new(), Vec::new());
let (mut winner_slot, mut oracle) = (Vec::new(), Vec::new());
for row in 0..KNOT_ROWS {
for col in 0..KNOT_COLS {
// 152 px column pitch against a widest padded box of 80 px, so knots are
// far apart in cells even though rivals inside one are 5 px apart.
let base = [90.0 + col as f32 * 152.0, 80.0 + row as f32 * 150.0];
let ranks = sh.perm(KNOT_RIVALS);
let sizes = sh.perm(KNOT_RIVALS);
let jx = sh.perm(KNOT_RIVALS);
let jy = sh.perm(KNOT_RIVALS);
let first = pos.len();
for s in 0..KNOT_RIVALS {
// +-2.8 / +-1.7 px: enough that no two rivals sit on identical
// pixels, small enough that all 8 padded boxes mutually overlap.
pos.push([base[0] + (jx[s] as f32 - 3.5) * 0.8, base[1] + (jy[s] as f32 - 3.5) * 0.5]);
// Ranks 2..=9 and glyph counts 3..=10, each a fresh permutation, so
// neither the winner's slot nor its size is guessable from the other.
rank.push(2 + ranks[s] as u8);
glyphs_n.push(3 + sizes[s] as u32);
}
let w = (0..KNOT_RIVALS).max_by_key(|&s| rank[first + s]).expect("a knot has rivals");
winner_slot.push(w);
oracle.push((first + w) as u32);
}
}
let inv = 1.0 / REPEAT_CELL_PX;
let cells: Vec<(i32, i32)> =
pos.iter().map(|p| ((p[0] * inv).floor() as i32, (p[1] * inv).floor() as i32)).collect();
let names = alias_free_names(&cells);
let specs: Vec<([f32; 2], u32, u8, u32, &str)> = (0..pos.len())
// `order` is the submission index; `rank` dominates it, which is the whole
// point — see the shuffle note above.
.map(|i| (pos[i], glyphs_n[i], rank[i], i as u32, names[i].as_str()))
.collect();
let (cands, glyphs) = build(&specs);
(cands, glyphs, oracle, winner_slot)
}
/// **THE DENSE PROOF — identity, not a count.** 320 labels in 40 contended knots;
/// the device must keep exactly the highest-priority rival of each, and the readback
/// must say WHICH.
///
/// Each assertion and what it alone can catch:
///
/// * the fixture really contends (every rival of a knot shares cells with every
/// other; no two knots share one) — an oracle spelled independently of the shader;
/// * priority order is not submission order — without this, `atomicMax` could be
/// deleted and the pass would still keep the right labels;
/// * `instance_count`, read out of the buffer `draw_indirect` reads, equals five
/// instances per glyph of exactly the expected 40 — a count assertion;
/// * the **id set** decoded from `out_glyphs` equals the expected 40 — the assertion
/// a count cannot make. Any wrong winner is a wrong id, even a wrong winner whose
/// glyph count happens to reproduce the right total;
/// * `labels_won == 40` — one winner per knot, so the pass neither kept a runner-up
/// nor culled a whole knot;
/// * the shared CPU rule agrees, which extends parity from 8 labels to 320.
#[test]
fn a_dense_contended_grid_keeps_exactly_the_top_priority_label_of_each_knot() {
let Some((device, queue, adapter, _gpu)) = headless_device() else {
eprintln!("[labels] no GPU adapter — skipping the dense arm (CPU-only CI)");
return;
};
let f = frame([1280.0, 768.0], 2);
let (cands, glyphs, oracle, winner_slot) = knot_fixture();
assert_eq!(cands.len(), KNOTS * KNOT_RIVALS, "320 candidates");
assert_eq!(oracle.len(), KNOTS);
// ── 0. MY OWN CHANNEL FIRST: the id encoding is injective over 320 labels ──
// `drawn_ids` decodes a f32 UV back to a label index. If two labels decoded to
// one id, or an id decoded to the wrong index, every identity assertion below
// would be noise that happened to agree.
for (i, c) in cands.iter().enumerate() {
let g = &glyphs[c.glyph_start as usize];
let back = drawn_ids(&[LabelGlyphInstance { uv_min: g.uv_min, ..Default::default() }]);
assert_eq!(
back.into_iter().collect::<Vec<_>>(),
vec![i as u32],
"the uv_min.y id channel must round-trip label {i} exactly"
);
}
// ── 1. THE FIXTURE CONTENDS — and only inside a knot ──────────────────
let p = f.grid_params();
let boxes: Vec<[u32; 4]> = cands.iter().map(|c| cell_box(c, &p)).collect();
for k in 0..KNOTS {
for a in 0..KNOT_RIVALS {
for b in a + 1..KNOT_RIVALS {
let (ia, ib) = (k * KNOT_RIVALS + a, k * KNOT_RIVALS + b);
assert!(
cells_meet(boxes[ia], boxes[ib]),
"knot {k}: rivals {a} and {b} must claim a common cell, or the knot \
has no race to resolve — {:?} vs {:?}",
boxes[ia],
boxes[ib]
);
}
}
}
for ka in 0..KNOTS {
for kb in ka + 1..KNOTS {
for a in 0..KNOT_RIVALS {
for b in 0..KNOT_RIVALS {
let (ia, ib) = (ka * KNOT_RIVALS + a, kb * KNOT_RIVALS + b);
assert!(
!cells_meet(boxes[ia], boxes[ib]),
"knots {ka} and {kb} must not touch, else the expected winner \
set is not one-per-knot"
);
}
}
}
}
// ── 2. PRIORITY ORDER IS NOT SUBMISSION ORDER (item 8's trap) ─────────
let mut by_priority: Vec<u32> = (0..cands.len() as u32).collect();
by_priority.sort_by_key(|&i| std::cmp::Reverse(cands[i as usize].priority));
assert_ne!(
by_priority,
(0..cands.len() as u32).collect::<Vec<_>>(),
"if priority order WERE submission order, atomicMax would be unobservable"
);
let distinct_slots: std::collections::BTreeSet<usize> = winner_slot.iter().copied().collect();
assert!(
distinct_slots.len() >= 4,
"the winning slot must vary across knots — got {distinct_slots:?}; all-0 would \
mean 'the first submitted wins', which is what a dead atomicMax also produces"
);
let first_submitted_wins = winner_slot.iter().filter(|&&s| s == 0).count();
assert!(
first_submitted_wins <= KNOTS / 4,
"only {first_submitted_wins} of {KNOTS} knots may be won by their first-submitted \
rival; more than that and the fixture is not really shuffled"
);
// ── 3. THE DEVICE — IDENTITY FIRST ────────────────────────────────────
//
// The id assertion comes BEFORE the count assertions, and the order is
// load-bearing rather than tidy. Both fire under the mutation that inverts the
// claim comparison so the LOWEST priority wins each cell (a pass that still
// draws 40 labels, all of them the wrong ones). With the count asserted first,
// that mutation stopped the test at "1385 expected, got 1195" and the id
// assertion never executed — so nothing had shown it capable of going red, which
// is the definition of a decorative guard. Asserted first, it reds with the exact
// missing/unexpected label lists, and the count then corroborates it.
let c = collider(&device, &queue, &cands, &glyphs);
let (args, stats) = collide(&c, &device, &queue, &f);
let want: std::collections::BTreeSet<u32> = oracle.iter().copied().collect();
let want_instances: u32 =
oracle.iter().map(|&i| cands[i as usize].glyph_count).sum::<u32>() * INSTANCES_PER_GLYPH;
let emitted = c.glyphs_readback(&device, &queue, args.instance_count);
let got = drawn_ids(&emitted);
assert_eq!(
got, want,
"the surviving labels must be exactly the top-ranked rival of each knot.\n \
missing: {:?}\n unexpected: {:?}",
want.difference(&got).collect::<Vec<_>>(),
got.difference(&want).collect::<Vec<_>>()
);
// Every drawn label contributed ALL its glyphs, so the count and the identity
// are consistent rather than two readings that merely both look plausible.
for &i in &oracle {
let n = emitted
.iter()
.filter(|g| ((g.uv_min[1] - LABEL_UV_BASE) / LABEL_UV_STEP).round() as u32 == i)
.count() as u32;
assert_eq!(
n,
cands[i as usize].glyph_count * INSTANCES_PER_GLYPH,
"winner {i} must emit all {} of its glyphs x {INSTANCES_PER_GLYPH}",
cands[i as usize].glyph_count
);
}
// ── 4. …and the count the DRAW consumes agrees ────────────────────────
assert_eq!(
args.instance_count, want_instances,
"the indirect instance_count must be 5 x the glyphs of the 40 top-ranked \
labels ({want_instances}); got {}",
args.instance_count
);
assert_eq!(
stats.labels_won as usize, KNOTS,
"exactly one winner per knot: {KNOTS} expected, the device counted {}",
stats.labels_won
);
assert_eq!(stats.labels_drawn as usize, KNOTS, "and all {KNOTS} fit the ink budget");
assert_eq!(stats.glyph_overflow, 0);
// ── 5. …and the shared CPU rule agrees, at 320 labels rather than 8 ───
let kept = resolve(&screen_labels(&cands, &f), &p);
assert_eq!(
kept.iter().copied().collect::<std::collections::BTreeSet<u32>>(),
want,
"the fail-safe CPU lane must resolve the same 40 labels the device did"
);
eprintln!(
"[labels] DENSE — {} candidates in {KNOTS} knots of {KNOT_RIVALS}; device kept {} \
(instance_count={}, winning slots {:?}) on {adapter}",
cands.len(),
got.len(),
args.instance_count,
winner_slot
);
crate::testmatrix::emit(
"facett-core",
"a_dense_contended_grid_keeps_exactly_the_top_priority_label_of_each_knot",
got == want && args.instance_count == want_instances && stats.labels_won as usize == KNOTS,
&format!(
"{} candidates, {} knots: kept {} of {}, instance_count={} ({adapter})",
cands.len(),
KNOTS,
got.len(),
cands.len(),
args.instance_count
),
);
}
/// A lattice of `n` labels whose padded boxes are separated by more than one grid
/// cell in both axes, so **nothing contends**. Same alias-free naming as the knots,
/// so a red here is spatial and never a name collision.
fn lattice_fixture(n: usize) -> (Vec<LabelCandidate>, Vec<GlyphSrc>) {
// Widest padded box is 10 glyphs x 7 px + 10 = 80 px, tallest 13 + 6 = 19 px.
// Pitch 88 x 25 leaves gaps of >= 8 px and 6 px, against cells of 5 x 4 px at
// 1280x768 — so adjacent boxes cannot even share a cell, which is the condition
// that matters here, not mere non-overlap of the rects.
let mut sh = Shuffle(0x85EB_CA6B);
let pos: Vec<[f32; 2]> =
(0..n).map(|i| [44.0 + (i % 14) as f32 * 88.0, 12.0 + (i / 14) as f32 * 25.0]).collect();
let counts: Vec<u32> = (0..n).map(|_| 3 + (sh.next() % 8)).collect();
let inv = 1.0 / REPEAT_CELL_PX;
let cells: Vec<(i32, i32)> =
pos.iter().map(|p| ((p[0] * inv).floor() as i32, (p[1] * inv).floor() as i32)).collect();
let names = alias_free_names(&cells);
let specs: Vec<([f32; 2], u32, u8, u32, &str)> =
(0..n).map(|i| (pos[i], counts[i], 5, i as u32, names[i].as_str())).collect();
build(&specs)
}
/// **NON-VACUITY AT DENSITY, both ways.** A declutter pass that only ever culls is
/// exactly as useless as one that never culls, and at 8 labels neither failure is
/// visible. So:
///
/// * **48 labels, nothing overlapping ⇒ all 48 draw, and the readback names all
/// 48.** Not "most", not "enough" — the whole set, by identity.
/// * **300 labels, nothing overlapping ⇒ the rule culls NONE of them** (`labels_won
/// == 300`, a device counter) **while the ink budget holds the draw to 48.** That
/// pair is the only arm in this file that can see the difference between "won the
/// rule" and "reached the paper"; `labels_drawn` used to be incremented before the
/// budget gate and so reported 300, a number that reads correct on every fixture
/// under the budget.
#[test]
fn a_dense_field_with_no_overlap_culls_nothing() {
let Some((device, queue, adapter, _gpu)) = headless_device() else {
eprintln!("[labels] no GPU adapter — skipping the no-overlap arm (CPU-only CI)");
return;
};
let f = frame([1280.0, 768.0], 2);
let p = f.grid_params();
// ── 1. exactly the ink budget: ALL of them draw, by identity ───────────
let n = MAX_VISIBLE_LABELS as usize;
let (cands, glyphs) = lattice_fixture(n);
// The oracle, spelled independently: no two cell ranges may meet.
let boxes: Vec<[u32; 4]> = cands.iter().map(|c| cell_box(c, &p)).collect();
for a in 0..n {
for b in a + 1..n {
assert!(!cells_meet(boxes[a], boxes[b]), "the lattice must not contend: {a} vs {b}");
}
}
let c = collider(&device, &queue, &cands, &glyphs);
let (args, stats) = collide(&c, &device, &queue, &f);
let all: u32 = cands.iter().map(|c| c.glyph_count).sum::<u32>() * INSTANCES_PER_GLYPH;
assert_eq!(args.instance_count, all, "nothing overlaps, so all {n} labels draw ({all} instances)");
assert_eq!(stats.labels_won as usize, n, "the rule culled nothing");
assert_eq!(stats.labels_drawn as usize, n);
let got = drawn_ids(&c.glyphs_readback(&device, &queue, args.instance_count));
assert_eq!(
got,
(0..n as u32).collect::<std::collections::BTreeSet<u32>>(),
"every one of the {n} labels must appear in the emitted instances, by id — \
missing {:?}",
(0..n as u32).filter(|i| !got.contains(i)).collect::<Vec<_>>()
);
// ── 2. over the budget: the rule keeps all 300, the paper takes 48 ────
let big = 300usize;
let (cands2, glyphs2) = lattice_fixture(big);
let boxes2: Vec<[u32; 4]> = cands2.iter().map(|c| cell_box(c, &p)).collect();
for a in 0..big {
for b in a + 1..big {
assert!(!cells_meet(boxes2[a], boxes2[b]), "the 300-lattice must not contend: {a} vs {b}");
}
}
let c2 = collider(&device, &queue, &cands2, &glyphs2);
let (args2, stats2) = collide(&c2, &device, &queue, &f);
assert_eq!(
stats2.labels_won as usize, big,
"the collision rule must keep every one of {big} non-overlapping labels — it \
counted {}; a pass that culls under load is the failure this arm exists for",
stats2.labels_won
);
assert_eq!(
stats2.labels_drawn, MAX_VISIBLE_LABELS,
"…and the ink budget must hold the DRAW to {MAX_VISIBLE_LABELS}, got {}",
stats2.labels_drawn
);
assert_eq!(stats2.glyph_overflow, 0, "the budget, not the capacity guard, is what bounded it");
// Which 48 is arrival order and so nondeterministic, but they must be 48
// DISTINCT labels each contributing all of its glyphs — the count and the
// identity have to be the same story.
let got2 = drawn_ids(&c2.glyphs_readback(&device, &queue, args2.instance_count));
assert_eq!(
got2.len(),
MAX_VISIBLE_LABELS as usize,
"{} distinct label ids in the emitted instances, not {}",
MAX_VISIBLE_LABELS,
got2.len()
);
let sum: u32 = got2.iter().map(|&i| cands2[i as usize].glyph_count).sum::<u32>() * INSTANCES_PER_GLYPH;
assert_eq!(
args2.instance_count, sum,
"instance_count must be exactly the glyphs of the labels the readback names"
);
eprintln!(
"[labels] NON-VACUITY — {n}/{n} drawn by identity ({all} instances); \
{big} offered -> labels_won={} labels_drawn={} instance_count={} on {adapter}",
stats2.labels_won, stats2.labels_drawn, args2.instance_count
);
crate::testmatrix::emit(
"facett-core",
"a_dense_field_with_no_overlap_culls_nothing",
args.instance_count == all
&& got.len() == n
&& stats2.labels_won as usize == big
&& stats2.labels_drawn == MAX_VISIBLE_LABELS,
&format!(
"{n} separated -> all {n} drawn; {big} separated -> won={} drawn={} ({adapter})",
stats2.labels_won, stats2.labels_drawn
),
);
}
/// **The lane really paints.** The indirect draw puts ink on a real target: pixels
/// appear where the surviving label is and stay empty where the suppressed one
/// was. `instance_count` alone cannot prove the draw consumed it — item 2's
/// finding was a pixel assertion passing at 486 px with **0 device frames**,
/// because a different painter was drawing. Here the only painter is this one.
#[test]
fn the_indirect_draw_inks_the_surviving_label_and_not_the_suppressed_one() {
let Some((device, queue, _adapter, _gpu)) = headless_device() else {
eprintln!("[labels] no GPU adapter — skipping ink proof (CPU-only CI)");
return;
};
let (w, h) = (256u32, 128u32);
let f = frame([w as f32, h as f32], 2);
// A winner on the left, a loser stacked on it, and a second winner far right.
let (cands, glyphs) = build(&[
([60.0, 64.0], 3, 9, 0, "win"),
([62.0, 65.0], 3, 2, 1, "lose"),
([200.0, 64.0], 3, 9, 2, "far"),
]);
let c = collider(&device, &queue, &cands, &glyphs);
let target = device.create_texture(&wgpu::TextureDescriptor {
label: Some("label_ink_target"),
size: wgpu::Extent3d { width: w, height: h, depth_or_array_layers: 1 },
mip_level_count: 1,
sample_count: crate::render::gpu::NO_MSAA_SAMPLES,
dimension: wgpu::TextureDimension::D2,
format: wgpu::TextureFormat::Rgba8Unorm,
usage: wgpu::TextureUsages::RENDER_ATTACHMENT | wgpu::TextureUsages::COPY_SRC,
view_formats: &[],
});
let view = target.create_view(&wgpu::TextureViewDescriptor::default());
let mut enc = device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
c.run(&queue, &mut enc, &f);
c.render_to(&mut enc, &view, false);
queue.submit(Some(enc.finish()));
device.poll(wgpu::PollType::wait_indefinitely()).ok();
// The two winners have 3 glyphs each; the loser draws nothing.
let args = c.state_readback(&device, &queue).args;
assert_eq!(args.instance_count, 2 * 3 * INSTANCES_PER_GLYPH, "two winners of 3 glyphs each");
let px = read_texture(&device, &queue, &target, w, h);
let lit = |x0: u32, x1: u32| -> usize {
(0..h)
.flat_map(|y| (x0..x1).map(move |x| (x, y)))
.filter(|&(x, y)| px[((y * w + x) * 4 + 3) as usize] > 8)
.count()
};
let left = lit(30, 95);
let right = lit(170, 235);
let gap = lit(110, 160);
assert!(left > 20, "the left winner inked pixels (got {left})");
assert!(right > 20, "the right winner inked pixels (got {right})");
assert_eq!(gap, 0, "the empty band between them stayed empty (got {gap})");
// …and the ink is the ink colour, not the halo and not the atlas white.
let strongest = px
.as_chunks::<4>().0.iter()
.filter(|p| p[3] > 200)
.map(|p| [p[0], p[1], p[2]])
.max_by_key(|c| u32::from(c[2]));
assert!(strongest.is_some(), "at least one fully-opaque ink pixel landed");
eprintln!("[labels] ink proof — left={left} px right={right} px gap={gap} px, instance_count={}", args.instance_count);
}
/// Texture → tight RGBA, un-padding the 256-byte row alignment.
fn read_texture(
device: &wgpu::Device,
queue: &wgpu::Queue,
tex: &wgpu::Texture,
w: u32,
h: u32,
) -> Vec<u8> {
let unpadded = w * 4;
let padded = unpadded.div_ceil(wgpu::COPY_BYTES_PER_ROW_ALIGNMENT) * wgpu::COPY_BYTES_PER_ROW_ALIGNMENT;
let buf = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("label_ink_readback"),
size: u64::from(padded * h),
usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ,
mapped_at_creation: false,
});
let mut enc = device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
enc.copy_texture_to_buffer(
wgpu::TexelCopyTextureInfo {
texture: tex,
mip_level: 0,
origin: wgpu::Origin3d::ZERO,
aspect: wgpu::TextureAspect::All,
},
wgpu::TexelCopyBufferInfo {
buffer: &buf,
layout: wgpu::TexelCopyBufferLayout {
offset: 0,
bytes_per_row: Some(padded),
rows_per_image: Some(h),
},
},
wgpu::Extent3d { width: w, height: h, depth_or_array_layers: 1 },
);
queue.submit(Some(enc.finish()));
let slice = buf.slice(..);
let (tx, rx) = std::sync::mpsc::channel();
slice.map_async(wgpu::MapMode::Read, move |r| {
let _ = tx.send(r);
});
device.poll(wgpu::PollType::wait_indefinitely()).ok();
let _ = rx.recv();
let data = slice.get_mapped_range().expect("wgpu: get_mapped_range on a mapped buffer");
let mut out = Vec::with_capacity((w * h * 4) as usize);
for row in 0..h {
let s = (row * padded) as usize;
out.extend_from_slice(&data[s..s + unpadded as usize]);
}
drop(data);
buf.unmap();
out
}
}