concinnity-device 0.19.119

//! GPU-driven cull support for the Metal frame encoder: per-frame object /
//! draw-args / joint buffer construction, the cull compute pass, and the
//! bindless texture argument buffer.
#![deny(unsafe_op_in_unsafe_fn)]

use concinnity_core::gfx::cull_status::CullStatus;
use concinnity_core::gfx::frustum::Frustum;
use concinnity_core::gfx::lod;
use concinnity_core::gfx::render_types;
use concinnity_core::render::error::{RenderError, RenderResult};
use concinnity_core::render::model_history::HistoryMode;
use concinnity_core::render::shader_programs::metal::bindless_textures;
use concinnity_core::transform::IDENTITY;
use objc2::rc::Retained;
use objc2::runtime::ProtocolObject;
use objc2_foundation::{NSString, ns_string};
use objc2_metal::{
    MTLArgumentEncoder, MTLCommandBuffer as _, MTLComputePassDescriptor, MTLComputePipelineState,
    MTLDevice as _, MTLFunction as _, MTLLibrary as _, MTLRenderPipelineState,
};

use super::builtin_shaders::compute_pipeline;
use super::context::*;
use super::encode::ComputeEncode;
use super::init::pipelines::BucketPipelines;
use super::pipeline::{cull_encode_library, ns_str};
use super::scoped_encoder::ScopedEncoder;
use concinnity_core::render::uniforms::metal::*;

// One shadow pass's GPU-driven views: an ICB with one region of `cull_count()`
// command slots per view (view `v` at base `v * cull_count()`), the argument
// buffer the encode kernel writes it through, and the status buffer each view's
// decision dispatch fills. The three grow in lockstep; empty until first sized.
#[derive(Default)]
pub(crate) struct ShadowViewIcb {
    pub icb: Option<Retained<ProtocolObject<dyn objc2_metal::MTLIndirectCommandBuffer>>>,
    pub arg_buffer: Option<Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>>,
    pub status: Option<Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>>,
    // Command capacity of `icb` across every view; 0 until first built.
    pub capacity: usize,
}

// The views of one set a frame refreshes, as (region, light frustum) pairs, plus
// how many regions the set holds and which of them refresh.
struct ViewSet<'a> {
    views: &'a [(usize, Frustum)],
    region_count: usize,
    region_mask: u32,
}

// All GPU-driven main pass + cull state grouped into one feature unit: the
// main and world-shader pipelines the indirect draws execute under, the
// phase-1 + phase-2 cull pipelines, their indirect command buffers + argument
// encoders/buffers, the per-object status buffer, the two-pass-occlusion
// toggle, and the Hi-Z depth pyramid + the view-projection snapshots the
// occlusion test reprojects through. All `Some`/active only when the world
// has 3D scene content; a UI-only world leaves every field `None` / default.
pub(crate) struct CullState {
    // True when the GPU-driven main pass exists: a world with 3D scene
    // content. The static draw loop then reads each object from the per-frame
    // `GpuObjectData` buffer and the bindless texture pool. False for a world
    // with no scene content, whose Main pass is a bare clear.
    pub bindless: bool,
    // Bucket 0's main-pass PBR and G-buffer pre-pass pipelines. None for a
    // world with no 3D scene content: the Main pass then encodes as a bare
    // clear and every geometry sub-path early-outs.
    pub main_pipeline: Option<BucketPipelines>,
    // Pipelines for the material-referenced world shaders, indexed by
    // `shader_bucket - 1` (bucket 0 is `main_pipeline`). Each executes its
    // bucket's ICB in the main pass and the pre-pass; empty for single-shader
    // worlds. `None` while the bucket's Shader is not resident -- init defers a
    // shader owned by a scene other than the start scene, and
    // `install_world_shader` builds it when that scene pins. Draws carrying a
    // `None` bucket are skipped.
    pub world_pipelines: concinnity_core::render::world_pipelines::WorldPipelines<BucketPipelines>,
    // The phase-1 decision kernel (`cull.hlsl`). `Some` only when `bindless`
    // is set; every other pipeline here is `Some` exactly when it is.
    pub pipeline: Option<Retained<ProtocolObject<dyn MTLComputePipelineState>>>,
    // The ICB encode kernel (`cull_encode.metal`), dispatched after every
    // decision dispatch on the same encoder.
    pub encode_pipeline: Option<Retained<ProtocolObject<dyn MTLComputePipelineState>>>,
    // How many shader buckets this world routes between: the world Shader
    // count, clamped to MAX_SHADER_BUCKETS. Fixed at init.
    pub bucket_count: usize,
    // One indirect command buffer per shader bucket (index = bucket); the
    // cull kernel encodes each record's draw into its bucket's ICB and resets
    // its slot everywhere else. Empty until `ensure_icb_capacity` builds them.
    pub icbs: Vec<Retained<ProtocolObject<dyn objc2_metal::MTLIndirectCommandBuffer>>>,
    // Encoder that writes an ICB array into the encode kernel's argument
    // buffer. Every ICB argument buffer (main, phase 2, shadow, mirror) has the
    // same layout, so this one encoder serves them all.
    pub icb_arg_encoder: Option<Retained<ProtocolObject<dyn MTLArgumentEncoder>>>,
    // Argument buffer holding the encoded references to `icbs`.
    pub icb_arg_buffer: Option<Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>>,
    // Command capacity of each of `icbs`; 0 until first built. `icbs_2` +
    // `status_buffer` grow in lockstep with it.
    pub icb_capacity: usize,
    // Second-pass decision kernel for two-pass occlusion. `Some` whenever
    // `pipeline` is; used only when `two_pass_occlusion` is on.
    pub pipeline_phase2: Option<Retained<ProtocolObject<dyn MTLComputePipelineState>>>,
    // Per-bucket phase-2 (disocclusion) indirect command buffers.
    pub icbs_2: Vec<Retained<ProtocolObject<dyn objc2_metal::MTLIndirectCommandBuffer>>>,
    // Argument buffer wiring `icbs_2` into the encode kernel.
    pub icb_2_arg_buffer: Option<Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>>,
    // Per-object status buffer (one `u32` each): phase-1 cull writes it,
    // phase-2 cull reads it. Private storage, never CPU-touched.
    pub status_buffer: Option<Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>>,
    // Two-pass Hi-Z occlusion toggle, resolved from
    // `PostProcessConfig.occlusion_two_pass`.
    pub two_pass_occlusion: bool,
    // Hi-Z (depth-mip pyramid) used by the cull kernel for occlusion culling.
    // Built at the end of `draw_frame`, consumed by the next frame's cull
    // dispatch (projected through `prev_view_proj`).
    pub hiz: Option<super::hiz::HiZResources>,
    // Previous frame's un-jittered view-projection, captured every Hi-Z build.
    // The next frame's cull kernel projects AABBs through it. Distinct from the
    // velocity pre-pass's `view_history`.
    pub prev_view_proj: [[f32; 4]; 4],
    // This frame's un-jittered view-projection, captured before `execute_graph`
    // so the phase-2 cull can project AABBs against the freshly built pyramid.
    pub cur_view_proj: [[f32; 4]; 4],
    // `false` on the first frame and after a resize; while false the cull
    // kernel skips the Hi-Z test. Flipped `true` after the first build.
    pub hiz_valid: bool,
    // GPU-driven shadow views. Both pipelines are `Some` only on the bindless
    // path with shadows enabled (`bindless && shadow.map_size > 0`). The
    // frustum-only shadow decision kernel (`shadow_pipeline`) writes each
    // view's outcomes into its region of a view set's status buffer, one encode
    // dispatch per set turns them into commands in that set's ICB, and the
    // depth-only `shadow_bindless_pipeline` issues each view's range.
    // `shadow_views` holds the cascades, `spot_views` the spot slices: separate
    // ICBs, so the two shadow passes share no command slots.
    pub shadow_pipeline: Option<Retained<ProtocolObject<dyn MTLComputePipelineState>>>,
    pub shadow_bindless_pipeline: Option<Retained<ProtocolObject<dyn MTLRenderPipelineState>>>,
    pub shadow_views: ShadowViewIcb,
    pub spot_views: ShadowViewIcb,
    // Per-planar-slot mirror cull ICBs. The planar reflection pass re-runs the
    // phase-1 cull kernel with each plane's reflected-camera frustum into its own
    // ICB, so geometry visible only in the reflection (outside the main frustum)
    // is captured. One slot per distinct planar plane (<= MAX_PLANAR_PLANES); each
    // holds an ICB + its argument buffer, grown in lockstep with the main ICB by
    // `ensure_mirror_icb_capacity`. Empty when the world has no planar set or
    // no scene content. The status buffer is shared scratch: the mirror cull is single-pass, so its
    // per-object status is written but never read.
    pub mirror_slots: Vec<MirrorCullSlot>,
    pub mirror_status: Option<Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>>,
    // Command capacity of every `mirror_slots` ICB; 0 until first built.
    pub mirror_icb_capacity: usize,
}

