concinnity-device 0.18.64

// src/metal/raytrace.rs
//
// Hardware ray-tracing acceleration structures for the Metal backend. Builds,
// from the shared static vertex / index buffers and the `DrawObject` list, the
// bottom- and top-level acceleration structures (BLAS / TLAS) the RT-reflection
// kernel traces against, plus a per-instance geometry table the kernel uses to
// fetch the hit triangle and shade it.
//
// One primitive BLAS per object over its slice of the shared buffers, one
// instance in the TLAS per object (transform = the object's model matrix,
// instance_id = the object index). The BLAS describe object-space geometry and
// never change for a rigid transform; only the TLAS instance transforms (and
// the geometry table's per-instance model matrices the kernel shades with) move
// when a prop moves.
//
// Dynamic transforms (`RtDynamicMode`) update the structures per frame. The
// per-frame skinned update (`rebuild_skinned`) keeps the persistent static +
// cluster BLAS, re-skins the current pose, and rebuilds only the skinned BLAS +
// TLAS + geometry table. It is fully asynchronous: NO `waitUntilCompleted`. The
// three GPU steps are committed on the one shared queue in dependency order: skin
// compute (writes the deformed buffer), then the BLAS/TLAS build (reads it), all
// in `rt_dynamic_update`, strictly before the reflection-trace command buffer
// (committed later in `execute_graph`). Same-queue FIFO commit order runs them
// skin → build → trace; the render graph already depends on exactly this
// event-free ordering for every cross-pass read. Faults (which can no longer be
// caught synchronously) are surfaced from completion handlers.
//
// (Historical note: an earlier bisect concluded no GPU-side primitive orders these
// steps and kept the rebuild synchronous. That predated the fix for the actual
// fault (a CPU/GPU `RtGeomEntry` struct-layout mismatch that made the trace read
// out of bounds), so those fault observations were the layout bug, not an ordering
// failure. With it fixed, same-queue commit order alone orders the rebuild.)
//
// The one-time seed build and the incremental topology refresh allocate fresh
// and park the outgoing structures / Shared buffers in a frame-tagged
// deferred-free pool (`RetirePool`) until the frames-in-flight fence retires the
// frames whose still-in-flight trace could read them. The per-frame skinned
// update instead rebuilds in place in a ring slot (`rt_ring`), which is sound
// precisely because it runs on every frame: see that module's header for why the
// static paths cannot use the same trick.
#![deny(unsafe_op_in_unsafe_fn)]

use std::ptr::NonNull;

use objc2::rc::Retained;
use objc2::runtime::ProtocolObject;
use objc2_foundation::NSArray;
use objc2_metal::{
    MTLAccelerationStructure, MTLAccelerationStructureCommandEncoder,
    MTLAccelerationStructureGeometryDescriptor, MTLAccelerationStructureInstanceDescriptor,
    MTLAccelerationStructureInstanceDescriptorType, MTLAccelerationStructureInstanceOptions,
    MTLAccelerationStructureTriangleGeometryDescriptor, MTLAccelerationStructureUsage,
    MTLAttributeFormat, MTLBuffer, MTLCommandBuffer as _, MTLCommandBufferStatus,
    MTLCommandEncoder as _, MTLCommandQueue as _, MTLComputeCommandEncoder as _,
    MTLComputePipelineState, MTLDevice as _, MTLIndexType,
    MTLInstanceAccelerationStructureDescriptor, MTLPackedFloat3, MTLPackedFloat4x3,
    MTLPrimitiveAccelerationStructureDescriptor, MTLRenderCommandEncoder, MTLRenderPipelineState,
    MTLRenderStages, MTLResource, MTLResourceOptions, MTLResourceUsage, MTLSize,
};

use super::context::write_buffer_slice;
use super::encode::ComputeEncode;
use super::rt_ring::{BlasUpdate, RtFrameRing, SkinnedBlasSet, SkinnedShape, TlasKey};
use super::transient::RetirePool;
use crate::gfx::render_types::{DrawObject, InstancedCluster, RtGeomEntry, SkinnedDrawObject};
use crate::gfx::rt_geom::{cluster_geom_entry, geom_entry, skinned_geom_entry};
use crate::gfx::rt_reflections::RtReflectionSettings;
// The dynamic-update mode ladder lives in concinnity-render; re-exported so the
// `super::raytrace::RtDynamicMode` path (init + draw) keeps resolving.
pub(crate) use crate::gfx::rt_geom::RtDynamicMode;
// Shared with the Vulkan and DirectX hosts: one `.slang` declares it now.
use concinnity_render::uniforms::SkinParams;

// Byte stride of a `Vertex` in the shared vertex buffer (pos + normal + tangent
// + colour + uv = 14 floats). The RT kernel reads positions at this stride; the
// main-pass skinned fold sizes its deformed buffer by it too.
pub(in crate::metal) const VERTEX_STRIDE: usize = 56;

// All hardware-ray-traced-reflection state grouped into one feature unit: the
// resolved tunables, the scene acceleration structure, the dynamic-update
// mode + failure-streak flag, and the resolve / textured-resolve / skinning
// pipelines. `settings`/`accel`/the pipelines are `Some` only when RT
// reflections are on and the GPU supports ray tracing (see the per-field
// docs on [`MtlContext`](super::context::MtlContext) for the exact gates).
pub(crate) struct RtState {
    // Resolved + clamped tunables. `Some` only when the world's
    // `PostProcessConfig` sets `ray_traced_reflections` AND the GPU supports
    // ray tracing; gates the RT pass. RT takes precedence over SSR resolve
    // and reuses `ssr.targets.output` as its resolve target.
    pub settings: Option<RtReflectionSettings>,
    // Scene acceleration structure (BLAS/TLAS) + geometry table. `Some` only
    // when RT reflections are on and the scene has resident geometry; updated
    // per frame when `dynamic_mode` is dynamic. Resolution-independent.
    pub accel: Option<RtAccelData>,
    // How the acceleration structure is kept current as props move (the
    // launch's `--rt-dynamic` request; `Auto` by default).
    pub dynamic_mode: RtDynamicMode,
    // Whether skinned meshes join the BVH (the launch's
    // `--rt-skinned-geometry` request; in by default). Clear it and the BVH
    // covers static + instanced geometry only, isolating the skinned trace path.
    pub skinned_geometry: bool,
    // Whether the per-frame BVH update is currently in a failure streak. A
    // transient rebuild failure is non-fatal (keep last frame's BVH) and
    // logged once per streak rather than every frame.
    pub update_failed: bool,
    // Set when an operation changes the RT-relevant draw set (a streamed chunk
    // added/removed, a prop cloned, a material edit that flips RT participation)
    // since the last update. The per-frame update consumes it to refresh the
    // BLAS topology -- reusing every unchanged BLAS and building only the new
    // ones -- instead of either ignoring the change (the default `Auto` path
    // only watches transforms of the prior set) or rebuilding every BLAS.
    pub topology_dirty: bool,
    // Resolve pipeline, flat-tint hit shading. Used for non-bindless worlds
    // (no albedo pool).
    pub pipeline: Option<Retained<ProtocolObject<dyn MTLRenderPipelineState>>>,
    // Resolve pipeline, textured hit shading (samples the bindless albedo pool
    // at buffer(7)). Preferred over `pipeline` when the bindless texture
    // argument buffer is available this frame.
    pub pipeline_textured: Option<Retained<ProtocolObject<dyn MTLRenderPipelineState>>>,
    // Compute-skinning pipeline that deforms skinned vertices into a buffer the
    // BVH can trace. Consumed each frame by `rebuild_rt_accel` to pose skinned
    // geometry before the skinned BLAS build.
    pub skin_pipeline: Option<Retained<ProtocolObject<dyn MTLComputePipelineState>>>,
}

// Identifies the geometry a draw-object BLAS traces, on the shared
// vertex/index buffers. Two draw objects with the same signature trace identical
// geometry, so a topology refresh can reuse the existing BLAS instead of
// building a new one. Sound because the shared buffer *objects* are stable once
// streaming is set up (`add_chunk_mesh` / `remove_chunk_mesh` write regions in
// place; a buffer swap goes through a full rebuild, not this path). A streamed
// mesh returns on whatever slice the sub-allocator hands out, so the signature
// moves with it and the BLAS is rebuilt rather than wrongly reused.
// `base_vertex` + `index_offset` + `index_count` are exactly the inputs
// `prim_desc_for` uses; `vertex_offset` is carried too so a static draw (whose
// `base_vertex` is 0) still distinguishes distinct vertex regions.
//
// The slice location alone is not enough: an asset hot-reload rewrites a slot's
// bytes in place at unchanged offsets, which leaves every field above equal.
// `generation` (the draw object's `geometry_generation`) moves on each such
// rewrite so the stale BLAS is rebuilt instead of reused. Mirrors
// `concinnity_render::rt_topology::GeomSig`.
#[derive(Clone, Copy, PartialEq, Eq)]
pub(crate) struct GeomSig {
    base_vertex: i32,
    vertex_offset: usize,
    index_offset: usize,
    index_count: usize,
    generation: u32,
}

impl GeomSig {
    fn of(obj: &DrawObject) -> Self {
        Self {
            base_vertex: obj.base_vertex,
            vertex_offset: obj.vertex_offset,
            index_offset: obj.index_offset,
            index_count: obj.index_count,
            generation: obj.geometry_generation,
        }
    }
}

// Per-new-slot decision for a topology refresh of the draw-object BLAS head.
struct TopologyPlan {
    // `reuse[j] == Some(k)`: new draw slot `j` reuses the old draw BLAS at index
    // `k` (its geometry is unchanged). `None`: build a fresh BLAS for slot `j`.
    reuse: Vec<Option<usize>>,
    // Old draw BLAS indices no longer referenced by any new slot -- retire them.
    retire: Vec<usize>,
}

// Decide, for the draw-object BLAS head only, which BLAS to reuse, which to
// build, and which to retire when the participating draw set changes. Matches
// old and new slots by `draw.objects` index AND geometry signature: a slot whose
// geometry moved (a chunk slot recycled for a different chunk) does not match, so
// it rebuilds. Pure so it is unit-testable without Metal.
fn plan_topology_refresh(
    old_indices: &[usize],
    old_sigs: &[GeomSig],
    new_indices: &[usize],
    new_sigs: &[GeomSig],
) -> TopologyPlan {
    use std::collections::HashMap;
    // draw.objects index -> (position in the old draw BLAS head, its signature).
    // `object_indices` entries are unique (one per draw slot), so this is 1:1.
    let mut by_idx: HashMap<usize, (usize, GeomSig)> = HashMap::with_capacity(old_indices.len());
    for (k, (&idx, &sig)) in old_indices.iter().zip(old_sigs).enumerate() {
        by_idx.insert(idx, (k, sig));
    }
    let mut used = vec![false; old_indices.len()];
    let mut reuse = Vec::with_capacity(new_indices.len());
    for (&idx, &sig) in new_indices.iter().zip(new_sigs) {
        match by_idx.get(&idx) {
            Some(&(k, old_sig)) if old_sig == sig && !used[k] => {
                used[k] = true;
                reuse.push(Some(k));
            }
            _ => reuse.push(None),
        }
    }
    let retire = used
        .iter()
        .enumerate()
        .filter(|&(_, &u)| !u)
        .map(|(k, _)| k)
        .collect();
    TopologyPlan { reuse, retire }
}

// The acceleration structures + geometry table for hardware ray tracing. Held
// on the context behind an `Option`; present only when the world enables RT
// reflections, the GPU supports ray tracing, and the scene has geometry.
pub(crate) struct RtAccelData {
    // Bottom-level acceleration structures, in build order: one per
    // participating `DrawObject` (in `object_indices` order), then one per
    // instanced cluster, then one per skinned object. This Vec is the sole CPU
    // owner that keeps every BLAS alive: a TLAS does not retain the structures it
    // references, and the `useResource` the kernel encoder issues only declares
    // residency, not lifetime, so a BLAS must stay owned here for as long as any
    // in-flight trace can reach it through the TLAS. A skinned rebuild produces a
    // whole fresh `RtAccelData`; the outgoing one is parked in the context's
    // retire pool until the frames-in-flight fence retires the frames that could
    // still trace it.
    pub blas: Vec<Retained<ProtocolObject<dyn MTLAccelerationStructure>>>,
    // How many leading entries of `blas` are the persistent static + cluster
    // BLAS, built once and never rebuilt (a rigid transform leaves object-space
    // geometry unchanged). Skinned BLAS occupy `blas[static_blas_count..]` and
    // are rebuilt each frame from the current pose; a skinned object's
    // `accelerationStructureIndex` is `static_blas_count + si`. Lets the
    // per-frame skinned update rebuild only the skinned tail and keep the head.
    static_blas_count: usize,
    // The top-level (instance) acceleration structure the kernel traces.
    pub tlas: Retained<ProtocolObject<dyn MTLAccelerationStructure>>,
    // `[RtGeomEntry; instance_count]`, indexed by the intersector's
    // `instance_id`. Lets the kernel find the hit triangle + shade it. Carries
    // each instance's model matrix, which the kernel uses to bring the hit
    // normal to world space, so it moves in lockstep with the TLAS transforms.
    pub geom_table: Retained<ProtocolObject<dyn MTLBuffer>>,

