facett-core 0.1.19

//! **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: &[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().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
            .chunks_exact(4)
            .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();
        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
    }
}