// One planar slot's mirror cull output: an ICB the cull kernel writes the
// reflected-frustum draws into, plus the argument buffer wiring that ICB into
// the kernel. Built by `ensure_mirror_icb_capacity`, consumed by
// `encode_mirror_cull` (writes) + the planar face render (executes).
pub(crate) struct MirrorCullSlot {
    pub icb: Retained<ProtocolObject<dyn objc2_metal::MTLIndirectCommandBuffer>>,
    pub arg_buffer: Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>,
}

// Buffer index the encode kernel expects its `ICBContainer` argument buffer
// at. The argument encoder that fills that argument buffer is built for this
// same index.
pub(super) const CULL_ICB_BUFFER_INDEX: usize = 4;
// Buffer index of the per-object status the decision kernel writes and the
// encode kernel reads back.
const CULL_STATUS_BUFFER_INDEX: usize = 5;
// Buffer index of the encode kernel's inline `EncodeParams`.
const CULL_ENCODE_PARAMS_INDEX: usize = 7;

// The four texture-pool indices the Metal bindless fragment shader
// (`fragment_main_bindless`, main_bindless.hlsl) reads for one surface. Albedo,
// normal, and every optional map share ONE handle-indexed pool (`tex_pool`,
// [real textures..][flat-normal fallback]), which the shader indexes directly
// (`tex_pool[obj.normal_index]`), so each index is the texture's own handle --
// no albedo-count bias. A normal-less draw's `normal` is the flat-normal
// fallback slot (`normal_pool_index`, at `texture_count`). Every index is
// clamped to the `BINDLESS_TEXTURE_COUNT` cap (the slots the host writes, so an
// over-cap index would read past them). Shared by the per-frame static fill
// (`build_object_buffer`) and the init-time instanced records
// (`metal_instance_records`) so a folded instance addresses the pool identically
// to a static object.
pub(super) struct FlatPoolIndices {
    pub albedo: u32,
    pub normal: u32,
    pub emissive: u32,
    pub orm: u32,
}

impl FlatPoolIndices {
    // `rec` with its emissive and ORM indices replaced by the flat-pool ones:
    // the core packers copy the material's raw indices, which Metal caps.
    fn onto(&self, mut rec: render_types::GpuObjectData) -> render_types::GpuObjectData {
        rec.emissive_map_index = self.emissive;
        rec.orm_map_index = self.orm;
        rec
    }
}

pub(super) fn metal_flat_pool_indices(
    texture_count: usize,
    texture_slot: usize,
    normal_map_slot: usize,
    material: &render_types::MaterialUniforms,
) -> FlatPoolIndices {
    use concinnity_core::gfx::render_types::{albedo_pool_index, normal_pool_index};
    let cap = (super::context::BINDLESS_TEXTURE_COUNT as u32).saturating_sub(1);
    let tc = texture_count as u32;
    let clamp = |i: u32| i.min(cap);
    // The emissive / ORM indices already carry their final shared-pool handle
    // (cook / the graphics system resolve them, `0` = unset), so they index the
    // pool directly.
    FlatPoolIndices {
        albedo: clamp(albedo_pool_index(texture_slot, tc)),
        normal: clamp(normal_pool_index(normal_map_slot, tc)),
        emissive: clamp(material.emissive_map_index),
        orm: clamp(material.orm_map_index),
    }
}

// Build the GPU-driven bindless instanced records: one `GpuObjectData` per
// cluster instance, in cluster-then-instance order, addressing the bindless pool
// with the SAME convention as `build_object_buffer`'s static fill (via
// `metal_flat_pool_indices`). Bounds are the cluster's mesh-local AABB
// transformed by each instance's model, so the cull kernel frustum/distance/
// Hi-Z-tests each instance independently. Built once at init (instances are
// placed at world load and never move) and appended to the per-frame object
// buffer after the static records. Kept separate from the shared core
// `instance_object_records` (which the DX/VK folds use) because Metal resolves
// all four texture indices through `metal_flat_pool_indices` and clamps them to
// its fixed-size MSL pool; the addressing itself now matches DX/VK.
pub(super) fn metal_instance_records(
    clusters: &[render_types::InstancedCluster],
    texture_count: usize,
) -> Vec<render_types::GpuObjectData> {
    let total: usize = clusters.iter().map(|c| c.instances.len()).sum();
    let mut records = Vec::with_capacity(total);
    for cluster in clusters {
        let idx = metal_flat_pool_indices(
            texture_count,
            cluster.texture_slot,
            cluster.normal_map_slot,
            &cluster.material,
        );
        for &model in &cluster.instances {
            let rec = render_types::pack_instance_record(cluster, model, idx.albedo, idx.normal);
            records.push(idx.onto(rec));
        }
    }
    records
}

// Build one GPU-driven bindless record for a skinned object. Reuses
// the core `pack_skinned_record` for the padded bind-pose AABB + model +
// material, then overwrites the texture indices with Metal's flat-pool
// convention: the core helper (like `pack_object_record`) leaves the emissive /
// ORM indices raw for DX/VK's in-shader bias, but Metal's bindless shader
// indexes the flat pool directly, so they must be pre-biased + capped via
// `metal_flat_pool_indices` -- the same convention the static + instanced
// records use.
pub(super) fn metal_skinned_record(
    obj: &render_types::SkinnedDrawObject,
    texture_count: usize,
) -> render_types::GpuObjectData {
    let idx = metal_flat_pool_indices(
        texture_count,
        obj.texture_slot,
        obj.normal_map_slot,
        &obj.material,
    );
    idx.onto(render_types::pack_skinned_record(
        obj, idx.albedo, idx.normal,
    ))
}

// The per-frame cull IO buffers: the packed DrawObject records the kernel tests
// and the indirect draw-args table it reads each survivor's index range from.
struct CullSceneBuffers<'a> {
    object_buffer: &'a ProtocolObject<dyn objc2_metal::MTLBuffer>,
    draw_args_buffer: &'a ProtocolObject<dyn objc2_metal::MTLBuffer>,
    // The record set the two buffers hold. Drives the dispatch width, so the
    // kernel never reads past a snapshot the live draw list has since outgrown.
    counts: crate::metal::context::DrawRecordCounts,
}