    // Per-frame update state.
    // Indices into the frame's `draw.objects` for the objects that participate,
    // in BLAS / instance order. Lets an update re-read current transforms in
    // the exact order the BLAS were built, and detect a changed draw list.
    object_indices: Vec<usize>,
    // The geometry signature each draw-object BLAS (`blas[..object_indices.len()]`)
    // was built from, parallel to `object_indices`. A topology refresh compares
    // these against the current draw set to reuse every unchanged BLAS and build
    // only the new / changed ones.
    draw_blas_sigs: Vec<GeomSig>,
    // Each participating object's model matrix as baked into the current TLAS,
    // in `object_indices` order. The `Auto` dirty check compares the live draw
    // list against these to decide whether a rebuild is needed.
    cached_models: Vec<[[f32; 4]; 4]>,
    // The TLAS instance descriptors for every instanced-cluster instance, in
    // the order they follow the draw-object instances. Clusters are baked
    // static into the BVH, so a per-frame TLAS rebuild re-appends these
    // verbatim after the freshly-transformed draw-object instances (their
    // `accelerationStructureIndex` points at the cluster BLAS, which never
    // move in `blas`). Empty when the world declares no `InstancedProp`.
    cluster_instances: Vec<MTLAccelerationStructureInstanceDescriptor>,
    // The geometry-table entries for the cluster instances, parallel to
    // `cluster_instances`. Re-appended alongside them on a rebuild.
    cluster_geom: Vec<RtGeomEntry>,
    // Private scratch buffer sized for the largest of every BLAS build and the
    // TLAS build, reused by the per-frame TLAS rebuild (the instance count is
    // fixed across rebuilds, so the init sizing always suffices).
    scratch: Retained<ProtocolObject<dyn MTLBuffer>>,
    // The TLAS instance-descriptor buffer of the last *asynchronous* static
    // build. Only the TLAS build reads it (the built TLAS bakes the instances),
    // so it is not bound to the trace; it is held here because a topology
    // refresh commits without waiting and the build keeps reading it after the
    // call returns. The per-frame skinned update does not use it -- its instance
    // buffer is a ring slot the ring keeps alive.
    instance_buffer: Retained<ProtocolObject<dyn MTLBuffer>>,

    // Deformed (posed) skinned vertices in the static 56-byte `Vertex` layout,
    // written by the `rt_skin` compute pass and traced by the skinned BLAS. The
    // reflection kernel reads it (buffer 5) for skinned hits. A 1-element dummy
    // when the scene has no skinned geometry, so the encoder always has a buffer
    // to bind. The skinned rebuild allocates a fresh one each frame and retires
    // the old through `retire_pool` (it cannot overwrite in place: a prior
    // frame's trace may still be reading it).
    pub deformed_verts: Retained<ProtocolObject<dyn MTLBuffer>>,
    // The shared skinned index buffer, cloned here so the reflection kernel
    // can bind it (buffer 6) for skinned hits. A 1-element dummy when there is
    // no skinned geometry.
    pub skinned_indices: Retained<ProtocolObject<dyn MTLBuffer>>,
    // Outgoing structures / Shared buffers from prior rebuilds, held alive until
    // the frames-in-flight fence retires the frames whose still-in-flight trace
    // could read them. Drained once per frame in `rt_dynamic_update`. The seed
    // build and an incremental topology refresh allocate fresh and park the old
    // here rather than freeing in place; the per-frame skinned update only pushes
    // when it takes over from one of them (see `ring_published`).
    retire_pool: RetirePool<RetiredRt>,

    // Per-in-flight-frame storage the skinned update rebuilds in place instead of
    // allocating fresh. One slot per frame in flight; see `super::rt_ring`.
    ring: RtFrameRing,
    // A 1-vertex Shared buffer bound at the deformed-vertex slot whenever no
    // skinned geometry is being traced, so the encoder always has a buffer for
    // the binding the shader declares (the skinned branch is never taken then).
    // Also what `release_skinned` falls back to when the ring stops publishing.
    deformed_dummy: Retained<ProtocolObject<dyn MTLBuffer>>,
    // A single-identity joint palette the skin dispatch binds for an object with
    // no pose, so it deforms to bind pose rather than reading whatever the
    // frame's pre-built palette buffer happens to hold.
    identity_palette: Retained<ProtocolObject<dyn MTLBuffer>>,
    // Bumped whenever the persistent BLAS head (`blas[..static_blas_count]`)
    // changes identity, so every ring slot's cached TLAS descriptor rebuilds.
    head_generation: u64,
    // Whether `tlas` / `geom_table` / `deformed_verts` / the skinned tail of
    // `blas` are currently ring-owned clones. When they are not -- right after
    // the seed build or a topology refresh built fresh ones -- the next skinned
    // update must park the outgoing handles in `retire_pool` instead of dropping
    // them, because a prior in-flight frame's trace can still reach them.
    ring_published: bool,
    // Persistent CPU scratch for the per-frame skinned update, swapped out with
    // `mem::take` so its heap capacity survives the frame.
    update_scratch: RtUpdateScratch,
}

// The scene-scaled `Vec`s the per-frame skinned update fills. Kept on the accel
// so each frame reuses the capacity instead of collecting fresh ones.
#[derive(Default)]
struct RtUpdateScratch {
    // Indices into the frame's skinned draw objects, for those visible with real
    // triangles, in skinned-BLAS order.
    skinned: Vec<usize>,
    // The geometry each of those objects' BLAS covers, parallel to `skinned`.
    shapes: Vec<SkinnedShape>,
    // This frame's TLAS instance descriptors and per-instance geometry entries,
    // in instance order.
    instances: Vec<MTLAccelerationStructureInstanceDescriptor>,
    geom: Vec<RtGeomEntry>,
}

// Outgoing RT resources parked by a skinned rebuild or an incremental topology
// refresh for deferred free. Never read again: they exist only to keep the Metal
// handles (and thus the GPU allocations) valid until `RetirePool` drops them,
// once the fence guarantees no in-flight trace can still reference them.
struct RetiredRt {
    #[expect(
        dead_code,
        reason = "held so the acceleration structures stay valid until RetirePool drops them"
    )]
    structures: Vec<Retained<ProtocolObject<dyn MTLAccelerationStructure>>>,
    #[expect(
        dead_code,
        reason = "held so the backing buffers stay valid until RetirePool drops them"
    )]
    buffers: Vec<Retained<ProtocolObject<dyn MTLBuffer>>>,
}

// The per-frame skinned-geometry inputs `build_rt_accel` needs to deform and
// add skinned objects to the BVH. Assembled from the context's skinned state;
// `None` skips skinned geometry entirely (the static-only path).
pub(crate) struct SkinnedRtInputs<'a> {
    // One entry per skinned mesh (only `visible`, real-triangle objects build).
    pub objects: &'a [SkinnedDrawObject],
    // Shared skinned vertex buffer (`SkinnedVertex`, 80-byte stride) the skin
    // kernel reads bind-pose vertices from.
    pub vertex_buffer: &'a Retained<ProtocolObject<dyn MTLBuffer>>,
    // Shared skinned index buffer (absolute indices) the skinned BLAS and
    // the reflection kernel address the deformed buffer with. Cloned into
    // `RtAccelData` so the reflection encoder can bind it.
    pub index_buffer: &'a Retained<ProtocolObject<dyn MTLBuffer>>,
    // Per-object joint palettes, parallel to `objects`; uploaded transiently
    // and consumed by the skin kernel.
    pub joint_matrices: &'a [Vec<[[f32; 4]; 4]>],
    // The compiled `rt_skin` compute pipeline.
    pub skin_pipeline: &'a ProtocolObject<dyn MTLComputePipelineState>,
}

// The device + queue every acceleration-structure build encodes on, plus the
// frames-in-flight depth the per-frame ring is sized to.
#[derive(Clone, Copy)]
pub(crate) struct RtGpu<'a> {
    pub device: &'a ProtocolObject<dyn objc2_metal::MTLDevice>,
    pub command_queue: &'a ProtocolObject<dyn objc2_metal::MTLCommandQueue>,
    pub frames_in_flight: usize,
}

// Which frame a per-frame RT update belongs to: the id outgoing resources are
// tagged with in the retire pool, and the ring slot whose storage this frame's
// skinned structures are rebuilt in. Both come from the same frame counter the
// rest of the backend's per-frame rings use.
#[derive(Clone, Copy)]
pub(crate) struct RtFrame {
    pub id: u64,
    pub ring_slot: usize,
}

// The shared static geometry buffers a BLAS build addresses: the u32-indexed
// vertex + index buffers that hold every non-skinned draw object and cluster.
#[derive(Clone, Copy)]
pub(crate) struct RtStaticGeometry<'a> {
    pub vertex_buffer: &'a ProtocolObject<dyn MTLBuffer>,
    pub index_buffer: &'a ProtocolObject<dyn MTLBuffer>,
}

// The scene geometry a full BVH build spans: the draw objects and the instanced
// clusters (both filtered to resident, real-triangle participants inside).
#[derive(Clone, Copy)]
pub(crate) struct RtSceneGeometry<'a> {
    pub draw_objects: &'a [DrawObject],
    pub clusters: &'a [InstancedCluster],
}

// The shared texture pool's real-texture count a geometry table resolves its
// per-object albedo / normal indices against (the flat-normal fallback sits at
// this index), so an index never points past the resident pool.
#[derive(Clone, Copy)]
pub(crate) struct RtTextureCounts {
    pub albedo_count: usize,
}

// The trailing knobs of an incremental topology refresh: whether see-through
// glass is excluded from the BLAS, whether to also rebuild the TLAS inline (the
// no-skinned path), and the frame id the retired resources are parked under.
#[derive(Clone, Copy)]
pub(crate) struct RtTopologyRefreshOptions {
    pub exclude_seethrough: bool,
    pub build_tlas: bool,
    pub frame_id: u64,
}

// Whether the GPU supports hardware ray tracing. Apple-silicon GPUs report
// `true`; Intel / most AMD Macs report `false`, in which case the caller falls
// back to SSR (or no reflections). Mirrors the capability gates the MetalFX /
// HDR paths use at init.
pub(crate) fn raytracing_supported(device: &ProtocolObject<dyn objc2_metal::MTLDevice>) -> bool {
    device.supportsRaytracing()
}

// Pack a column-major object-to-world `model` matrix into Metal's
// `MTLPackedFloat4x3` instance transform. The packed form is the first three
// rows of each of the four columns (the affine `[0,0,0,1]` bottom row is
// dropped), so `columns[c] = (model[c][0], model[c][1], model[c][2])`. Getting
// this transpose wrong silently mirrors / shears every reflection, so it is
// unit-tested.
pub(crate) fn pack_instance_transform(model: [[f32; 4]; 4]) -> MTLPackedFloat4x3 {
    let col = |c: usize| MTLPackedFloat3 {
        x: model[c][0],
        y: model[c][1],
        z: model[c][2],
    };
    MTLPackedFloat4x3 {
        columns: [col(0), col(1), col(2), col(3)],
    }
}

// A primitive (triangle) BLAS descriptor over a slice of the shared buffers.
// `vertexBufferOffset = base_vertex * stride` so a chunk with mesh-relative
// indices and a non-zero base vertex still resolves; static geometry and
// instanced clusters use base_vertex 0 (their indices are already absolute).
// `index_type` selects the index width; every shared buffer is `UInt32`.
// `usage` is `Refit` only for the skinned structures the
// per-frame update re-fits in place; Metal requires it at build time for a later
// refit to be legal, and it is left `None` everywhere else so the static
// structures keep the better-optimised default tree.
fn prim_desc_for(
    vertex_buffer: &ProtocolObject<dyn MTLBuffer>,
    index_buffer: &ProtocolObject<dyn MTLBuffer>,
    base_vertex: usize,
    index_offset: usize,
    index_count: usize,
    index_type: MTLIndexType,
    usage: MTLAccelerationStructureUsage,
) -> Retained<MTLPrimitiveAccelerationStructureDescriptor> {
    let index_bytes = match index_type {
        MTLIndexType::UInt16 => 2,
        _ => 4,
    };
    // SAFETY: plain descriptor property setters, all values in range.
    let geo = unsafe {
        let g = MTLAccelerationStructureTriangleGeometryDescriptor::descriptor();
        g.setVertexBuffer(Some(vertex_buffer));
        g.setVertexBufferOffset(base_vertex * VERTEX_STRIDE);
        g.setVertexStride(VERTEX_STRIDE);
        g.setVertexFormat(MTLAttributeFormat::Float3);
        g.setIndexBuffer(Some(index_buffer));
        g.setIndexBufferOffset(index_offset * index_bytes);
        g.setIndexType(index_type);
        g.setTriangleCount(index_count / 3);
        g
    };
    let geo_ref: &MTLAccelerationStructureGeometryDescriptor = &geo;
    let geos = NSArray::from_slice(&[geo_ref]);
    let prim = MTLPrimitiveAccelerationStructureDescriptor::descriptor();
    prim.setGeometryDescriptors(Some(&geos));
    prim.setUsage(usage);
    prim
}