// The camera the cull kernel tests records against: the frustum planes plus the
// eye position for the distance-based LOD pick.
struct CullView<'a> {
    frustum: &'a Frustum,
    cam_pos: [f32; 3],
}

// The output the cull kernel encodes survivors into: the target ICB, its
// argument buffer, and the per-object status scratch it writes each record's
// cull outcome to.
struct CullOutputTarget<'a> {
    // One ICB per shader bucket, encoded at matching indices in `arg_buf`;
    // the dispatch's `bucket_count` is this slice's length. Single-stream
    // dispatches (the mirror cull) pass one.
    icbs: &'a [Retained<ProtocolObject<dyn objc2_metal::MTLIndirectCommandBuffer>>],
    arg_buf: &'a Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>,
    status: &'a Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>,
}

// One planar mirror's cull: the frame's object records, the mirror frustum and
// reflected eye to test them against, the mirror slot whose ICB receives the
// survivors, and where the dispatch sits in the planar pass's timing span.
pub(in crate::metal) struct MirrorCull<'a> {
    pub(in crate::metal) object_buffer: &'a ProtocolObject<dyn objc2_metal::MTLBuffer>,
    pub(in crate::metal) draw_args_buffer: &'a ProtocolObject<dyn objc2_metal::MTLBuffer>,
    pub(in crate::metal) frustum: &'a Frustum,
    pub(in crate::metal) eye: [f32; 3],
    pub(in crate::metal) slot: usize,
    pub(in crate::metal) timer: super::pass_timing::PassTimer,
}

// Per-dispatch cull knobs: whether to consult the Hi-Z pyramid (off for the
// mirror cull, whose pyramid is the wrong screen space), where the dispatch sits
// in a pass's timing span, and the encoder debug label.
struct CullDispatchOptions<'a> {
    use_hiz: bool,
    timing: super::pass_timing::PassTimer,
    label: &'a NSString,
}

// One thread per slot over a non-uniform grid: no remainder branch beyond the
// kernel's own bounds guard.
fn dispatch_records(
    enc: &ProtocolObject<dyn objc2_metal::MTLComputeCommandEncoder>,
    pipeline: &ProtocolObject<dyn MTLComputePipelineState>,
    slots: usize,
) {
    use objc2_metal::{MTLComputeCommandEncoder as _, MTLComputePipelineState as _, MTLSize};
    let tg = pipeline.maxTotalThreadsPerThreadgroup().clamp(1, 64);
    enc.dispatchThreads_threadsPerThreadgroup(
        MTLSize {
            width: slots,
            height: 1,
            depth: 1,
        },
        MTLSize {
            width: tg,
            height: 1,
            depth: 1,
        },
    );
}

impl MtlContext {
    // Build the current-pose joint-palette buffers for the skinned passes, one
    // per skinned object, from the per-object ring at `ring_slot`. Returns an
    // empty vec when there are no skinned meshes. The ring reuses persistent
    // shared buffers across frames (the frames-in-flight fence guarantees the
    // slot is no longer GPU-read before it is overwritten) instead of minting
    // a fresh buffer per object per frame.
    pub(super) fn build_joint_buffers(
        &mut self,
        ring_slot: usize,
    ) -> RenderResult<Vec<Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>>> {
        self.rings.joint.write_all(
            &self.hw.device,
            ring_slot,
            &self.state.skinned.joint_matrices,
        )
    }

    // Build this frame's per-object morph-weight buffers, from the same ring
    // slot as the joint palettes. The skin kernel reads them as a plain buffer
    // rather than inline constants, so the encoder binds a resource that lives
    // for the whole frame, which a parallel per-pass encoder needs.
    pub(super) fn build_morph_weight_buffers(
        &mut self,
        ring_slot: usize,
    ) -> RenderResult<Vec<Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>>> {
        self.rings.joint.write_weights(
            &self.hw.device,
            ring_slot,
            &self.state.skinned.morph_weights,
        )
    }

    // Build the per-frame `GpuObjectData` buffer the GPU-driven pass consumes:
    // one record per static `DrawObject`, then one per folded instance, then
    // one per folded skinned object, each indexed by the id the draw call
    // passes as `[[base_instance]]`. Returns `None` when the world has nothing
    // to draw. Rebuilt every frame so `update_model` / `update_visibility`
    // changes are reflected; the committed command buffer keeps the transient
    // buffer alive until the GPU is done with it.
    pub(super) fn build_object_buffer(
        &mut self,
        ring_slot: usize,
    ) -> RenderResult<Option<Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>>> {
        if self.cull_count() == 0 {
            return Ok(None);
        }
        let texture_count = self.scene.textures.len();
        // Reuse a persistent scratch Vec across frames; `mem::take` lifts it out
        // so the build loop borrows only `draw.objects` while the ring + device
        // borrows below stay on disjoint fields.
        let mut objects = std::mem::take(&mut self.rings.object_scratch);
        objects.clear();
        for obj in &self.state.draw.objects {
            // Shared-pool indices (each the texture's own handle, a normal-less
            // draw's normal being the flat-normal fallback slot, all clamped to
            // the cap); the identical mapping the folded instance records use, so
            // static + instances address the pool the same way.
            let idx = metal_flat_pool_indices(
                texture_count,
                obj.texture_slot,
                obj.normal_map_slot,
                &obj.material,
            );
            objects.push(idx.onto(render_types::pack_object_record(
                obj, idx.albedo, idx.normal,
            )));
        }
        // Fold the instanced clusters into the same buffer: each instance's
        // pre-built record is appended after the static objects so one cull
        // dispatch + one indirect draw cover both (the ring auto-grows to the
        // written slice). The records are static (built once at init), so this
        // is a memcpy, and every instance draws the cluster base LOD. `objects`
        // was `mem::take`n, so this borrows only `instanced.records`, leaving
        // the other fields free.
        if self.state.draw.n_instances > 0 {
            objects.extend_from_slice(&self.instanced.records);
        }
        // Append a record per skinned object: the compute-deformed
        // geometry draws as rigid static geometry, so it folds into the same
        // cull. Rebuilt every frame (the record's AABB + model follow obj.model,
        // which animates), unlike the cached static instance records.
        if self.state.draw.n_skinned > 0 {
            for obj in &self.state.skinned.draw_objects {
                objects.push(metal_skinned_record(obj, texture_count));
            }
        }
        let result = self.rings.object.write(
            &self.hw.device,
            ring_slot,
            super::context::bytes_of_slice(&objects),
        );
        self.rings.object_scratch = objects;
        result.map(Some)
    }