// An instance descriptor with an explicit transform + BLAS index
// (`accelerationStructureIndex` selects which BLAS this instance uses). The
// shader indexes the geometry table by the intersector's `instance_id`, which
// for `MTLAccelerationStructureInstanceDescriptorType::Default` is the
// instance's position in the instance buffer (NOT the
// `accelerationStructureIndex`), so the table carries one entry per instance,
// in instance order (multiple cluster instances share one BLAS but get distinct
// entries). See `build_rt_accel`.
fn instance_desc_at(
    model: [[f32; 4]; 4],
    blas_index: u32,
) -> MTLAccelerationStructureInstanceDescriptor {
    MTLAccelerationStructureInstanceDescriptor {
        transformationMatrix: pack_instance_transform(model),
        options: MTLAccelerationStructureInstanceOptions::Opaque,
        mask: 0xFF,
        intersectionFunctionTableOffset: 0,
        accelerationStructureIndex: blas_index,
    }
}

// The instance descriptor for draw object `i` (its BLAS index == its position).
fn instance_desc(obj: &DrawObject, i: usize) -> MTLAccelerationStructureInstanceDescriptor {
    instance_desc_at(obj.model, i as u32)
}

// The TLAS descriptor over `blas_refs`, reading transforms from
// `instance_buffer`. Takes plain references so the BLAS array can be assembled
// from more than one source (e.g. persistent static BLAS followed by this
// frame's fresh skinned BLAS).
fn make_tlas_desc_from_refs(
    blas_refs: &[&ProtocolObject<dyn MTLAccelerationStructure>],
    instance_buffer: &ProtocolObject<dyn MTLBuffer>,
    instance_count: usize,
) -> Retained<MTLInstanceAccelerationStructureDescriptor> {
    let blas_array = NSArray::from_slice(blas_refs);
    let desc = MTLInstanceAccelerationStructureDescriptor::descriptor();
    desc.setInstancedAccelerationStructures(Some(&blas_array));
    desc.setInstanceCount(instance_count);
    desc.setInstanceDescriptorBuffer(Some(instance_buffer));
    desc.setInstanceDescriptorType(MTLAccelerationStructureInstanceDescriptorType::Default);
    desc
}

// The TLAS descriptor over `blas`, reading transforms from `instance_buffer`.
fn make_tlas_desc(
    blas: &[Retained<ProtocolObject<dyn MTLAccelerationStructure>>],
    instance_buffer: &ProtocolObject<dyn MTLBuffer>,
    instance_count: usize,
) -> Retained<MTLInstanceAccelerationStructureDescriptor> {
    let blas_refs: Vec<&ProtocolObject<dyn MTLAccelerationStructure>> =
        blas.iter().map(|b| b.as_ref()).collect();
    make_tlas_desc_from_refs(&blas_refs, instance_buffer, instance_count)
}

// Declare BLAS that a TLAS build references resident on the build encoder. A
// TLAS build reads the primitive structures its instances point to; unlike a
// direct buffer binding, that indirect reference does not make them resident,
// so Metal requires an explicit `useResource` (or `useHeap`) or the build can
// read a non-resident structure and fault. Only structures built on an *earlier*
// command buffer need this: BLAS built in the same encoder are already resident
// and ordered by same-encoder hazard tracking, so they must NOT be passed here.
fn declare_blas_resident<'a>(
    enc: &ProtocolObject<dyn MTLAccelerationStructureCommandEncoder>,
    blas: impl IntoIterator<Item = &'a Retained<ProtocolObject<dyn MTLAccelerationStructure>>>,
) {
    for b in blas {
        enc.useResource_usage(ProtocolObject::from_ref(&**b), MTLResourceUsage::Read);
    }
}

// Declare every BLAS resident for a fragment-stage trace in ONE batched
// `useResources` call rather than N per-BLAS ones. A trace render pass (the
// transparent glass/water pass, the RT-reflection resolve) reaches each BLAS
// indirectly through the TLAS, which does NOT make them resident, so the pass
// must declare them itself. Batching collapses the per-frame Obj-C message-send
// count on BLAS-heavy worlds (the driver records the same residency set either
// way, just in one call). A no-op when there are no BLAS.
pub(in crate::metal) fn use_blas_resident_fragment(
    enc: &ProtocolObject<dyn MTLRenderCommandEncoder>,
    blas: &[Retained<ProtocolObject<dyn MTLAccelerationStructure>>],
) {
    if blas.is_empty() {
        return;
    }
    let res: Vec<NonNull<ProtocolObject<dyn MTLResource>>> = blas
        .iter()
        .map(|b| NonNull::from(ProtocolObject::from_ref(&**b)))
        .collect();
    // SAFETY: `res` is a non-empty, contiguous array of `res.len()` live resource
    // pointers; the encoder reads it for the duration of the call only.
    unsafe {
        enc.useResources_count_usage_stages(
            NonNull::new(res.as_ptr() as *mut NonNull<ProtocolObject<dyn MTLResource>>)
                .expect("non-empty blas slice has a non-null pointer"),
            res.len(),
            MTLResourceUsage::Read,
            MTLRenderStages::Fragment,
        );
    }
}

// Attach a completion handler that logs the first GPU fault on an async RT
// command buffer (`what` names the stage: skin compute or BLAS/TLAS build). The
// per-frame skinned rebuild commits these without `waitUntilCompleted`, so a
// fault can no longer be caught synchronously by `check_build_status`; this
// surfaces it (once per process, so a wedged GPU does not spam) instead of
// leaving only the downstream trace victim to report.
fn attach_async_fault_logger(
    cmd: &ProtocolObject<dyn objc2_metal::MTLCommandBuffer>,
    what: &'static str,
) {
    static LOGGED: std::sync::atomic::AtomicBool = std::sync::atomic::AtomicBool::new(false);
    let handler = block2::RcBlock::new(
        move |cb: NonNull<ProtocolObject<dyn objc2_metal::MTLCommandBuffer>>| {
            // SAFETY: Metal hands the completion handler a live command buffer, and the borrow does
            // not escape the block.
            let cb = unsafe { cb.as_ref() };
            if cb.status() == MTLCommandBufferStatus::Error
                && !LOGGED.swap(true, std::sync::atomic::Ordering::Relaxed)
            {
                tracing::error!("RT {what} faulted (async): {:?}", cb.error());
            }
        },
    );
    // SAFETY: addCompletedHandler copies the block, so the RcBlock may drop here.
    unsafe {
        cmd.addCompletedHandler(block2::RcBlock::as_ptr(&handler));
    }
}

// Identity matrix used as a one-joint fallback palette so a skinned object with
// no pose yet still has a valid (undeformed) palette to dispatch against.
const IDENTITY4: [[f32; 4]; 4] = [
    [1.0, 0.0, 0.0, 0.0],
    [0.0, 1.0, 0.0, 0.0],
    [0.0, 0.0, 1.0, 0.0],
    [0.0, 0.0, 0.0, 1.0],
];

// Run the `rt_skin` compute pass: deform each skinned object's bind-pose
// vertices into `deformed_verts` (posed, model-space, 56-byte `Vertex` layout)
// using its joint palette. Runs on its OWN command buffer, committed and
// waited, so the deformed buffer is complete before the acceleration-structure
// build reads it: an AS build does not synchronize against a prior compute
// pass that wrote its input vertex buffer (it is outside the normal encoder
// hazard tracking), so without this wait the build would race the skinning and
// bake a BLAS from half-written vertices.
fn dispatch_skin(
    device: &ProtocolObject<dyn objc2_metal::MTLDevice>,
    command_queue: &ProtocolObject<dyn objc2_metal::MTLCommandQueue>,
    skinned: &SkinnedRtInputs,
    skinned_objects: &[usize],
    deformed_verts: &ProtocolObject<dyn MTLBuffer>,
) -> Result<(), String> {
    let skin_cmd = command_queue
        .commandBuffer()
        .ok_or("failed to create RT skin command buffer")?;
    let cenc = skin_cmd
        .computeCommandEncoder()
        .ok_or("failed to create RT skin compute encoder")?;
    // Transient palette buffers must outlive the GPU work; held until the wait
    // below completes.
    let palette_bufs = encode_skin_dispatch(
        &cenc,
        skinned,
        skinned_objects,
        deformed_verts,
        SkinPalettes::Upload(device),
    )?;
    cenc.endEncoding();
    skin_cmd.commit();
    skin_cmd.waitUntilCompleted();
    drop(palette_bufs);
    check_build_status(&skin_cmd, "skinning compute")
}

// Where the skin dispatch gets each object's joint palette.
enum SkinPalettes<'a> {
    // Upload one transient buffer per object. Used by the one-time seed, which
    // runs before any per-frame palette buffer exists; the caller keeps the
    // returned buffers alive across the commit-and-wait.
    Upload(&'a ProtocolObject<dyn objc2_metal::MTLDevice>),
    // Bind the palette buffers the main and shadow passes already built for this
    // frame's ring slot, which live for the whole frame. `identity` stands in for
    // an object with no pose, so it deforms to bind pose rather than reading
    // whatever that object's slot happens to hold.
    Prebuilt {
        buffers: &'a [Retained<ProtocolObject<dyn MTLBuffer>>],
        identity: &'a Retained<ProtocolObject<dyn MTLBuffer>>,
    },
}

// Encode the `rt_skin` dispatch for each skinned object into `cenc` (setting the
// pipeline state first) and return any transient joint-palette buffers it
// uploaded, which must outlive the GPU work. Empty on the per-frame path, whose
// palettes are owned by the frame. The caller owns command-buffer lifetime
// (commit + wait, or commit + a completion handler).
fn encode_skin_dispatch(
    cenc: &ProtocolObject<dyn objc2_metal::MTLComputeCommandEncoder>,
    skinned: &SkinnedRtInputs,
    skinned_objects: &[usize],
    deformed_verts: &ProtocolObject<dyn MTLBuffer>,
    palettes: SkinPalettes,
) -> Result<Vec<Retained<ProtocolObject<dyn MTLBuffer>>>, String> {
    cenc.set_pipeline(skinned.skin_pipeline);
    let threadgroup = skinned
        .skin_pipeline
        .maxTotalThreadsPerThreadgroup()
        .clamp(1, 64);
    let mut uploaded: Vec<Retained<ProtocolObject<dyn MTLBuffer>>> = Vec::new();
    for &obj_idx in skinned_objects {
        let obj = &skinned.objects[obj_idx];
        let matrices: &[[[f32; 4]; 4]] = skinned
            .joint_matrices
            .get(obj_idx)
            .map(|v| v.as_slice())
            .unwrap_or(&[]);
        // Empty pose -> a single identity joint (undeformed) so the dispatch
        // always has a valid palette to index.
        let joint_count = matrices.len().max(1);
        let palette = match &palettes {
            SkinPalettes::Upload(device) => {
                let slice = if matrices.is_empty() {
                    std::slice::from_ref(&IDENTITY4)
                } else {
                    matrices
                };
                let buf = upload_buffer(device, slice, "RT skin palette")?;
                uploaded.push(buf.clone());
                buf
            }
            SkinPalettes::Prebuilt { buffers, identity } => match buffers.get(obj_idx) {
                Some(buf) if !matrices.is_empty() => buf.clone(),
                _ => (*identity).clone(),
            },
        };
        encode_skin_object(
            cenc,
            SkinDispatchBuffers {
                vertex: skinned.vertex_buffer.as_ref(),
                deformed: deformed_verts,
                palette: palette.as_ref(),
            },
            obj,
            joint_count,
            threadgroup,
        );
    }
    Ok(uploaded)
}

// The three buffers one `rt_skin` dispatch binds: the shared bind-pose vertices
// it reads, the deformed buffer it writes, and the object's joint palette.
#[derive(Clone, Copy)]
struct SkinDispatchBuffers<'a> {
    vertex: &'a ProtocolObject<dyn MTLBuffer>,
    deformed: &'a ProtocolObject<dyn MTLBuffer>,
    palette: &'a ProtocolObject<dyn MTLBuffer>,
}