    // Build the per-frame `GpuDrawArgs` buffer for the GPU-driven cull pass:
    // one record per `DrawObject` (same indexing as the `GpuObjectData`
    // buffer), carrying the indexed-draw arguments the cull kernel encodes
    // into the indirect command buffer plus the per-frame cull-decision bits.
    // Returns `None` when there is nothing to draw. Rebuilt every frame so
    // `update_visibility` / streaming residency changes (*and* per-frame LOD
    // swaps driven by camera distance) take effect.
    pub(super) fn build_draw_args_buffer(
        &mut self,
        cam_pos: [f32; 3],
        ring_slot: usize,
        history: HistoryMode,
    ) -> RenderResult<Option<Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>>> {
        use concinnity_core::gfx::render_types::{GpuDrawArgs, draw_args_flags};
        if self.cull_count() == 0 {
            return Ok(None);
        }
        let n_cull = self.cull_count();
        self.state.model_history.get_mut().begin(history, n_cull);
        // A transparent glass mesh (Layer 2) is disabled in the opaque pass when
        // the RT path is live: it draws in the transparent pass instead. Clearing
        // ENABLED makes the cull kernel reset its ICB slot to a no-op (the same
        // path invisible / non-resident objects take), so it neither draws opaque
        // nor occludes the refraction snapshot. The object keeps its slot, so every
        // parallel cull / object-buffer / prev-model index stays intact.
        let mesh_glass_active = self.mesh_glass_active();
        let mut args = std::mem::take(&mut self.rings.draw_args_scratch);
        args.clear();
        for (i, obj) in self.state.draw.objects.iter().enumerate() {
            // Pick this frame's active LOD by camera distance: the bindless
            // main pass then renders the chosen slice with no shader-side
            // change. Objects with no alternates fall straight through to LOD0.
            let d = lod::camera_distance(obj, cam_pos);
            let (index_offset, index_count) = obj.active_lod(d);
            let opaque_visible =
                obj.visible && !(mesh_glass_active && obj.material.see_through != 0);

            args.push(GpuDrawArgs {
                index_count: index_count as u32,
                index_offset: index_offset as u32,
                base_vertex: obj.base_vertex as u32,
                // The record's shader bucket rides the upper flag bits so the
                // cull kernel can route its command into that bucket's ICB.
                flags: draw_args_flags(opaque_visible, obj.resident, obj.cullable())
                    | render_types::draw_args_bucket_bits(obj.shader_bucket)
                    | self.state.model_history.get_mut().draw_flags(i, i),
            });
        }
        // Append the instances' draw args in the SAME cluster-then-instance
        // order as `instanced.records`, so cull index `draw.objects.len() + k`
        // reads matching object + draw-args records. The cached args are each
        // cluster's base LOD slice, so this is a memcpy; clusters that declare
        // alternates then get their instances' active slice patched over it.
        if self.state.draw.n_instances > 0 {
            let instance_base = args.len();
            args.extend_from_slice(&self.instanced.draw_args);
            if self.instanced.any_lod {
                lod::for_each_instance_lod(
                    &self.instanced.clusters,
                    cam_pos,
                    |record, index_offset, index_count| {
                        let rec = &mut args[instance_base + record];
                        rec.index_offset = index_offset as u32;
                        rec.index_count = index_count as u32;
                    },
                );
            }
        }
        // Skinned draw args: one per skinned object, the active-LOD
        // slice into the skinned index buffer with base_vertex 0 (the
        // deformed buffer mirrors global skinned indexing). Cullable + gated on
        // obj.visible; rebuilt every frame (pose-driven LOD + visibility). The
        // cull kernel routes records at/after `skinned_record_base()` through the
        // skinned index buffer (see encode_cull's `skinned_base`).
        if self.state.draw.n_skinned > 0 {
            let base = args.len();
            for (k, obj) in self.state.skinned.draw_objects.iter().enumerate() {
                let d = lod::skinned_camera_distance(obj, cam_pos);
                let (index_offset, index_count) = obj.active_lod(d);
                args.push(GpuDrawArgs {
                    index_count: index_count as u32,
                    index_offset: index_offset as u32,
                    base_vertex: 0,
                    flags: draw_args_flags(obj.visible, true, true)
                        | self
                            .state
                            .model_history
                            .get_mut()
                            .skinned_flags(base + k, k),
                });
            }
        }
        let result = self.rings.draw_args.write(
            &self.hw.device,
            ring_slot,
            super::context::bytes_of_slice(&args),
        );
        self.rings.draw_args_scratch = args;
        result.map(Some)
    }

    // Encode the GPU-driven cull compute pass: one thread per
    // `DrawObject` frustum/distance-tests the object and either encodes an
    // indexed draw into `cull_icb` or resets that command slot to a no-op.
    // The bindless main pass then issues the whole buffer with one
    // `executeCommandsInBuffer`. A no-op when the cull pipeline / ICB are not
    // set up (non-bindless contexts) or there is no geometry.
    // pub(in crate::metal) so the render-graph executor in
    // metal/graph_exec.rs can dispatch this pass from a CompiledGraph.
    pub(in crate::metal) fn encode_cull(
        &self,
        cmd_buf: &ProtocolObject<dyn objc2_metal::MTLCommandBuffer>,
        object_buffer: &ProtocolObject<dyn objc2_metal::MTLBuffer>,
        draw_args_buffer: &ProtocolObject<dyn objc2_metal::MTLBuffer>,
        frustum: &Frustum,
        cam_pos: [f32; 3],
        counts: crate::metal::context::DrawRecordCounts,
    ) -> RenderResult<()> {
        let (Some(arg_buf), Some(status)) = (&self.cull.icb_arg_buffer, &self.cull.status_buffer)
        else {
            return Ok(());
        };
        if self.cull.icbs.is_empty() {
            return Ok(());
        }
        self.encode_cull_into(
            cmd_buf,
            CullSceneBuffers {
                object_buffer,
                draw_args_buffer,
                counts,
            },
            CullView { frustum, cam_pos },
            CullOutputTarget {
                icbs: &self.cull.icbs,
                arg_buf,
                status,
            },
            CullDispatchOptions {
                use_hiz: true,
                timing: super::pass_timing::PassTimer::Whole(super::pass_timing::PassId::Cull),
                label: ns_string!("cull phase1"),
            },
        )?;
        Ok(())
    }

    // Encode a phase-1 cull for one planar reflection slot: the same kernel as
    // the main cull, but tested against the reflected-camera `frustum` (so
    // geometry visible only in the mirror, outside the main frustum, is captured)
    // and written into that slot's dedicated mirror ICB. Hi-Z occlusion is forced
    // off because the pyramid lives in the main camera's screen space, not the
    // mirror's. A no-op when the slot or the shared mirror status is absent
    // (non-bindless, or `ensure_mirror_icb_capacity` has not run). The planar face
    // render then executes the slot's ICB.
    pub(in crate::metal) fn encode_mirror_cull(
        &self,
        cmd_buf: &ProtocolObject<dyn objc2_metal::MTLCommandBuffer>,
        cull: MirrorCull<'_>,
    ) -> RenderResult<()> {
        let MirrorCull {
            object_buffer,
            draw_args_buffer,
            frustum,
            eye,
            slot,
            timer,
        } = cull;
        let (Some(mirror), Some(status)) =
            (self.cull.mirror_slots.get(slot), &self.cull.mirror_status)
        else {
            return Ok(());
        };
        self.encode_cull_into(
            cmd_buf,
            CullSceneBuffers {
                object_buffer,
                draw_args_buffer,
                counts: self.draw_record_counts(),
            },
            CullView {
                frustum,
                cam_pos: eye,
            },
            CullOutputTarget {
                icbs: std::slice::from_ref(&mirror.icb),
                arg_buf: &mirror.arg_buffer,
                status,
            },
            CullDispatchOptions {
                use_hiz: false,
                timing: timer,
                label: ns_string!("mirror cull"),
            },
        )
    }

    // Shared body of every phase-1 cull dispatch (main + per-planar mirror): one
    // thread per `DrawObject` frustum/distance(/Hi-Z)-tests the record and either
    // encodes an indexed draw into `icb` or resets that slot to a no-op. The caller
    // supplies the target ICB + its argument buffer + a per-object status scratch,
    // whether to consult the Hi-Z pyramid (`use_hiz`), an optional pass-timing id,
    // and an encoder label. A no-op when the cull pipeline is absent (non-bindless)
    // or there is no geometry.
    fn encode_cull_into(
        &self,
        cmd_buf: &ProtocolObject<dyn objc2_metal::MTLCommandBuffer>,
        scene: CullSceneBuffers,
        view: CullView,
        target: CullOutputTarget,
        options: CullDispatchOptions,
    ) -> RenderResult<()> {
        let CullSceneBuffers {
            object_buffer,
            draw_args_buffer,
            counts,
        } = scene;
        let CullView { frustum, cam_pos } = view;
        let CullOutputTarget {
            icbs,
            arg_buf,
            status,
        } = target;
        let CullDispatchOptions {
            use_hiz,
            timing,
            label,
        } = options;
        use objc2_metal::{MTLComputeCommandEncoder as _, MTLResourceUsage};
        let (Some(pipeline), Some(encode)) = (&self.cull.pipeline, &self.cull.encode_pipeline)
        else {
            return Ok(());
        };
        // Static objects + folded instances: the kernel tests one thread per
        // record and encodes survivors into the ICB.
        let object_count = counts.total;
        if object_count == 0 {
            return Ok(());
        }

        // Pack the six already-normalized frustum planes for the kernel.
        let mut planes = [[0.0f32; 4]; 6];
        for (i, p) in frustum.planes.iter().enumerate() {
            planes[i] = [p.normal[0], p.normal[1], p.normal[2], p.d];
        }
        // Hi-Z occlusion metadata + binding. `hiz_enabled` is gated on `use_hiz`
        // (off for the mirror cull, whose pyramid would be the wrong screen
        // space), the per-context `hiz_valid` flag (false on the first frame and
        // right after a resize), and on whether a `HiZResources` exists. The
        // texture is bound unconditionally when present so the kernel's
        // `texture(0)` always resolves; the `hiz_enabled` flag gates the sample.
        let (hiz_tex, hiz_size, hiz_mip_count, hiz_enabled) = match self.cull.hiz.as_ref() {
            Some(h) => (
                Some(h.texture.as_ref()),
                [h.width as f32, h.height as f32],
                h.mip_count,
                if use_hiz && self.cull.hiz_valid {
                    1u32
                } else {
                    0u32
                },
            ),
            None => (None, [1.0, 1.0], 1, 0u32),
        };
        let cull_uniforms = CullUniforms {
            planes,
            cam_pos: [cam_pos[0], cam_pos[1], cam_pos[2], 0.0],
            prev_view_proj: self.cull.prev_view_proj,
            hiz_size,
            hiz_mip_count,
            hiz_enabled,
            object_count: object_count as u32,
            skinned_base: counts.skinned_base as u32,
            // Main + mirror cull write at `tid` (cascade_base 0); the shadow cull
            // is the only path that offsets by cascade.
            cascade_base: 0,
            bucket_count: icbs.len() as u32,
        };

        let cull_pass_desc = MTLComputePassDescriptor::new();
        if let Some(t) = &self.diagnostics.pass_timing {
            t.attach_compute_timer(&cull_pass_desc, timing);
        }
        let enc = ScopedEncoder::new(
            cmd_buf
                .computeCommandEncoderWithDescriptor(&cull_pass_desc)
                .ok_or_else(|| RenderError::Other("failed to get compute encoder".to_string()))?,
            label,
        );
        enc.set_pipeline(pipeline);
        enc.set_buffer(object_buffer, 0, 0);
        enc.set_buffer(draw_args_buffer, 0, 1);
        enc.set_value(&cull_uniforms, 2);
        enc.set_buffer(&self.scene.index_buffer, 0, 3);
        enc.set_buffer(arg_buf, 0, CULL_ICB_BUFFER_INDEX);
        // Per-object cull status: the decision kernel writes it, the encode
        // kernel reads it back. The main cull's status is read again by phase 2
        // under two-pass occlusion; the mirror cull's shared scratch is not.
        enc.set_buffer(status, 0, CULL_STATUS_BUFFER_INDEX);
        // Skinned index buffer at buffer(6): the encode kernel bakes it into
        // the indirect command for records at/after `skinned_base`. Bound
        // unconditionally (Metal requires a buffer the kernel references to be
        // bound even under a never-taken branch); the static index buffer is a
        // harmless placeholder when no skinned mesh is folded (skinned_base ==
        // object_count then, so the skinned branch never fires).
        enc.set_buffer(self.skinned_index_or_placeholder(), 0, 6);
        // Hi-Z depth pyramid at texture(0). Bound directly (not via an
        // argument buffer), so Metal tracks its residency automatically.
        // Always bound when present; `hiz_enabled` decides whether it's read.
        if let Some(tex) = hiz_tex {
            enc.set_texture(tex, 0);
        }
        // The encode kernel writes draw commands into the ICBs through the
        // argument buffer, so each must be declared resident for the pass.
        for icb in icbs {
            enc.useResource_usage(ProtocolObject::from_ref(&**icb), MTLResourceUsage::Write);
        }

        dispatch_records(&enc, pipeline, object_count);
        enc.set_pipeline(encode);
        enc.set_value(
            &EncodeParams {
                object_count: object_count as u32,
                region_count: 1,
                region_mask: 1,
                skinned_base: counts.skinned_base as u32,
                bucket_count: icbs.len() as u32,
                draw_status: CullStatus::DRAWN,
                region_base: 0,
                _pad: 0,
            },
            CULL_ENCODE_PARAMS_INDEX,
        );
        dispatch_records(&enc, encode, object_count);
        Ok(())
    }