// Encode one skinned object's `rt_skin` dispatch: one thread per vertex, writing
// its posed model-space `Vertex` into the deformed buffer.
fn encode_skin_object(
    cenc: &ProtocolObject<dyn objc2_metal::MTLComputeCommandEncoder>,
    bufs: SkinDispatchBuffers,
    obj: &SkinnedDrawObject,
    joint_count: usize,
    threadgroup: usize,
) {
    // The RT dispatch runs the base pose; morphing happens in the per-frame main
    // fold. `target_count == 0` leaves the morph slots unread, so a dummy binding
    // satisfies them.
    let params = SkinParams {
        vertex_base: obj.vertex_base,
        vertex_count: obj.vertex_count as u32,
        joint_count: joint_count as u32,
        target_count: 0,
    };
    cenc.set_buffer(bufs.vertex, 0, 0);
    cenc.set_buffer(bufs.deformed, 0, 1);
    cenc.set_buffer(bufs.palette, 0, 2);
    cenc.set_value(&params, 3);
    // Both morph slots take the vertex buffer as a dummy binding; the kernel
    // leaves them unread because `target_count` is zero.
    cenc.set_buffer(bufs.vertex, 0, 4);
    cenc.set_buffer(bufs.vertex, 0, 5);
    cenc.dispatchThreads_threadsPerThreadgroup(
        MTLSize {
            width: obj.vertex_count.max(1),
            height: 1,
            depth: 1,
        },
        MTLSize {
            width: threadgroup,
            height: 1,
            depth: 1,
        },
    );
}

// The per-object buffers the per-frame skin fold binds, both built from this
// frame's ring slot and live for the whole frame.
#[derive(Clone, Copy)]
pub(in crate::metal) struct MainSkinBuffers<'a> {
    pub joints: &'a [Retained<ProtocolObject<dyn MTLBuffer>>],
    pub morph_weights: &'a [Retained<ProtocolObject<dyn MTLBuffer>>],
}

impl crate::metal::context::MtlContext {
    // Per-frame pre-skin for the GPU-driven skinned fold: deform every
    // skinned object's bind-pose vertices into `deformed` (this frame's ring
    // slot) using the per-object joint-palette buffers the main / shadow passes
    // already build for the legacy skinned VS. Reuses the `rt_skin` kernel.
    //
    // Encoded into the Cull pass's command buffer (its own compute encoder),
    // which commits before the Main pass: Metal's automatic hazard tracking then
    // orders this compute write before the main pass's vertex read of `deformed`
    // (the same cross-command-buffer mechanism the static cull → ICB relies on).
    // Unlike the RT seed it binds the pre-built joint and morph-weight buffers
    // instead of uploading transient palettes -- the parallel per-pass encoder
    // cannot keep a transient buffer alive past the worker, while these live for
    // the whole frame. A no-op when the skin pipeline / skinned VB are unset.
    pub(in crate::metal) fn encode_main_skin(
        &self,
        cmd_buf: &ProtocolObject<dyn objc2_metal::MTLCommandBuffer>,
        deformed: &ProtocolObject<dyn MTLBuffer>,
        bufs: MainSkinBuffers<'_>,
    ) -> Result<(), String> {
        let MainSkinBuffers {
            joints: joint_bufs,
            morph_weights: weight_bufs,
        } = bufs;
        let (Some(skin_pipeline), Some(svb)) = (
            self.skinned.skin_pipeline.as_ref(),
            self.skinned.vertex_buffer.as_ref(),
        ) else {
            return Ok(());
        };
        if self.skinned.draw_objects.is_empty() {
            return Ok(());
        }
        let cenc = cmd_buf
            .computeCommandEncoder()
            .ok_or("failed to create main-skin compute encoder")?;
        cenc.set_pipeline(skin_pipeline);
        let tg = skin_pipeline.maxTotalThreadsPerThreadgroup().clamp(1, 64);
        for (i, obj) in self.skinned.draw_objects.iter().enumerate() {
            let Some(joint_buf) = joint_bufs.get(i) else {
                continue;
            };
            // Palette length = this object's matrix count (seeded to >= 1, and
            // `update_skinned_pose` never leaves it empty), matching the buffer
            // the kernel indexes.
            let joint_count = self
                .skinned
                .joint_matrices
                .get(i)
                .map(|m| m.len().max(1))
                .unwrap_or(1);
            // Morphing needs this object's deltas AND this frame's weights.
            // Both slots are bound unconditionally, so a missing one takes the
            // vertex buffer as a dummy and `target_count` goes to zero, which
            // leaves the kernel reading neither.
            let morph = self.skinned.morphs.get(i).and_then(|m| m.as_ref());
            let weights = weight_bufs.get(i);
            let params = SkinParams {
                vertex_base: obj.vertex_base,
                vertex_count: obj.vertex_count as u32,
                joint_count: joint_count as u32,
                target_count: match (morph, weights) {
                    (Some(m), Some(_)) => m.target_count,
                    _ => 0,
                },
            };
            cenc.set_buffer(svb.as_ref(), 0, 0);
            cenc.set_buffer(deformed, 0, 1);
            cenc.set_buffer(joint_buf.as_ref(), 0, 2);
            cenc.set_value(&params, 3);
            cenc.set_buffer(morph.map_or(svb.as_ref(), |m| m.buffer.as_ref()), 0, 4);
            cenc.set_buffer(weights.map_or(svb.as_ref(), |b| b.as_ref()), 0, 5);
            cenc.dispatchThreads_threadsPerThreadgroup(
                MTLSize {
                    width: obj.vertex_count.max(1),
                    height: 1,
                    depth: 1,
                },
                MTLSize {
                    width: tg,
                    height: 1,
                    depth: 1,
                },
            );
        }
        cenc.endEncoding();
        Ok(())
    }
}

// Build the BLAS / TLAS / geometry table for the scene. Returns `None` (not an
// error) when there is no resident triangle geometry to trace: the caller then
// leaves RT disabled and the pass falls back to the base scene.
//
// `skinned`, when present, adds skeletally-animated geometry: a compute pass
// deforms each skinned object's vertices into a fresh model-space buffer, and
// one BLAS per skinned object is built over that buffer. Because
// the pose changes every frame the whole structure is rebuilt per frame (the
// caller forces this); fresh allocations keep it hazard-free.
pub(crate) fn build_rt_accel(
    gpu: RtGpu,
    static_geometry: RtStaticGeometry,
    scene: RtSceneGeometry,
    texture_counts: RtTextureCounts,
    skinned: Option<SkinnedRtInputs>,
    // Layer 2 see-through: when set, see-through glass meshes are left out of the
    // BLAS (they trace their own per-pixel reflection in the transparent pass, and
    // excluding them means glass does not reflect glass). Off keeps every
    // transparent mesh IN the BVH so Layer 1 opaque glass reflects + is reflected
    // like any other surface. Driven by `seethrough_meshes_enabled` (opt-in per
    // `Material::see_through`), not a global flag.
    exclude_seethrough: bool,
) -> Result<Option<RtAccelData>, String> {
    let RtGpu {
        device,
        command_queue,
        frames_in_flight,
    } = gpu;
    let RtStaticGeometry {
        vertex_buffer,
        index_buffer,
    } = static_geometry;
    let RtSceneGeometry {
        draw_objects,
        clusters,
    } = scene;
    let RtTextureCounts { albedo_count } = texture_counts;
    // Only resident draw objects with real triangles take part. When the Layer 2
    // see-through path is enabled, see-through glass meshes are excluded (they
    // route through the transparent pass with their own per-pixel trace, so glass
    // does not reflect glass and the trace never self-hits); otherwise (Layer 1)
    // they stay in so opaque glass reflects + is reflected normally. Track the
    // participating indices into `draw.objects` so a per-frame update re-reads
    // transforms in BLAS-build order.
    let object_indices: Vec<usize> = draw_objects
        .iter()
        .enumerate()
        .filter(|(_, o)| {
            o.resident && o.index_count >= 3 && !(exclude_seethrough && o.material.see_through != 0)
        })
        .map(|(i, _)| i)
        .collect();
    // Instanced clusters that carry real geometry and at least one instance.
    let cluster_list: Vec<&InstancedCluster> = clusters
        .iter()
        .filter(|c| c.index_count >= 3 && !c.instances.is_empty())
        .collect();
    // Skinned objects that are visible and carry real triangles, as indices into
    // the skinned draw list so the skin dispatch finds each one's pose. Clusters
    // and skinned geometry coexist in the BVH; the combination once page-faulted
    // the trace, but that was a per-frame VRAM leak (no autorelease pool around
    // the frame), fixed separately.
    let skinned_list: &[SkinnedDrawObject] = skinned.as_ref().map_or(&[], |s| s.objects);
    let skinned_objects: Vec<usize> = skinned_list
        .iter()
        .enumerate()
        .filter(|(_, o)| o.visible && o.index_count >= 3)
        .map(|(i, _)| i)
        .collect();
    if object_indices.is_empty() && cluster_list.is_empty() && skinned_objects.is_empty() {
        return Ok(None);
    }
    let objects: Vec<&DrawObject> = object_indices.iter().map(|&i| &draw_objects[i]).collect();

    // Deformed-vertex buffer for skinned geometry: the `rt_skin` kernel writes
    // posed model-space `Vertex`s here, mirroring the skinned vertex buffer's
    // indexing so the skinned index buffer addresses it directly. Sized to
    // the highest vertex the skinned objects reach; a 1-vertex dummy when there
    // is no skinned geometry (so the encoder always has a buffer to bind).
    let deformed_extent: usize = skinned_objects
        .iter()
        .map(|&i| skinned_list[i].vertex_base as usize + skinned_list[i].vertex_count)
        .max()
        .unwrap_or(0);
    let deformed_bytes = (deformed_extent * VERTEX_STRIDE).max(VERTEX_STRIDE);
    // Shared, not Private: the buffer is written by the skin compute pass and
    // then read both by the acceleration-structure build and (per hit) by the
    // reflection fragment shader, which run in *separate* command buffers. A
    // Private buffer in that cross-command-buffer producer/consumer pattern was
    // observed to GPU page-fault on the fragment read under the parallel per-
    // pass encoder; Shared is always host-resident and coherent, sidestepping it.
    //
    // The 1-vertex dummy is allocated unconditionally: it is what the encoder
    // binds whenever no skinned geometry is traced, both here and after the
    // per-frame update stops publishing its ring slot.
    let deformed_dummy = device
        .newBufferWithLength_options(VERTEX_STRIDE, MTLResourceOptions::StorageModeShared)
        .ok_or("failed to allocate RT deformed-vertex dummy buffer")?;
    let deformed_verts = if skinned_objects.is_empty() {
        deformed_dummy.clone()
    } else {
        device
            .newBufferWithLength_options(deformed_bytes, MTLResourceOptions::StorageModeShared)
            .ok_or("failed to allocate RT deformed-vertex buffer")?
    };
    // The shared skinned index buffer the kernel + skinned BLAS address; a
    // dummy when there is no skinned geometry. The dummy is one u32 rather than
    // one u16 because the trace reads the buffer as packed u32 words (two
    // indices each) -- Metal API validation rejects a 2-byte buffer bound to a
    // 4-byte element type, even where the branch that reads it never runs.
    let skinned_indices: Retained<ProtocolObject<dyn MTLBuffer>> = match &skinned {
        Some(s) if !skinned_objects.is_empty() => s.index_buffer.clone(),
        _ => device
            .newBufferWithLength_options(
                std::mem::size_of::<u32>(),
                MTLResourceOptions::StorageModePrivate,
            )
            .ok_or("failed to allocate RT skinned-index dummy buffer")?,
    };

    // One BLAS per draw object, then one per cluster, then one per skinned
    // object. `blas[i]` for i < draw_blas_count is draw object i; the next
    // `cluster_list.len()` are clusters; the rest are skinned objects.
    let draw_blas_count = objects.len();
    let skinned_blas_base = draw_blas_count + cluster_list.len();
    let mut prim_descs: Vec<Retained<MTLPrimitiveAccelerationStructureDescriptor>> =
        Vec::with_capacity(skinned_blas_base + skinned_objects.len());
    for obj in &objects {
        prim_descs.push(prim_desc_for(
            vertex_buffer,
            index_buffer,
            obj.base_vertex as usize,
            obj.index_offset,
            obj.index_count,
            MTLIndexType::UInt32,
            MTLAccelerationStructureUsage::None,
        ));
    }
    for c in &cluster_list {
        prim_descs.push(prim_desc_for(
            vertex_buffer,
            index_buffer,
            0,
            c.index_offset,
            c.index_count,
            MTLIndexType::UInt32,
            MTLAccelerationStructureUsage::None,
        ));
    }
    // Skinned BLAS trace the deformed buffer (absolute indices, base_vertex
    // 0). The buffer's contents are written by the compute pass on the same
    // command buffer below, before this BLAS builds.
    for &i in &skinned_objects {
        let obj = &skinned_list[i];
        prim_descs.push(prim_desc_for(
            deformed_verts.as_ref(),
            skinned_indices.as_ref(),
            0,
            obj.index_offset,
            obj.index_count,
            MTLIndexType::UInt32,
            MTLAccelerationStructureUsage::Refit,
        ));
    }

    // Allocate each BLAS and track the largest scratch requirement so a single
    // shared scratch buffer covers the whole build (reused serially).
    let mut blas: Vec<Retained<ProtocolObject<dyn MTLAccelerationStructure>>> =
        Vec::with_capacity(prim_descs.len());
    let mut max_scratch: usize = 0;
    for prim in &prim_descs {
        let sizes = device.accelerationStructureSizesWithDescriptor(prim);
        let acc = device
            .newAccelerationStructureWithSize(sizes.accelerationStructureSize)
            .ok_or("failed to allocate BLAS")?;
        max_scratch = max_scratch.max(sizes.buildScratchBufferSize);
        blas.push(acc);
    }

    // The geometry table is indexed PER INSTANCE, by the intersector's
    // `instance_id`, which is the instance's position in the instance buffer
    // (NOT the `accelerationStructureIndex`). So there is exactly one entry per
    // TLAS instance, in instance order: draw objects, then every cluster
    // instance, then skinned objects.
    let mut instance_descs: Vec<MTLAccelerationStructureInstanceDescriptor> = objects
        .iter()
        .enumerate()
        .map(|(i, obj)| instance_desc(obj, i))
        .collect();
    let mut geom_entries: Vec<RtGeomEntry> = objects
        .iter()
        .map(|obj| geom_entry(obj, albedo_count as u32))
        .collect();

    // Clusters: one TLAS instance + one geometry entry per cluster instance, all
    // referencing the cluster's single BLAS (via `accelerationStructureIndex`)
    // but each with its own transform + geometry entry (so per-instance normals
    // are correct). Stored on `RtAccelData` so a per-frame TLAS rebuild
    // re-appends them verbatim (clusters are baked static into the BVH).
    let mut cluster_instances: Vec<MTLAccelerationStructureInstanceDescriptor> = Vec::new();
    let mut cluster_geom: Vec<RtGeomEntry> = Vec::new();
    for (ci, c) in cluster_list.iter().enumerate() {
        let blas_index = (draw_blas_count + ci) as u32;
        for model in &c.instances {
            cluster_instances.push(instance_desc_at(*model, blas_index));
            cluster_geom.push(cluster_geom_entry(c, *model, albedo_count as u32));
        }
    }
    instance_descs.extend_from_slice(&cluster_instances);
    geom_entries.extend_from_slice(&cluster_geom);

    // Skinned objects: one TLAS instance + one geometry entry each (each skinned
    // object has its own BLAS). The deformed verts are in model space, so the
    // instance transform (= the object's model matrix) brings the trace to world
    // space, like the static path.
    for (si, &i) in skinned_objects.iter().enumerate() {
        let obj = &skinned_list[i];
        let blas_index = (skinned_blas_base + si) as u32;
        instance_descs.push(instance_desc_at(obj.model, blas_index));
        geom_entries.push(skinned_geom_entry(obj, albedo_count as u32));
    }

    let instance_buffer = upload_buffer(device, &instance_descs, "RT instance descriptors")?;
    let geom_table = upload_buffer(device, &geom_entries, "RT geometry table")?;

    let tlas_desc = make_tlas_desc(&blas, &instance_buffer, instance_descs.len());
    let tlas_sizes = device.accelerationStructureSizesWithDescriptor(&tlas_desc);
    let tlas = device
        .newAccelerationStructureWithSize(tlas_sizes.accelerationStructureSize)
        .ok_or("failed to allocate TLAS")?;
    // Size the scratch for the largest of every BLAS build and the TLAS build
    // so the per-frame TLAS rebuild can reuse the same buffer.
    max_scratch = max_scratch.max(tlas_sizes.buildScratchBufferSize);

    let scratch = device
        .newBufferWithLength_options(max_scratch.max(1), MTLResourceOptions::StorageModePrivate)
        .ok_or("failed to allocate RT scratch buffer")?;

    // Skin first (on its own committed-and-waited command buffer) so the
    // deformed buffer is complete before the BLAS build reads it; see
    // `dispatch_skin` for why the wait is required.
    if let Some(s) = &skinned
        && !skinned_objects.is_empty()
    {
        dispatch_skin(
            device,
            command_queue,
            s,
            &skinned_objects,
            deformed_verts.as_ref(),
        )?;
    }

    // Build each BLAS in its own acceleration-structure encoder, then the TLAS in
    // a final encoder. Metal does not order (or prevent overlap of) builds within
    // a single encoder, so a one-encoder build of "all BLAS then the TLAS" lets
    // the TLAS read half-built BLAS and lets builds sharing the scratch buffer
    // stomp on it. Separate encoders within the command buffer serialize, which
    // both orders the TLAS after its BLAS and makes the shared scratch safe.
    // Synchronous: wait for the GPU so the structures (and the skinning that feeds
    // the skinned BLAS) are ready before the first frame traces them.
    let cmd = command_queue
        .commandBuffer()
        .ok_or("failed to create RT build command buffer")?;
    for (acc, prim) in blas.iter().zip(prim_descs.iter()) {
        let enc = cmd
            .accelerationStructureCommandEncoder()
            .ok_or("failed to create acceleration-structure encoder")?;
        enc.buildAccelerationStructure_descriptor_scratchBuffer_scratchBufferOffset(
            acc, prim, &scratch, 0,
        );
        enc.endEncoding();
    }
    // The TLAS references every BLAS, all built on earlier encoders, so declare
    // them resident in this encoder.
    let enc = cmd
        .accelerationStructureCommandEncoder()
        .ok_or("failed to create acceleration-structure encoder")?;
    declare_blas_resident(&enc, &blas);
    enc.buildAccelerationStructure_descriptor_scratchBuffer_scratchBufferOffset(
        &tlas, &tlas_desc, &scratch, 0,
    );
    enc.endEncoding();
    cmd.commit();
    cmd.waitUntilCompleted();
    check_build_status(&cmd, "acceleration-structure build")?;

    let cached_models = objects.iter().map(|o| o.model).collect();
    let draw_blas_sigs = objects.iter().map(|o| GeomSig::of(o)).collect();
    let identity_palette = upload_buffer(device, &[IDENTITY4], "RT identity palette")?;

    Ok(Some(RtAccelData {
        blas,
        static_blas_count: skinned_blas_base,
        tlas,
        geom_table,
        object_indices,
        draw_blas_sigs,
        cached_models,
        cluster_instances,
        cluster_geom,
        scratch,
        instance_buffer,
        deformed_verts,
        skinned_indices,
        retire_pool: RetirePool::new(),
        ring: RtFrameRing::new(frames_in_flight),
        deformed_dummy,
        identity_palette,
        head_generation: 0,
        // Everything above is a fresh allocation, not a ring clone, so the first
        // skinned update has to retire it rather than drop it.
        ring_published: false,
        update_scratch: RtUpdateScratch::default(),
    }))
}

// Allocate one BLAS per skinned object over the deformed buffer (absolute
// indices, base_vertex 0), with the descriptors they were sized from and the
// largest build scratch any of them needs. `Refit` usage is required at build
// time for the per-frame in-place refit to be legal.
fn allocate_skinned_blas(
    device: &ProtocolObject<dyn objc2_metal::MTLDevice>,
    deformed_verts: &ProtocolObject<dyn MTLBuffer>,
    skinned_indices: &ProtocolObject<dyn MTLBuffer>,
    shapes: &[SkinnedShape],
) -> Result<SkinnedBlasSet, String> {
    let mut blas = Vec::with_capacity(shapes.len());
    let mut descs = Vec::with_capacity(shapes.len());
    let mut scratch_bytes = 0usize;
    for shape in shapes {
        let prim = prim_desc_for(
            deformed_verts,
            skinned_indices,
            0,
            shape.index_offset,
            shape.index_count,
            MTLIndexType::UInt32,
            MTLAccelerationStructureUsage::Refit,
        );
        let sizes = device.accelerationStructureSizesWithDescriptor(&prim);
        let acc = device
            .newAccelerationStructureWithSize(sizes.accelerationStructureSize)
            .ok_or("failed to allocate skinned BLAS")?;
        acc.setLabel(Some(&crate::metal::pipeline::ns_str("rt_skinned_blas")));
        scratch_bytes = scratch_bytes.max(sizes.buildScratchBufferSize);
        blas.push(acc);
        descs.push(prim);
    }
    Ok(SkinnedBlasSet {
        blas,
        descs,
        scratch_bytes,
    })
}

// One frame's build-scratch requirement: the largest of every skinned BLAS build
// and the TLAS build, which share the slot's one scratch buffer (separate
// encoders serialize them). At least one byte so Metal never sees a zero-length
// buffer. Pure so the sizing is unit-testable.
fn slot_scratch_bytes(blas_scratch: usize, tlas_scratch: usize) -> usize {
    blas_scratch.max(tlas_scratch).max(1)
}

// Whether every participating draw-object index still resolves to a resident,
// real-triangle object. `false` means the draw list changed shape, and the
// caller leaves the structure as-is for this frame (a full rebuild is the path
// that handles a changed object set). Free-standing so it can be called while
// another field of the accel is mutably borrowed.
fn objects_current(object_indices: &[usize], draw_objects: &[DrawObject]) -> bool {
    object_indices.iter().all(|&idx| {
        draw_objects
            .get(idx)
            .is_some_and(|o| o.resident && o.index_count >= 3)
    })
}

// The participating draw objects in BLAS order, without materialising a `Vec`.
// Only meaningful once `objects_current` has passed; a stale index is skipped
// rather than panicking.
fn objects_in_blas_order<'a>(
    object_indices: &'a [usize],
    draw_objects: &'a [DrawObject],
) -> impl Iterator<Item = &'a DrawObject> + Clone {
    object_indices
        .iter()
        .filter_map(move |&idx| draw_objects.get(idx))
}

impl RtAccelData {
    // Keep the static BLAS; rebuild the TLAS + geometry table from current
    // transforms with fresh allocations, then build on a separate command
    // buffer (committed and waited). Fresh allocations mean no prior in-flight
    // frame can observe a half-updated structure: the old TLAS / table stay
    // alive (retained by their command buffers) until those frames complete.
    pub(crate) fn rebuild_tlas(
        &mut self,
        device: &ProtocolObject<dyn objc2_metal::MTLDevice>,
        command_queue: &ProtocolObject<dyn objc2_metal::MTLCommandQueue>,
        draw_objects: &[DrawObject],
        albedo_count: usize,
    ) -> Result<(), String> {
        if !objects_current(&self.object_indices, draw_objects) {
            return Ok(());
        }
        // Freshly-transformed draw-object instances, then the cluster instances
        // re-appended verbatim (clusters are baked static; their BLAS never move
        // in `self.blas`, so the stored `accelerationStructureIndex` stays
        // valid). The geometry table stays per-BLAS: draw entries (per object)
        // then the per-cluster entries. Built into the persistent scratch so the
        // per-frame `Auto` rebuild reuses its capacity.
        let mut scratch = std::mem::take(&mut self.update_scratch);
        scratch.instances.clear();
        scratch.geom.clear();
        for (i, obj) in objects_in_blas_order(&self.object_indices, draw_objects).enumerate() {
            scratch.instances.push(instance_desc(obj, i));
            scratch.geom.push(geom_entry(obj, albedo_count as u32));
        }
        scratch.instances.extend_from_slice(&self.cluster_instances);
        scratch.geom.extend_from_slice(&self.cluster_geom);
        let instance_descs = &scratch.instances;
        let geom_entries = &scratch.geom;

        let instance_buffer = upload_buffer(device, instance_descs, "RT instance descriptors")?;
        let geom_table = upload_buffer(device, geom_entries, "RT geometry table")?;
        let tlas_desc = make_tlas_desc(&self.blas, &instance_buffer, instance_descs.len());
        let sizes = device.accelerationStructureSizesWithDescriptor(&tlas_desc);
        let tlas = device
            .newAccelerationStructureWithSize(sizes.accelerationStructureSize)
            .ok_or("failed to allocate TLAS")?;
        // Reuse the scratch sized at init (the prior frame's build completed
        // before we got here, so it is free) -- but a topology refresh can change
        // the instance count, and a larger TLAS needs more build scratch than the
        // init sizing. Grow it when so. Replacing the handle is safe: this path
        // is synchronous (commit + wait), and an in-flight command buffer that
        // still references the old scratch retains it independently of this Vec.
        if (sizes.buildScratchBufferSize as u64) > self.scratch.length() as u64 {
            self.scratch = device
                .newBufferWithLength_options(
                    sizes.buildScratchBufferSize.max(1),
                    MTLResourceOptions::StorageModePrivate,
                )
                .ok_or("failed to grow RT scratch buffer")?;
        }

        let cmd = command_queue
            .commandBuffer()
            .ok_or("failed to create RT rebuild command buffer")?;
        let enc = cmd
            .accelerationStructureCommandEncoder()
            .ok_or("failed to create acceleration-structure encoder")?;
        // Every BLAS the rebuilt TLAS references was built on an earlier command
        // buffer (none are rebuilt here), so all must be declared resident.
        declare_blas_resident(&enc, &self.blas);
        enc.buildAccelerationStructure_descriptor_scratchBuffer_scratchBufferOffset(
            &tlas,
            &tlas_desc,
            &self.scratch,
            0,
        );
        enc.endEncoding();
        cmd.commit();
        cmd.waitUntilCompleted();
        check_build_status(&cmd, "TLAS rebuild")?;

        self.tlas = tlas;
        self.geom_table = geom_table;
        // Fresh allocations, so a later skinned update has to retire rather than
        // drop them.
        self.ring_published = false;
        // Snapshot the transforms now baked into the TLAS so the next frame's
        // dirty check compares against what was actually built.
        self.cached_models.clear();
        self.cached_models
            .extend(objects_in_blas_order(&self.object_indices, draw_objects).map(|o| o.model));
        self.update_scratch = scratch;
        Ok(())
    }