    // Encode the phase-2 GPU cull for two-pass occlusion. Runs after the Hi-Z
    // pyramid has been rebuilt mid-frame from phase-1 depth (`encode_hiz_build`
    // dispatched as the `HizBuild` graph pass). One thread per `DrawObject`
    // re-tests the objects phase 1 marked `HIZ_CANDIDATE` against the fresh
    // pyramid, projecting through *this* frame's view-projection
    // (`cull_cur_view_proj`), and marks the visible ones `REDRAW`; the encode
    // dispatch then writes exactly those into `icbs_2`, which `Main2` issues. A
    // no-op when the phase-2 pipeline / ICB are not set up (two-pass off, or
    // non-bindless).
    pub(in crate::metal) fn encode_cull_phase2(
        &self,
        cmd_buf: &ProtocolObject<dyn objc2_metal::MTLCommandBuffer>,
        object_buffer: &ProtocolObject<dyn objc2_metal::MTLBuffer>,
        draw_args_buffer: &ProtocolObject<dyn objc2_metal::MTLBuffer>,
        frustum: &Frustum,
        cam_pos: [f32; 3],
    ) -> RenderResult<u32> {
        use objc2_metal::{MTLComputeCommandEncoder as _, MTLResourceUsage};
        let (Some(pipeline), Some(encode), Some(arg_buf), Some(status), Some(hiz)) = (
            &self.cull.pipeline_phase2,
            &self.cull.encode_pipeline,
            &self.cull.icb_2_arg_buffer,
            &self.cull.status_buffer,
            self.cull.hiz.as_ref(),
        ) else {
            return Ok(0);
        };
        if self.cull.icbs_2.is_empty() {
            return Ok(0);
        }
        // Static objects + folded instances re-tested against the fresh Hi-Z
        // pyramid; same record count as phase 1.
        let object_count = self.cull_count();
        if object_count == 0 {
            return Ok(0);
        }

        // Frustum planes are unused by the phase-2 kernel (candidates already
        // passed the frustum test in phase 1) but the uniform layout is shared,
        // so pack them anyway for a clean struct.
        let mut planes = [[0.0f32; 4]; 6];
        for (i, p) in frustum.planes.iter().enumerate() {
            planes[i] = [p.normal[0], p.normal[1], p.normal[2], p.d];
        }
        // Project AABBs through this frame's un-jittered VP: it matches the
        // pyramid we just rebuilt from this frame's depth. `hiz_enabled = 1`:
        // the `HizBuild` pass always precedes this dispatch in the graph, so a
        // valid pyramid is guaranteed (the kernel still guards defensively).
        let skinned_base = self.skinned_record_base() as u32;
        let cull_uniforms = CullUniforms {
            planes,
            cam_pos: [cam_pos[0], cam_pos[1], cam_pos[2], 0.0],
            prev_view_proj: self.cull.cur_view_proj,
            hiz_size: [hiz.width as f32, hiz.height as f32],
            hiz_mip_count: hiz.mip_count,
            hiz_enabled: 1,
            object_count: object_count as u32,
            skinned_base,
            cascade_base: 0,
            bucket_count: self.cull.icbs_2.len() as u32,
        };

        let cull_pass_desc = MTLComputePassDescriptor::new();
        if let Some(t) = &self.diagnostics.pass_timing {
            t.attach_compute(&cull_pass_desc, super::pass_timing::PassId::Cull2);
        }
        let enc = ScopedEncoder::new(
            cmd_buf
                .computeCommandEncoderWithDescriptor(&cull_pass_desc)
                .ok_or_else(|| RenderError::Other("failed to get compute encoder".to_string()))?,
            ns_string!("cull phase2"),
        );
        enc.set_pipeline(pipeline);
        enc.set_buffer(object_buffer, 0, 0);
        enc.set_buffer(draw_args_buffer, 0, 1);
        enc.set_value(&cull_uniforms, 2);
        enc.set_buffer(&self.scene.index_buffer, 0, 3);
        enc.set_buffer(arg_buf, 0, CULL_ICB_BUFFER_INDEX);
        enc.set_buffer(status, 0, CULL_STATUS_BUFFER_INDEX);
        // Skinned index buffer at buffer(6); see encode_cull. Phase 2
        // of two-pass occlusion re-tests the same records, so the skinned
        // tail is handled here too.
        enc.set_buffer(self.skinned_index_or_placeholder(), 0, 6);
        enc.set_texture(hiz.texture.as_ref(), 0);
        // The encode kernel writes draw commands into the phase-2 ICBs through
        // the argument buffer, so each must be declared resident here too.
        for icb in &self.cull.icbs_2 {
            enc.useResource_usage(ProtocolObject::from_ref(&**icb), MTLResourceUsage::Write);
        }

        dispatch_records(&enc, pipeline, object_count);
        enc.set_pipeline(encode);
        enc.set_value(
            &EncodeParams {
                object_count: object_count as u32,
                region_count: 1,
                region_mask: 1,
                skinned_base,
                bucket_count: self.cull.icbs_2.len() as u32,
                draw_status: CullStatus::REDRAW,
                region_base: 0,
                _pad: 0,
            },
            CULL_ENCODE_PARAMS_INDEX,
        );
        dispatch_records(&enc, encode, object_count);
        Ok(0)
    }

    // Encode the GPU-driven shadow culls: for the cascades and then the spot
    // slices, one shadow decision dispatch per re-rendered view (gated by
    // `shadow.render_mask` / `spot_shadow.render_mask`), each frustum-testing
    // every record against that view's LIGHT frustum into the view's region of
    // its set's status buffer (`cascade_base = v * cull_count()`), then one
    // encode dispatch per set turning the refreshed regions into commands in
    // that set's ICB. Hi-Z + distance are off (frustum only). A no-op when the
    // shadow-bindless path is inactive or there is no geometry.
    //
    // Runs as a compute prologue in the SAME command buffer as the main `Cull`
    // pass (dispatched right after `encode_cull` from the graph executor's Cull
    // arm), so the ICB writes land in a command buffer committed before the
    // `Shadow` and `SpotShadow` render passes' command buffers -- the exact
    // cross-command-buffer FIFO ordering the main cull -> main ICB already
    // relies on, pinned by both passes' graph edge after Cull. No explicit
    // barrier (Metal has none); residency is declared with `useResource`.
    // pub(in crate::metal) so the graph executor can dispatch it.
    pub(in crate::metal) fn encode_shadow_culls(
        &self,
        cmd_buf: &ProtocolObject<dyn objc2_metal::MTLCommandBuffer>,
        object_buffer: &ProtocolObject<dyn objc2_metal::MTLBuffer>,
        draw_args_buffer: &ProtocolObject<dyn objc2_metal::MTLBuffer>,
    ) -> RenderResult<()> {
        use concinnity_core::gfx::render_types::{MAX_SHADOWED_SPOTS, NUM_SHADOW_CASCADES};
        let Some(pipeline) = &self.cull.shadow_pipeline else {
            return Ok(());
        };
        if self.cull_count() == 0
            || (self.cull.shadow_views.icb.is_none() && self.cull.spot_views.icb.is_none())
        {
            return Ok(());
        }

        let cull_pass_desc = MTLComputePassDescriptor::new();
        let enc = ScopedEncoder::new(
            cmd_buf
                .computeCommandEncoderWithDescriptor(&cull_pass_desc)
                .ok_or_else(|| {
                    RenderError::Other("failed to get shadow cull compute encoder".to_string())
                })?,
            ns_string!("shadow cull"),
        );
        enc.set_buffer(object_buffer, 0, 0);
        enc.set_buffer(draw_args_buffer, 0, 1);
        enc.set_buffer(&self.scene.index_buffer, 0, 3);
        // Skinned index buffer at buffer(6); the encode kernel bakes it into
        // the skinned-tail commands exactly like the main cull.
        enc.set_buffer(self.skinned_index_or_placeholder(), 0, 6);

        // Same cascade set the shadow render pass refreshes this frame; a skipped
        // cascade keeps its prior depth slice, so its cull dispatch + ICB region
        // are left untouched.
        let all = (1u32 << NUM_SHADOW_CASCADES) - 1;
        let mask = if self.shadow.render_mask == 0 {
            all
        } else {
            self.shadow.render_mask
        };
        let mut cascades = [(0, Frustum::from_shadow(IDENTITY)); NUM_SHADOW_CASCADES];
        let mut kept = 0;
        for (c, light_vp) in self.shadow.uniforms.light_vps.iter().enumerate() {
            if mask & (1u32 << c) == 0 {
                continue;
            }
            // The caster-extent near push baked into light_vps survives, so
            // off-screen / tall casters are kept.
            cascades[kept] = (c, Frustum::from_shadow(*light_vp));
            kept += 1;
        }
        self.encode_view_culls(
            &enc,
            pipeline,
            &self.cull.shadow_views,
            ViewSet {
                views: &cascades[..kept],
                region_count: NUM_SHADOW_CASCADES,
                region_mask: mask,
            },
        );

        let mut spots = [(0, Frustum::from_shadow(IDENTITY)); MAX_SHADOWED_SPOTS];
        let mut kept = 0;
        let mut spot_mask = 0u32;
        for slice in self.spot_shadow.refreshed_slices() {
            spots[kept] = (slice as usize, self.spot_shadow.frusta[slice as usize]);
            spot_mask |= 1 << slice;
            kept += 1;
        }
        self.encode_view_culls(
            &enc,
            pipeline,
            &self.cull.spot_views,
            ViewSet {
                views: &spots[..kept],
                region_count: self.spot_shadow.count as usize,
                region_mask: spot_mask,
            },
        );
        Ok(())
    }