    // Whether the BVH has no draw-object and no cluster geometry left. After a
    // topology refresh removes the last of both (every chunk streamed out, with
    // no clusters), the caller drops the structure so a later add re-seeds it
    // rather than building a degenerate zero-instance TLAS. (Skinned geometry is
    // handled on its own per-frame path and never reaches the refresh, so it is
    // not consulted here.)
    pub(crate) fn is_empty(&self) -> bool {
        self.object_indices.is_empty() && self.cluster_instances.is_empty()
    }

    // Incrementally bring the draw-object BLAS head in line with the current
    // participating draw set: reuse every BLAS whose geometry is unchanged, build
    // only the new / changed ones, retire the orphans. The cluster + skinned tails
    // of `blas` are preserved verbatim. When `build_tlas` is set (the no-skinned
    // path), also rebuilds the TLAS + geometry table over [refreshed draw head +
    // clusters] in the same command buffer; when clear (the skinned path) only the
    // head is refreshed and the caller's `rebuild_skinned` rebuilds the TLAS over
    // the head + the fresh skinned tail. Used when streamed chunks are
    // added/removed, props are cloned, or a material edit changes RT
    // participation; a full rebuild of every BLAS would be too costly when most
    // are unchanged.
    //
    // Fully asynchronous, mirroring `rebuild_skinned`: NO `waitUntilCompleted`.
    // The new BLAS (and, when `build_tlas`, the TLAS) build on one command buffer
    // committed on the shared queue ahead of this frame's reflection-trace command
    // buffer, ordered by same-queue FIFO commit -- the same mechanism the skinned
    // rebuild and the whole render graph rely on -- so the trace reads
    // fully-built structures with no CPU stall. All outgoing / transient resources
    // (orphan BLAS, the old TLAS + geometry table + instance buffer when
    // `build_tlas`, and the build scratch) are parked in `retire_pool` rather than
    // freed in place: `useResource` declares residency not lifetime, and the build
    // keeps reading the scratch / instance buffer after this returns, so they must
    // outlive the frames whose still-in-flight trace could reach them.
    pub(crate) fn refresh_static_topology(
        &mut self,
        gpu: RtGpu,
        static_geometry: RtStaticGeometry,
        draw_objects: &[DrawObject],
        texture_counts: RtTextureCounts,
        options: RtTopologyRefreshOptions,
    ) -> Result<(), String> {
        let RtGpu {
            device,
            command_queue,
            ..
        } = gpu;
        let RtStaticGeometry {
            vertex_buffer,
            index_buffer,
        } = static_geometry;
        let RtTextureCounts { albedo_count } = texture_counts;
        let RtTopologyRefreshOptions {
            exclude_seethrough,
            build_tlas,
            frame_id,
        } = options;
        // The current participating draw set, by the same predicate as the full
        // build. (Clusters + skinned never change here, so they are not re-filtered.)
        let new_indices: Vec<usize> = draw_objects
            .iter()
            .enumerate()
            .filter(|(_, o)| {
                o.resident
                    && o.index_count >= 3
                    && !(exclude_seethrough && o.material.see_through != 0)
            })
            .map(|(i, _)| i)
            .collect();
        let new_sigs: Vec<GeomSig> = new_indices
            .iter()
            .map(|&i| GeomSig::of(&draw_objects[i]))
            .collect();

        let plan = plan_topology_refresh(
            &self.object_indices,
            &self.draw_blas_sigs,
            &new_indices,
            &new_sigs,
        );

        let old_draw_count = self.object_indices.len();
        // Clusters occupy `blas[old_draw_count..static_blas_count]`; skinned the
        // tail past `static_blas_count`. Both are preserved across the refresh.
        let cluster_count = self.static_blas_count - old_draw_count;

        // Allocate (but do not yet build) a fresh BLAS for every slot the plan did
        // not match to an existing one. Each is parked at its new-slot position so
        // the assembly below can interleave reused and built BLAS in `new_indices`
        // order.
        let mut fresh: Vec<Option<Retained<ProtocolObject<dyn MTLAccelerationStructure>>>> =
            (0..new_indices.len()).map(|_| None).collect();
        let mut build_jobs: Vec<(usize, Retained<MTLPrimitiveAccelerationStructureDescriptor>)> =
            Vec::new();
        let mut max_scratch: usize = 0;
        for (j, reuse) in plan.reuse.iter().enumerate() {
            if reuse.is_some() {
                continue;
            }
            let obj = &draw_objects[new_indices[j]];
            let prim = prim_desc_for(
                vertex_buffer,
                index_buffer,
                obj.base_vertex as usize,
                obj.index_offset,
                obj.index_count,
                MTLIndexType::UInt32,
                MTLAccelerationStructureUsage::None,
            );
            let sizes = device.accelerationStructureSizesWithDescriptor(&prim);
            let acc = device
                .newAccelerationStructureWithSize(sizes.accelerationStructureSize)
                .ok_or("failed to allocate topology-refresh BLAS")?;
            acc.setLabel(Some(&crate::metal::pipeline::ns_str("rt_topology_blas")));
            max_scratch = max_scratch.max(sizes.buildScratchBufferSize);
            fresh[j] = Some(acc);
            build_jobs.push((j, prim));
        }

        // Assemble the new BLAS array: [refreshed draw head, clusters, skinned],
        // pulling each draw slot from the reused old BLAS or its freshly-built one.
        let old_blas = std::mem::take(&mut self.blas);
        let mut new_blas: Vec<Retained<ProtocolObject<dyn MTLAccelerationStructure>>> =
            Vec::with_capacity(new_indices.len() + (old_blas.len() - old_draw_count));
        for (j, reuse) in plan.reuse.iter().enumerate() {
            match reuse {
                Some(k) => new_blas.push(old_blas[*k].clone()),
                None => new_blas.push(fresh[j].clone().expect("fresh BLAS built above")),
            }
        }
        // Clusters then skinned, verbatim.
        for b in &old_blas[old_draw_count..] {
            new_blas.push(b.clone());
        }

        // Outgoing structures / buffers to retire once the frames-in-flight fence
        // clears them. Orphaned draw BLAS go here always: the current (not yet
        // replaced) TLAS, which an in-flight trace may still be reading, references
        // them, and `useResource` is residency not lifetime.
        let mut retire_structures: Vec<Retained<ProtocolObject<dyn MTLAccelerationStructure>>> =
            plan.retire.iter().map(|&k| old_blas[k].clone()).collect();
        let mut retire_buffers: Vec<Retained<ProtocolObject<dyn MTLBuffer>>> = Vec::new();
        drop(old_blas);

        // When asked, rebuild the TLAS + geometry table over the refreshed draw
        // head + the cluster instances (re-appended verbatim), with the current
        // transforms. Skipped if the set is empty (the caller drops the BVH rather
        // than build a degenerate zero-instance TLAS). The structures are
        // allocated here and built on the command buffer below.
        let do_tlas = build_tlas && !(new_indices.is_empty() && cluster_count == 0);
        let tlas_build = if do_tlas {
            let objects: Vec<&DrawObject> = new_indices.iter().map(|&i| &draw_objects[i]).collect();
            let mut instance_descs: Vec<MTLAccelerationStructureInstanceDescriptor> = objects
                .iter()
                .enumerate()
                .map(|(i, obj)| instance_desc(obj, i))
                .collect();
            let mut geom_entries: Vec<RtGeomEntry> = objects
                .iter()
                .map(|obj| geom_entry(obj, albedo_count as u32))
                .collect();
            instance_descs.extend_from_slice(&self.cluster_instances);
            geom_entries.extend_from_slice(&self.cluster_geom);

            let instance_buffer =
                upload_buffer(device, &instance_descs, "RT instance descriptors")?;
            let geom_table = upload_buffer(device, &geom_entries, "RT geometry table")?;
            let tlas_desc = make_tlas_desc(&new_blas, &instance_buffer, instance_descs.len());
            let tlas_sizes = device.accelerationStructureSizesWithDescriptor(&tlas_desc);
            max_scratch = max_scratch.max(tlas_sizes.buildScratchBufferSize);
            let tlas = device
                .newAccelerationStructureWithSize(tlas_sizes.accelerationStructureSize)
                .ok_or("failed to allocate TLAS")?;
            tlas.setLabel(Some(&crate::metal::pipeline::ns_str("rt_tlas")));
            let cached_models: Vec<[[f32; 4]; 4]> = objects.iter().map(|o| o.model).collect();
            Some((tlas, tlas_desc, instance_buffer, geom_table, cached_models))
        } else {
            None
        };

        // Build everything on ONE command buffer, committed WITHOUT waiting. Each
        // new BLAS in its own encoder (Metal does not order builds within an
        // encoder, and they share the scratch); then, when building the TLAS, a
        // final encoder that declares every referenced BLAS resident (all were
        // built on this or an earlier command buffer, so the TLAS build needs the
        // explicit `useResource`, exactly as the full build does).
        if !build_jobs.is_empty() || tlas_build.is_some() {
            let scratch = device
                .newBufferWithLength_options(
                    max_scratch.max(1),
                    MTLResourceOptions::StorageModePrivate,
                )
                .ok_or("failed to allocate topology-refresh scratch buffer")?;
            let cmd = command_queue
                .commandBuffer()
                .ok_or("failed to create topology-refresh command buffer")?;
            cmd.setLabel(Some(&crate::metal::pipeline::ns_str("rt_topology_build")));
            for (j, prim) in &build_jobs {
                let acc = fresh[*j].as_ref().expect("fresh BLAS allocated above");
                let enc = cmd
                    .accelerationStructureCommandEncoder()
                    .ok_or("failed to create acceleration-structure encoder")?;
                enc.buildAccelerationStructure_descriptor_scratchBuffer_scratchBufferOffset(
                    acc, prim, &scratch, 0,
                );
                enc.endEncoding();
            }
            if let Some((tlas, tlas_desc, _, _, _)) = &tlas_build {
                let enc = cmd
                    .accelerationStructureCommandEncoder()
                    .ok_or("failed to create acceleration-structure encoder")?;
                declare_blas_resident(&enc, &new_blas);
                enc.buildAccelerationStructure_descriptor_scratchBuffer_scratchBufferOffset(
                    tlas, tlas_desc, &scratch, 0,
                );
                enc.endEncoding();
            }
            attach_async_fault_logger(&cmd, "RT topology build");
            cmd.commit();
            // The async build keeps reading the scratch after this returns.
            retire_buffers.push(scratch);
        }

        // Swap in the refreshed structures; park the outgoing ones for deferred
        // free. The new BLAS stay owned by `self.blas` (the build references them
        // by residency, not retention), so they must not be dropped here.
        self.blas = new_blas;
        self.static_blas_count = new_indices.len() + cluster_count;
        self.object_indices = new_indices;
        self.draw_blas_sigs = new_sigs;
        // The persistent BLAS head changed identity, so every ring slot's cached
        // TLAS descriptor (which pins the array of referenced structures) is stale.
        self.head_generation = self.head_generation.wrapping_add(1);
        if let Some((tlas, _, instance_buffer, geom_table, cached_models)) = tlas_build {
            retire_structures.push(std::mem::replace(&mut self.tlas, tlas));
            retire_buffers.push(std::mem::replace(&mut self.geom_table, geom_table));
            retire_buffers.push(std::mem::replace(
                &mut self.instance_buffer,
                instance_buffer,
            ));
            // Snapshot the transforms baked into the new TLAS for the next dirty check.
            self.cached_models = cached_models;
            // Fresh allocations, so a later skinned update has to retire rather
            // than drop them.
            self.ring_published = false;
        }
        // When `build_tlas` is clear (the skinned path), `cached_models` is rebuilt
        // by the caller's `rebuild_skinned` over the refreshed `object_indices`.
        if !retire_structures.is_empty() || !retire_buffers.is_empty() {
            self.retire_pool.push(
                frame_id,
                RetiredRt {
                    structures: retire_structures,
                    buffers: retire_buffers,
                },
            );
        }
        Ok(())
    }