    // One view set's decision dispatches (one per view, into the view's region
    // of the set's status buffer) and its single encode dispatch over every
    // region, writing the refreshed regions' commands into the set's ICB. The
    // caller has bound the record, index and skinned-index buffers. A no-op for
    // a set with no ICB or no view to refresh.
    fn encode_view_culls(
        &self,
        enc: &ProtocolObject<dyn objc2_metal::MTLComputeCommandEncoder>,
        pipeline: &ProtocolObject<dyn MTLComputePipelineState>,
        set: &ShadowViewIcb,
        views: ViewSet<'_>,
    ) {
        use objc2_metal::{MTLComputeCommandEncoder as _, MTLResourceUsage};
        let (Some(encode), Some(icb), Some(arg_buf), Some(status)) = (
            &self.cull.encode_pipeline,
            &set.icb,
            &set.arg_buffer,
            &set.status,
        ) else {
            return;
        };
        if views.views.is_empty() {
            return;
        }
        let object_count = self.cull_count();
        let skinned_base = self.skinned_record_base() as u32;
        enc.set_pipeline(pipeline);
        enc.set_buffer(arg_buf, 0, CULL_ICB_BUFFER_INDEX);
        enc.set_buffer(status, 0, CULL_STATUS_BUFFER_INDEX);
        // The encode kernel writes draw commands into the ICB through the
        // argument buffer, so it must be resident for the compute pass.
        enc.useResource_usage(ProtocolObject::from_ref(&**icb), MTLResourceUsage::Write);

        for (region, frustum) in views.views {
            let mut planes = [[0.0f32; 4]; 6];
            for (i, p) in frustum.planes.iter().enumerate() {
                planes[i] = [p.normal[0], p.normal[1], p.normal[2], p.d];
            }
            let cull_uniforms = CullUniforms {
                planes,
                // Unused by the shadow kernel (no distance cull); kept zero.
                cam_pos: [0.0; 4],
                // Unused (Hi-Z disabled); identity keeps the struct clean.
                prev_view_proj: IDENTITY,
                hiz_size: [1.0, 1.0],
                hiz_mip_count: 1,
                hiz_enabled: 0,
                object_count: object_count as u32,
                skinned_base,
                cascade_base: (region * object_count) as u32,
                // The shadow cull writes one depth-only stream: icbs[0].
                bucket_count: 1,
            };
            enc.set_value(&cull_uniforms, 2);
            dispatch_records(enc, pipeline, object_count);
        }
        // One encode dispatch over the span of refreshed regions; the mask
        // leaves a skipped region's commands exactly as its last cull left them.
        let (region_base, span) =
            EncodeParams::encoded_span(views.region_mask, views.region_count as u32);
        if span == 0 {
            return;
        }
        enc.set_pipeline(encode);
        enc.set_value(
            &EncodeParams {
                object_count: object_count as u32,
                region_count: views.region_count as u32,
                region_mask: views.region_mask,
                skinned_base,
                bucket_count: 1,
                draw_status: CullStatus::DRAWN,
                region_base,
                _pad: 0,
            },
            CULL_ENCODE_PARAMS_INDEX,
        );
        dispatch_records(enc, encode, span as usize * object_count);
    }

    // The `BindlessTextures` block's members ahead of the pool, with their ids:
    // the shadow maps, the two IBL cubes, the AO (the blurred occlusion when
    // SSAO is on, else 1x1 white), the probe cube array and the LTC tables.
    fn bindless_fixed_members(
        &self,
    ) -> [(usize, &ProtocolObject<dyn objc2_metal::MTLTexture>); bindless_textures::FIXED] {
        use bindless_textures as ids;
        [
            (ids::SHADOW_MAP, &*self.shadow.map),
            (ids::IRRADIANCE_CUBE, &*self.scene.env_map.irradiance),
            (ids::PREFILTER_CUBE, &*self.scene.env_map.prefilter),
            (ids::SSAO, self.ao_output_texture()),
            (ids::PROBE_CUBES, self.probe.cubes.texture()),
            (ids::SPOT_SHADOW_MAP, &*self.spot_shadow.map),
            (ids::LTC_MATRIX, &*self.scene.ltc_matrix_texture),
            (ids::LTC_MAGNITUDE, &*self.scene.ltc_magnitude_texture),
        ]
    }

    // The identity of everything the `BindlessTextures` block names outside the
    // shared pool, plus the pool's own generation. Two frames that agree on
    // this would write byte-identical argument buffers, so a ring slot already
    // holding it can be left alone.
    pub(super) fn bindless_texture_signature(&self) -> u64 {
        let mut sig = super::bindless_args::Signature::new();
        sig.push_u64(self.arg_buffers.texture_epoch);
        sig.push_u64(self.scene.textures.len() as u64);
        for tex in &self.scene.fallback_textures {
            sig.push_texture(tex.as_ref());
        }
        for (_, tex) in self.bindless_fixed_members() {
            sig.push_texture(tex);
        }
        sig.finish()
    }

    // The pool's unused tail, every slot of it the same white texture. It is
    // most of the block on a world that fills a fraction of the pool, and it
    // moves only when the pool's length or that texture does, so it is gated
    // apart from the contents a stream-in changes every time.
    fn bindless_tail_signature(&self) -> u64 {
        let mut sig = super::bindless_args::Signature::new();
        sig.push_u64(self.scene.textures.len() as u64);
        if let Some(tex) = self.scene.fallback_textures.get(1) {
            sig.push_texture(tex.as_ref());
        }
        sig.finish()
    }

    // Build the per-frame `BindlessTextures` argument buffer for the bindless
    // static pass: the fixed members at their ids, then the albedo + normal-map
    // pool (every one of the `BINDLESS_TEXTURE_COUNT` slots filled: overflow and
    // trailing empty slots fall back to the white albedo texture). The bindless
    // fragment shader can only reach textures through this argument buffer
    // because discrete texture bindings make it incompatible with indirect
    // command buffers.
    //
    // The buffer is a ring slot, one per frame in flight, so the GPU never
    // reads a slot the CPU is rewriting. Its contents move only on a texture
    // stream/evict, a probe bake, an env-map swap or a transient-pool repack,
    // and a signature gate per slot skips the write entirely on the frames
    // between: a change is rewritten once into each slot and then stops. `sig`
    // is this frame's `bindless_texture_signature`. `None` for non-bindless
    // contexts. The committed command buffer keeps the buffer alive.
    pub(super) fn build_bindless_texture_args(
        &mut self,
        ring_slot: usize,
        sig: u64,
    ) -> RenderResult<Option<Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>>> {
        use bindless_textures as ids;
        if !self.cull.bindless {
            return Ok(None);
        }
        let tail_sig = self.bindless_tail_signature();
        let count = super::context::BINDLESS_TEXTURE_COUNT;
        let len = super::bindless_args::bindless_block_len(count);
        // Ring slot, grown to the block's length instead of a fresh allocation
        // each frame, and rewritten in place; the fence guarantees the prior
        // user of this slot has retired on the GPU.
        let (buf, allocated) =
            self.rings
                .bindless_tex
                .slot_fresh(&self.hw.device, ring_slot, len)?;
        if allocated {
            self.arg_buffers.bindless_tex_gates.invalidate(ring_slot);
            self.arg_buffers.bindless_tail_gates.invalidate(ring_slot);
        }
        let write_tail = self
            .arg_buffers
            .bindless_tail_gates
            .stale(ring_slot, tail_sig);
        let write_pool = self.arg_buffers.bindless_tex_gates.stale(ring_slot, sig);
        if !write_pool && !write_tail {
            return Ok(Some(buf));
        }
        let mut block = super::bindless_args::ResourceIdWriter::new(&buf);
        let texture_count = self.scene.textures.len();
        if write_pool {
            for (id, tex) in self.bindless_fixed_members() {
                block.set(id, tex);
            }
            // The shared pool: every real texture, then the reserved fallbacks --
            // flat-normal at `texture_count`, white at `texture_count + 1` -- so an
            // over-cap or clamped index still samples a valid texture.
            for (i, tex) in self.scene.textures.iter().take(count).enumerate() {
                block.set(ids::pool(i), tex);
            }
            for (i, tex) in (texture_count..count).zip(&self.scene.fallback_textures) {
                block.set(ids::pool(i), tex);
            }
        }
        if write_tail && let Some(white) = self.scene.fallback_textures.last() {
            // White across the unused tail, past the two reserved fallbacks.
            for i in (texture_count + self.scene.fallback_textures.len()).min(count)..count {
                block.set(ids::pool(i), white);
            }
        }
        Ok(Some(buf))
    }