    // Per-frame skinned update: keep the persistent static + cluster BLAS,
    // re-skin this frame's pose, update the skinned BLAS, and rebuild the TLAS +
    // geometry table over the static head plus the skinned tail.
    //
    // Fully asynchronous: NO `waitUntilCompleted`. The skin compute (which writes
    // `deformed_verts`) and the BLAS/TLAS build (which reads it) run on separate
    // command buffers, both committed on the shared queue in order (skin, then
    // build) ahead of this frame's reflection-trace command buffer. An
    // acceleration-structure build is outside Metal's automatic hazard tracking,
    // so skin → build and build → trace are ordered by same-queue FIFO commit order,
    // the same mechanism the render graph uses for every cross-pass read.
    // Per-frame GPU stalls are gone; faults are surfaced from completion handlers.
    //
    // Allocation-free in steady state. Every structure and buffer it writes lives
    // in this frame's ring slot (`super::rt_ring`) and is rebuilt in place: the
    // update runs on every frame, so the frames-in-flight fence guarantees the
    // slot's previous writer has retired. The skinned BLAS are re-fit rather than
    // rebuilt while the triangle set is unchanged, with a periodic full rebuild to
    // bound the traversal-quality drift, and the joint palettes are the buffers
    // the main pass already built for this frame.
    //
    // Returns `Ok(())` without touching the structures when the draw list changed
    // shape (a full rebuild handles that), and falls back to `rebuild_tlas` when
    // no skinned object is visible this frame.
    pub(crate) fn rebuild_skinned(
        &mut self,
        gpu: RtGpu,
        draw_objects: &[DrawObject],
        skinned: SkinnedRtInputs,
        joint_buffers: &[Retained<ProtocolObject<dyn MTLBuffer>>],
        texture_counts: RtTextureCounts,
        frame: RtFrame,
    ) -> Result<(), String> {
        let RtGpu {
            device,
            command_queue,
            ..
        } = gpu;
        let RtTextureCounts { albedo_count } = texture_counts;
        if !objects_current(&self.object_indices, draw_objects) {
            return Ok(());
        }
        // Persistent CPU scratch, swapped out so its heap capacity survives the
        // frame and put back before the successful return. An error path loses
        // the capacity, which is acceptable for an exceptional path.
        let mut scratch = std::mem::take(&mut self.update_scratch);
        let RtUpdateScratch {
            skinned: skinned_objects,
            shapes,
            instances,
            geom,
        } = &mut scratch;

        skinned_objects.clear();
        skinned_objects.extend(
            skinned
                .objects
                .iter()
                .enumerate()
                .filter(|(_, o)| o.visible && o.index_count >= 3)
                .map(|(i, _)| i),
        );
        // No skinned geometry visible this frame: stop publishing the ring's
        // structures (nothing may keep binding a slot a later skinned frame will
        // rewrite), keep the static BLAS, and just refresh the TLAS from current
        // transforms (the static path).
        if skinned_objects.is_empty() {
            self.update_scratch = scratch;
            self.release_skinned();
            return self.rebuild_tlas(device, command_queue, draw_objects, albedo_count);
        }

        // The deformed buffer mirrors the skinned vertex buffer's indexing, so it
        // spans the highest vertex any visible skinned object reaches.
        let deformed_extent = skinned_objects
            .iter()
            .map(|&i| skinned.objects[i].vertex_base as usize + skinned.objects[i].vertex_count)
            .max()
            .unwrap_or(0);
        let deformed_bytes = (deformed_extent * VERTEX_STRIDE).max(VERTEX_STRIDE);

        shapes.clear();
        shapes.extend(skinned_objects.iter().map(|&i| SkinnedShape {
            index_offset: skinned.objects[i].index_offset,
            index_count: skinned.objects[i].index_count,
        }));

        // TLAS instances + geometry table, in instance order: static draw objects
        // (current transforms), then the cluster instances verbatim, then one per
        // skinned object. Skinned BLAS follow the static/cluster head, so their
        // `accelerationStructureIndex` is `static_blas_count + si`. Built before
        // the ring slot is borrowed so the reads of `self` stay disjoint from it.
        let static_blas_count = self.static_blas_count;
        let head_generation = self.head_generation;
        instances.clear();
        geom.clear();
        for (i, obj) in objects_in_blas_order(&self.object_indices, draw_objects).enumerate() {
            instances.push(instance_desc(obj, i));
            geom.push(geom_entry(obj, albedo_count as u32));
        }
        instances.extend_from_slice(&self.cluster_instances);
        geom.extend_from_slice(&self.cluster_geom);
        for (si, &oi) in skinned_objects.iter().enumerate() {
            let obj = &skinned.objects[oi];
            instances.push(instance_desc_at(obj.model, (static_blas_count + si) as u32));
            geom.push(skinned_geom_entry(obj, albedo_count as u32));
        }

        let skinned_indices = skinned.index_buffer.clone();
        let slot = self.ring.slot(frame.ring_slot);

        // A (re)grown deformed buffer invalidates every descriptor built over the
        // old one, so it counts as a shape change even when the triangles did not
        // move.
        let (deformed_verts, deformed_fresh) = slot.deformed(device, deformed_bytes)?;
        let shape_changed = deformed_fresh || !slot.shape_matches(shapes);
        if shape_changed {
            slot.set_skinned(
                allocate_skinned_blas(
                    device,
                    deformed_verts.as_ref(),
                    skinned_indices.as_ref(),
                    shapes,
                )?,
                shapes,
            );
        }

        // This frame's instance descriptors + geometry entries, written straight
        // into the slot's upload buffers.
        let instance_buffer = slot.instances(device, std::mem::size_of_val(&instances[..]))?;
        let geom_table = slot.geom_table(device, std::mem::size_of_val(&geom[..]))?;
        write_buffer_slice(&instance_buffer, instances)?;
        write_buffer_slice(&geom_table, geom)?;

        // The TLAS descriptor pins the BLAS array, the instance buffer and the
        // instance count; while none of those change the cached one drives every
        // rebuild, so the per-frame `Vec` of BLAS references is only built when
        // something actually moved.
        let key = TlasKey {
            head_generation,
            slot_generation: slot.generation(),
            instance_count: instances.len(),
        };
        let cached = slot.tlas_desc(key);
        let tlas_desc = match cached {
            Some(desc) => desc,
            None => {
                let refs: Vec<&ProtocolObject<dyn MTLAccelerationStructure>> = self.blas
                    [..static_blas_count]
                    .iter()
                    .map(|b| b.as_ref())
                    .chain(slot.skinned_blas().iter().map(|b| b.as_ref()))
                    .collect();
                let desc = make_tlas_desc_from_refs(&refs, &instance_buffer, instances.len());
                slot.set_tlas_desc(key, desc.clone());
                desc
            }
        };
        let tlas_sizes = device.accelerationStructureSizesWithDescriptor(&tlas_desc);
        let tlas = slot.tlas(device, tlas_sizes.accelerationStructureSize)?;
        let scratch_buffer = slot.scratch(
            device,
            slot_scratch_bytes(slot.blas_scratch(), tlas_sizes.buildScratchBufferSize),
        )?;

        // Stage 1: skin compute on its own command buffer, committed WITHOUT
        // waiting. Same-queue commit order runs it before the build below (which
        // reads the deformed buffer it writes), the same FIFO ordering the build →
        // trace step and the whole render graph rely on. The palettes it binds are
        // this frame's pre-built joint buffers, which live for the whole frame, so
        // nothing transient has to outlive the async dispatch. A fault can no
        // longer be caught synchronously, so it is logged from a completion handler.
        {
            let skin_cmd = command_queue
                .commandBuffer()
                .ok_or("failed to create RT skin command buffer")?;
            skin_cmd.setLabel(Some(&crate::metal::pipeline::ns_str("rt_skin")));
            let cenc = skin_cmd
                .computeCommandEncoder()
                .ok_or("failed to create RT skin compute encoder")?;
            encode_skin_dispatch(
                &cenc,
                &skinned,
                skinned_objects,
                deformed_verts.as_ref(),
                SkinPalettes::Prebuilt {
                    buffers: joint_buffers,
                    identity: &self.identity_palette,
                },
            )?;
            cenc.endEncoding();
            attach_async_fault_logger(&skin_cmd, "skinning compute");
            skin_cmd.commit();
        }

        // Settle build-or-refit last, once every fallible step above has passed:
        // recording a build the encoder never ran would leave the slot claiming a
        // tree a later refit could not update.
        let update = slot.plan_blas_update(shape_changed);

        // Stage 2: skinned BLAS + TLAS update, committed WITHOUT waiting:
        // same-queue commit order runs it after the skin compute above and before
        // this frame's reflection trace (committed later on the shared queue), the
        // same FIFO ordering the render graph relies on for every cross-pass read.
        //
        // Each BLAS gets its OWN acceleration-structure encoder. Metal does not
        // guarantee the order (or non-overlap) of builds within a single encoder,
        // so a TLAS that references BLAS built in the same encoder can read them
        // half-built, and builds sharing one scratch buffer can stomp on it.
        // Separate encoders within the command buffer serialize, which both orders
        // the TLAS after its BLAS and makes the shared scratch safe to reuse. (The
        // static-only `rebuild_tlas` path never hit this because its BLAS were all
        // built on earlier command buffers.)
        {
            let cmd = command_queue
                .commandBuffer()
                .ok_or("failed to create RT skinned rebuild command buffer")?;
            cmd.setLabel(Some(&crate::metal::pipeline::ns_str("rt_build")));
            for (acc, prim) in slot.skinned_blas().iter().zip(slot.skinned_descs()) {
                let enc = cmd
                    .accelerationStructureCommandEncoder()
                    .ok_or("failed to create acceleration-structure encoder")?;
                match update {
                    BlasUpdate::Build => {
                        enc.buildAccelerationStructure_descriptor_scratchBuffer_scratchBufferOffset(
                            acc,
                            prim,
                            &scratch_buffer,
                            0,
                        );
                    }
                    // A nil destination refits in place, which is legal here
                    // because the slot's previous writer has retired and the
                    // structure was built with `MTLAccelerationStructureUsage::Refit`.
                    // SAFETY: `acc` and `scratch_buffer` are owned by the ring for
                    // longer than this command buffer runs, the descriptor is the
                    // one `acc` was built from, and the scratch was sized from that
                    // same descriptor's reported build size.
                    BlasUpdate::Refit => unsafe {
                        enc.refitAccelerationStructure_descriptor_destination_scratchBuffer_scratchBufferOffset(
                            acc,
                            prim,
                            None,
                            Some(&scratch_buffer),
                            0,
                        );
                    },
                }
                enc.endEncoding();
            }
            // The TLAS, in its own encoder after every BLAS is updated. It
            // references the persistent static/cluster head AND this frame's
            // skinned BLAS, all built on earlier encoders / command buffers, so
            // every one must be declared resident here.
            let enc = cmd
                .accelerationStructureCommandEncoder()
                .ok_or("failed to create acceleration-structure encoder")?;
            declare_blas_resident(
                &enc,
                self.blas[..static_blas_count]
                    .iter()
                    .chain(slot.skinned_blas().iter()),
            );
            enc.buildAccelerationStructure_descriptor_scratchBuffer_scratchBufferOffset(
                &tlas,
                &tlas_desc,
                &scratch_buffer,
                0,
            );
            enc.endEncoding();
            attach_async_fault_logger(&cmd, "skinned BLAS + TLAS build");
            cmd.commit();
        }

        // Publish this slot's structures. Only the skinned tail of `blas` rotates;
        // the static/cluster head is untouched. Handles that were NOT ring-owned
        // (the seed build's, or a topology refresh's) are parked in the retire pool
        // rather than dropped, because a prior in-flight frame's trace can still
        // reach them; once the ring owns them there is nothing left to retire.
        let takeover = !self.ring_published;
        let old_skinned = if takeover {
            self.blas.split_off(static_blas_count)
        } else {
            self.blas.truncate(static_blas_count);
            Vec::new()
        };
        self.blas.extend(slot.skinned_blas().iter().cloned());
        let old_tlas = std::mem::replace(&mut self.tlas, tlas);
        let old_geom_table = std::mem::replace(&mut self.geom_table, geom_table);
        let old_deformed = std::mem::replace(&mut self.deformed_verts, deformed_verts);
        if takeover {
            let mut structures = old_skinned;
            structures.push(old_tlas);
            self.retire_pool.push(
                frame.id,
                RetiredRt {
                    structures,
                    buffers: vec![old_geom_table, old_deformed],
                },
            );
        }
        self.skinned_indices = skinned_indices;
        self.cached_models.clear();
        self.cached_models
            .extend(objects_in_blas_order(&self.object_indices, draw_objects).map(|o| o.model));
        self.ring_published = true;
        self.update_scratch = scratch;
        Ok(())
    }

    // Stop publishing the ring's skinned structures: drop the skinned BLAS tail,
    // fall back to the persistent dummy deformed buffer, and let every slot forget
    // the trees it built. A ring slot may be rewritten in place only because the
    // frame that wrote it is the only frame that binds it, so the moment the
    // skinned path stops running its handles have to go with it.
    fn release_skinned(&mut self) {
        self.blas.truncate(self.static_blas_count);
        self.deformed_verts = self.deformed_dummy.clone();
        self.ring.release_all();
    }

    // Drop resources parked by prior skinned rebuilds that the frames-in-flight
    // fence now guarantees no in-flight frame can still read (`depth` =
    // frames-in-flight; see [`RetirePool::collect`]). Called once per frame.
    pub(crate) fn retire_completed(&mut self, frame_id: u64, depth: usize) {
        self.retire_pool.collect(frame_id, depth as u64);
    }