    // Rebuild the bindless pass's residency set when the textures the
    // `BindlessTextures` block names change, as `sig` (this frame's
    // `bindless_texture_signature`) says. Called once per frame before any pass
    // declares them; a frame whose signature is unchanged does nothing.
    pub(super) fn refresh_bindless_residency(&mut self, sig: u64) {
        // Taken out so the iterator below can borrow the rest of `self`.
        let mut set = core::mem::replace(
            &mut self.arg_buffers.bindless_residency,
            super::bindless_args::ResidencySet::new(),
        );
        set.refresh(
            sig,
            self.scene
                .textures
                .iter()
                .chain(self.scene.fallback_textures.iter())
                .map(|t| t.as_ref())
                .chain(self.bindless_fixed_members().map(|(_, tex)| tex)),
        );
        self.arg_buffers.bindless_residency = set;
    }

    // Declare every texture the bindless pass samples resident for the
    // indirect command buffer. The textures are referenced through the
    // `BindlessTextures` argument buffer rather than bound on the encoder, so
    // the indirect execution cannot see them unless they are explicitly used.
    // One batched call over the set `refresh_bindless_residency` cached: the
    // driver records the same residency either way.
    pub(super) fn use_bindless_textures(
        &self,
        encoder: &ProtocolObject<dyn objc2_metal::MTLRenderCommandEncoder>,
    ) {
        self.arg_buffers
            .bindless_residency
            .declare_fragment(encoder);
    }
}

// The GPU-driven cull's pipelines: the phase-1 and phase-2 decision kernels
// from `cull.hlsl`, the ICB encode kernel from `cull_encode.metal`, and the
// argument encoder every ICB argument buffer is written with (the encode
// kernel is the one function that declares the `ICBContainer`).
pub(super) struct CullPipeline {
    pub decide: Retained<ProtocolObject<dyn MTLComputePipelineState>>,
    pub decide_phase2: Retained<ProtocolObject<dyn MTLComputePipelineState>>,
    pub encode: Retained<ProtocolObject<dyn MTLComputePipelineState>>,
    pub icb_arg_encoder: Retained<ProtocolObject<dyn MTLArgumentEncoder>>,
}

// Build the GPU-driven cull pipelines. The decision kernels run one thread
// per `DrawObject` and record each record's outcome in the status buffer; the
// encode kernel turns a status into an indexed draw in the indirect command
// buffer or resets the slot. The render pass then issues the whole buffer with
// one `executeCommandsInBuffer`, so the CPU never walks the draw list.
pub(super) fn build_cull_pipeline(
    device: &ProtocolObject<dyn objc2_metal::MTLDevice>,
    hot_reload: bool,
) -> RenderResult<CullPipeline> {
    let decide = compute_pipeline(device, &super::builtin_shaders::CULL_PHASE1, hot_reload)?;
    let decide_phase2 = compute_pipeline(device, &super::builtin_shaders::CULL_PHASE2, hot_reload)?;
    let library = cull_encode_library(device, hot_reload)?;
    let encode_fn = library
        .newFunctionWithName(&ns_str("cull_encode"))
        .ok_or_else(|| {
            RenderError::ShaderCompile("cull_encode not found in cull_encode library".to_string())
        })?;
    let encode = device
        .newComputePipelineStateWithFunction_error(&encode_fn)
        .map_err(|e| RenderError::ShaderCompile(format!("cull encode pipeline: {e:?}")))?;
    // SAFETY: CULL_ICB_BUFFER_INDEX is the static buffer index the encode
    // kernel declares its argument-buffer parameter at.
    let icb_arg_encoder =
        unsafe { encode_fn.newArgumentEncoderWithBufferIndex(CULL_ICB_BUFFER_INDEX) };
    Ok(CullPipeline {
        decide,
        decide_phase2,
        encode,
        icb_arg_encoder,
    })
}

// Build the GPU-driven cascaded-shadow decision pipeline (a separate build
// keeps the call site shadow-gated rather than always building it on the
// bindless path). Its outcomes are encoded by the shared encode kernel; the
// render-side depth-only shadow pipeline is built separately in
// `init/pipelines.rs::build_shadow_bindless_pipeline`.
pub(super) fn build_shadow_cull_pipeline(
    device: &ProtocolObject<dyn objc2_metal::MTLDevice>,
    hot_reload: bool,
) -> RenderResult<Retained<ProtocolObject<dyn MTLComputePipelineState>>> {
    compute_pipeline(device, &super::builtin_shaders::CULL_SHADOW, hot_reload)
}

#[cfg(test)]
mod tests {
    use super::metal_flat_pool_indices;
    use concinnity_core::gfx::render_types::{MaterialUniforms, NO_NORMAL_MAP_SLOT};

    #[test]
    fn flat_pool_indices_share_one_handle_indexed_pool() {
        // Albedo, normal, and every optional map index the shared pool by their
        // own handle -- no albedo-count bias. 8 real textures.
        let material = MaterialUniforms {
            emissive_map_index: 3,
            orm_map_index: 0,
            ..MaterialUniforms::DEFAULT
        };
        let idx = metal_flat_pool_indices(8, 2, 1, &material);
        assert_eq!(idx.albedo, 2);
        assert_eq!(idx.normal, 1);
        assert_eq!(idx.emissive, 3);
        assert_eq!(idx.orm, 0);
    }

    #[test]
    fn flat_pool_indices_map_a_missing_normal_to_the_fallback() {
        // A draw with no normal map addresses the flat-normal fallback slot
        // (one past the last real texture), so the shader samples flat (0,0,1).
        let idx = metal_flat_pool_indices(8, 3, NO_NORMAL_MAP_SLOT, &MaterialUniforms::DEFAULT);
        assert_eq!(idx.albedo, 3);
        assert_eq!(idx.normal, 8);
    }

    #[test]
    fn flat_pool_indices_clamp_out_of_range_and_cap() {
        // Out-of-range albedo / normal slots clamp to the last real texture; a
        // pool larger than the fixed-size MSL array caps every index at the last
        // valid slot so `tex_pool` is never indexed past its end.
        let idx = metal_flat_pool_indices(4, 99, 99, &MaterialUniforms::DEFAULT);
        assert_eq!(idx.albedo, 3); // clamped to last real texture
        assert_eq!(idx.normal, 3);
        let cap = super::super::context::BINDLESS_TEXTURE_COUNT;
        let idx = metal_flat_pool_indices(cap + 5, 9999, 9999, &MaterialUniforms::DEFAULT);
        assert_eq!(idx.albedo, (cap - 1) as u32);
        assert_eq!(idx.normal, (cap - 1) as u32);
    }
}