    // Whether any participating object's model matrix differs from the one
    // baked into the current TLAS. The cheap per-frame check that gates the
    // `Auto` rebuild so a static scene never rebuilds. A changed draw-list
    // shape (missing index) reads as dirty: the conservative answer.
    pub(crate) fn transforms_dirty(&self, draw_objects: &[DrawObject]) -> bool {
        models_dirty(&self.object_indices, &self.cached_models, |idx| {
            draw_objects.get(idx).map(|o| o.model)
        })
    }
}

// Pure dirty test: true if `current(idx)` differs from the cached model for any
// `(idx, cached)` pair, or `current` has no entry for an index. Split out from
// `transforms_dirty` so it is unit-testable without a `DrawObject`.
fn models_dirty(
    object_indices: &[usize],
    cached_models: &[[[f32; 4]; 4]],
    current: impl Fn(usize) -> Option<[[f32; 4]; 4]>,
) -> bool {
    if object_indices.len() != cached_models.len() {
        return true;
    }
    object_indices
        .iter()
        .zip(cached_models.iter())
        .any(|(&idx, cached)| current(idx) != Some(*cached))
}

// Build the compute pipeline that deforms skinned vertices for ray tracing
// (`rt_skin.slang`). Compiled only when RT reflections are on and the GPU
// supports ray tracing, alongside the reflection pipelines.
pub(crate) fn build_rt_skin_pipeline(
    device: &ProtocolObject<dyn objc2_metal::MTLDevice>,
    hot_reload: bool,
) -> Result<Retained<ProtocolObject<dyn objc2_metal::MTLComputePipelineState>>, String> {
    use objc2_metal::{MTLDevice as _, MTLLibrary as _};
    let library = crate::metal::slang_shaders::RT_SKIN.library(device, hot_reload)?;
    let func = library
        .newFunctionWithName(&crate::metal::pipeline::ns_str("rt_skin"))
        .ok_or("rt_skin kernel not found")?;
    device
        .newComputePipelineStateWithFunction_error(&func)
        .map_err(|e| format!("failed to create RT skin pipeline: {:?}", e))
}

// Fail if a command buffer faulted on the GPU. `waitUntilCompleted` returns
// regardless of success, so without this a faulted build/skin would leave a
// corrupt structure the trace then reads. Surfacing it as an `Err` lets the
// non-fatal per-frame update skip the frame (keeping the last good BVH) instead
// of tracing garbage. `what` names the stage so a fault points at the actual
// culprit (the skin compute vs the acceleration-structure build) rather than a
// generic message, and a downstream `SubmissionsIgnored` cascade is
// distinguishable from an original fault by its error code.
fn check_build_status(
    cmd: &ProtocolObject<dyn objc2_metal::MTLCommandBuffer>,
    what: &str,
) -> Result<(), String> {
    if cmd.status() == MTLCommandBufferStatus::Error {
        return Err(format!("RT {what} faulted on the GPU: {:?}", cmd.error()));
    }
    Ok(())
}

// Upload a `#[repr(C)]` slice to a new shared GPU buffer.
fn upload_buffer<T: Copy>(
    device: &ProtocolObject<dyn objc2_metal::MTLDevice>,
    data: &[T],
    what: &str,
) -> Result<Retained<ProtocolObject<dyn MTLBuffer>>, String> {
    let bytes = std::mem::size_of_val(data);
    if bytes == 0 {
        // Metal rejects a zero-length buffer, and `data.as_ptr()` on an empty
        // slice is dangling, so there is no byte to copy from. Hand back a
        // 1-byte buffer the GPU never reads instead.
        return device
            .newBufferWithLength_options(1, MTLResourceOptions::StorageModeShared)
            .ok_or_else(|| format!("failed to allocate buffer for {what}"));
    }
    let ptr = std::ptr::NonNull::new(data.as_ptr() as *mut std::ffi::c_void)
        .ok_or_else(|| format!("{what}: null data pointer"))?;
    // SAFETY: `ptr`/`bytes` describe the live, non-empty `data` slice, and
    // Metal copies those bytes into the new buffer before the call returns.
    unsafe {
        device.newBufferWithBytes_length_options(ptr, bytes, MTLResourceOptions::StorageModeShared)
    }
    .ok_or_else(|| format!("failed to allocate buffer for {what}"))
}

#[cfg(test)]
mod tests {
    use super::*;

    #[test]
    fn pack_instance_transform_drops_affine_row_and_keeps_columns() {
        // A model with a clear translation column and a scaled basis. The packed
        // 4x3 keeps the top three rows of each column; the [0,0,0,1] row is gone.
        let model = [
            [2.0, 0.0, 0.0, 0.0], // column 0: scaled x basis
            [0.0, 3.0, 0.0, 0.0], // column 1: scaled y basis
            [0.0, 0.0, 4.0, 0.0], // column 2: scaled z basis
            [5.0, 6.0, 7.0, 1.0], // column 3: translation
        ];
        let p = pack_instance_transform(model);
        assert_eq!(
            (p.columns[0].x, p.columns[0].y, p.columns[0].z),
            (2.0, 0.0, 0.0)
        );
        assert_eq!(
            (p.columns[1].x, p.columns[1].y, p.columns[1].z),
            (0.0, 3.0, 0.0)
        );
        assert_eq!(
            (p.columns[2].x, p.columns[2].y, p.columns[2].z),
            (0.0, 0.0, 4.0)
        );
        // The translation lands in the fourth column, not a transposed row.
        assert_eq!(
            (p.columns[3].x, p.columns[3].y, p.columns[3].z),
            (5.0, 6.0, 7.0)
        );
    }

    // A distinct geometry signature keyed off `tag` (used as the index offset),
    // so two slots with different tags never compare equal.
    fn sig(tag: usize) -> GeomSig {
        GeomSig {
            base_vertex: tag as i32,
            vertex_offset: tag * 100,
            index_offset: tag,
            index_count: 3,
            generation: 0,
        }
    }

    #[test]
    fn topology_plan_reuses_an_unchanged_set() {
        let old_i = [2usize, 5, 7];
        let old_s = [sig(2), sig(5), sig(7)];
        let plan = plan_topology_refresh(&old_i, &old_s, &old_i, &old_s);
        assert_eq!(plan.reuse, vec![Some(0), Some(1), Some(2)]);
        assert!(plan.retire.is_empty());
    }

    #[test]
    fn topology_plan_builds_only_the_added_slot() {
        let old_i = [2usize, 5];
        let old_s = [sig(2), sig(5)];
        let new_i = [2usize, 5, 9];
        let new_s = [sig(2), sig(5), sig(9)];
        let plan = plan_topology_refresh(&old_i, &old_s, &new_i, &new_s);
        // The two existing slots reuse; the new one (9) builds fresh.
        assert_eq!(plan.reuse, vec![Some(0), Some(1), None]);
        assert!(plan.retire.is_empty());
    }

    #[test]
    fn topology_plan_retires_a_removed_slot() {
        let old_i = [2usize, 5, 7];
        let old_s = [sig(2), sig(5), sig(7)];
        let new_i = [2usize, 7];
        let new_s = [sig(2), sig(7)];
        let plan = plan_topology_refresh(&old_i, &old_s, &new_i, &new_s);
        assert_eq!(plan.reuse, vec![Some(0), Some(2)]);
        assert_eq!(plan.retire, vec![1]); // slot 5's old BLAS is orphaned
    }

    #[test]
    fn topology_plan_rebuilds_a_recycled_slot_whose_geometry_moved() {
        // Same draw index, different geometry signature: a chunk slot recycled for
        // a different chunk. The old BLAS must NOT be reused; it is retired and a
        // fresh one is built.
        let old_i = [5usize];
        let old_s = [sig(5)];
        let new_i = [5usize];
        let new_s = [sig(8)]; // moved geometry under the same draw index
        let plan = plan_topology_refresh(&old_i, &old_s, &new_i, &new_s);
        assert_eq!(plan.reuse, vec![None]);
        assert_eq!(plan.retire, vec![0]);
    }

    #[test]
    fn topology_plan_reuses_across_reorder_by_index() {
        // The participating set is the same but its order changed; each slot still
        // reuses its BLAS by draw index (the reuse points at the old position).
        let old_i = [2usize, 5];
        let old_s = [sig(2), sig(5)];
        let new_i = [5usize, 2];
        let new_s = [sig(5), sig(2)];
        let plan = plan_topology_refresh(&old_i, &old_s, &new_i, &new_s);
        assert_eq!(plan.reuse, vec![Some(1), Some(0)]);
        assert!(plan.retire.is_empty());
    }

    #[test]
    fn rt_geom_entry_is_128_bytes() {
        // The kernel's matching struct relies on this exact size + 16-byte
        // alignment for the array stride to agree. tint+roughness fill one
        // float4; metallic + emissive[3] fill the next so the float4x4 model
        // lands on a 16-byte boundary, exactly as MSL lays the struct out
        // (emissive is a `packed_float3` there, matching `[f32; 3]` here).
        assert_eq!(std::mem::size_of::<RtGeomEntry>(), 128);
    }

    #[test]
    fn slot_scratch_covers_the_largest_build_and_never_reaches_zero() {
        // One scratch buffer serves every skinned BLAS and the TLAS, so it takes
        // the larger requirement whichever side it comes from.
        assert_eq!(slot_scratch_bytes(4096, 1024), 4096);
        assert_eq!(slot_scratch_bytes(1024, 4096), 4096);
        // Metal rejects a zero-length buffer, so a scene whose structures need no
        // scratch still asks for a byte.
        assert_eq!(slot_scratch_bytes(0, 0), 1);
    }

    // A resident draw object carrying real triangles, tagged by `generation` so
    // a test can tell two of them apart.
    fn draw_object(generation: u32) -> DrawObject {
        DrawObject {
            vertex_offset: 0,
            vertex_count: 8,
            index_offset: 0,
            index_count: 6,
            base_vertex: 0,
            geometry_generation: generation,
            shader_bucket: 0,
            model: [[0.0; 4]; 4],
            texture_slot: 0,
            normal_map_slot: 0,
            material: crate::gfx::render_types::MaterialUniforms::DEFAULT,
            visible: true,
            resident: true,
            bb_min: [0.0; 3],
            bb_max: [1.0; 3],
            cull_distance: 0.0,
            lod_alternates: Vec::new(),
        }
    }

    #[test]
    fn objects_current_rejects_a_changed_draw_list() {
        let objects = vec![draw_object(0), draw_object(1), draw_object(2)];
        assert!(objects_current(&[0, 2], &objects));
        // An index past the end: the draw list shrank.
        assert!(!objects_current(&[0, 5], &objects));

        // A slot that streamed out is no longer resident.
        let mut evicted = vec![draw_object(0), draw_object(1)];
        evicted[1].resident = false;
        assert!(!objects_current(&[0, 1], &evicted));

        // A degenerate slot carries no triangles to trace.
        let mut degenerate = vec![draw_object(0), draw_object(1)];
        degenerate[1].index_count = 0;
        assert!(!objects_current(&[0, 1], &degenerate));
    }

    #[test]
    fn objects_in_blas_order_follows_the_index_list() {
        let objects = vec![draw_object(0), draw_object(1), draw_object(2)];
        let seen: Vec<u32> = objects_in_blas_order(&[2, 0], &objects)
            .map(|o| o.geometry_generation)
            .collect();
        assert_eq!(seen, vec![2, 0]);
        // A stale index is skipped rather than panicking; `objects_current` is
        // the guard that keeps a caller from reaching this state.
        let seen: Vec<u32> = objects_in_blas_order(&[1, 9], &objects)
            .map(|o| o.geometry_generation)
            .collect();
        assert_eq!(seen, vec![1]);
    }

    #[test]
    fn models_dirty_detects_moves_and_shape_changes() {
        let ident = [
            [1.0, 0.0, 0.0, 0.0],
            [0.0, 1.0, 0.0, 0.0],
            [0.0, 0.0, 1.0, 0.0],
            [0.0, 0.0, 0.0, 1.0],
        ];
        let mut moved = ident;
        moved[3][0] = 5.0; // translate one object along x

        let indices = vec![0usize, 2usize];
        let cached = vec![ident, ident];

        // All transforms unchanged -> not dirty.
        assert!(!models_dirty(&indices, &cached, |idx| match idx {
            0 | 2 => Some(ident),
            _ => None,
        }));
        // One object moved -> dirty.
        assert!(models_dirty(&indices, &cached, |idx| match idx {
            0 => Some(moved),
            2 => Some(ident),
            _ => None,
        }));
        // An index that no longer resolves (draw list shrank) -> dirty.
        assert!(models_dirty(&indices, &cached, |idx| match idx {
            0 => Some(ident),
            _ => None,
        }));
        // A cached/indices length mismatch -> dirty.
        assert!(models_dirty(&[0usize], &cached, |_| Some(ident)));
    }
}