concinnity-device 0.18.66

// src/metal/cull.rs
//
// 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::transform::IDENTITY;
use objc2::rc::Retained;
use objc2::runtime::ProtocolObject;
use objc2_metal::{
    MTLArgumentEncoder, MTLCommandBuffer as _, MTLComputePassDescriptor, MTLComputePipelineState,
    MTLDevice as _, MTLFunction as _, MTLLibrary as _, MTLRenderCommandEncoder as _,
    MTLRenderPipelineState,
};

use super::context::*;
use super::encode::ComputeEncode;
use super::pipeline::{ns_str, shader_library};
use super::scoped_encoder::ScopedEncoder;
use super::uniforms::*;

// Re-export the camera-distance helper under the legacy local name so the
// existing draw_args builder reads naturally; the actual implementation
// lives on the backend-agnostic `gfx::lod` module.
use crate::gfx::lod::camera_distance as lod_camera_distance;

// All GPU-driven cull state grouped into one feature unit: 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 on the bindless
// path; non-bindless shaders keep the legacy per-draw CPU loop and leave
// every field `None` / default.
pub(crate) struct CullState {
    // GPU-driven cull pipeline. `Some` only when `bindless` is set.
    pub 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 the `icbs` array into the kernel's argument buffer.
    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 cull pipeline 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 encoder + buffer wiring `icbs_2` into the phase-2 kernel.
    pub icb_2_arg_encoder: Option<Retained<ProtocolObject<dyn MTLArgumentEncoder>>>,
    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
    // TAA `prev_view_proj`.
    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 cascaded shadow. All `Some` only on the bindless
    // path with shadows enabled (`bindless && shadow.map_size > 0`); non-bindless
    // / custom-shader worlds leave them `None` and keep the legacy per-cascade
    // CPU shadow loop. The frustum-only `cull_encode_shadow` kernel
    // (`shadow_pipeline`) writes per-cascade indirect commands into one shadow
    // ICB holding `NUM_SHADOW_CASCADES * cull_count()` slots (cascade `c` at base
    // `c * cull_count()`); the depth-only `shadow_bindless_pipeline` then issues
    // each cascade's range. The ICB + its argument buffer grow in lockstep via
    // `ensure_shadow_icb_capacity`.
    pub shadow_pipeline: Option<Retained<ProtocolObject<dyn MTLComputePipelineState>>>,
    pub shadow_bindless_pipeline: Option<Retained<ProtocolObject<dyn MTLRenderPipelineState>>>,
    pub shadow_icb: Option<Retained<ProtocolObject<dyn objc2_metal::MTLIndirectCommandBuffer>>>,
    pub shadow_icb_arg_encoder: Option<Retained<ProtocolObject<dyn MTLArgumentEncoder>>>,
    pub shadow_icb_arg_buffer: Option<Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>>,
    // Command capacity of `shadow_icb` (across all cascades); 0 until first built.
    pub shadow_icb_capacity: usize,
    // 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 is
    // not bindless (the legacy path keeps reusing the main visible set). 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>>,
}

// Kernel buffer index the cull 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;

// The six texture-pool indices the Metal bindless fragment shader
// (`fragment_main_bindless`, main.metal) 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 (a fixed-size MSL array, so an
// over-cap index would read past it). 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 albedo_secondary: u32,
    pub normal_secondary: u32,
    pub emissive: u32,
    pub orm: u32,
}

pub(super) fn metal_flat_pool_indices(
    texture_count: usize,
    texture_slot: usize,
    normal_map_slot: usize,
    material: &crate::gfx::render_types::MaterialUniforms,
) -> FlatPoolIndices {
    use crate::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 secondary / emissive / ORM indices already carry their final shared-pool
    // handle (cook / graphics_system resolve them, `0` = unset for the optional
    // maps, `texture_count` = flat-normal for an unset secondary normal), so they
    // index the pool directly.
    FlatPoolIndices {
        albedo: clamp(albedo_pool_index(texture_slot, tc)),
        normal: clamp(normal_pool_index(normal_map_slot, tc)),
        albedo_secondary: clamp(material.albedo_secondary_index),
        normal_secondary: clamp(material.normal_secondary_index),
        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 six 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: &[crate::gfx::render_types::InstancedCluster],
    texture_count: usize,
) -> Vec<crate::gfx::render_types::GpuObjectData> {
    use crate::gfx::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 (bb_min, bb_max) = crate::gfx::frustum::transform_aabb(
                cluster.local_bb_min,
                cluster.local_bb_max,
                model,
            );
            records.push(GpuObjectData {
                model,
                tint: cluster.material.tint,
                roughness: cluster.material.roughness,
                emissive: cluster.material.emissive,
                metallic: cluster.material.metallic,
                albedo_index: idx.albedo,
                normal_index: idx.normal,
                macro_variation: cluster.material.macro_variation,
                terrain_blend: cluster.material.terrain_blend,
                bb_min,
                cull_distance: cluster.cull_distance,
                bb_max,
                secondary_blend_sharpness: cluster.material.secondary_blend_sharpness,
                albedo_secondary_index: idx.albedo_secondary,
                normal_secondary_index: idx.normal_secondary,
                emissive_map_index: idx.emissive,
                orm_map_index: idx.orm,
                alpha_cutoff: cluster.material.alpha_cutoff,
                _pad: [0.0; 3],
            });
        }
    }
    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 secondary
// / normal-secondary / 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: &crate::gfx::render_types::SkinnedDrawObject,
    texture_count: usize,
) -> crate::gfx::render_types::GpuObjectData {
    let idx = metal_flat_pool_indices(
        texture_count,
        obj.texture_slot,
        obj.normal_map_slot,
        &obj.material,
    );
    let mut rec = crate::gfx::render_types::pack_skinned_record(obj, idx.albedo, idx.normal);
    rec.albedo_secondary_index = idx.albedo_secondary;
    rec.normal_secondary_index = idx.normal_secondary;
    rec.emissive_map_index = idx.emissive;
    rec.orm_map_index = idx.orm;
    rec
}

// 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 crate::gfx::frustum::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>>,
}

// Per-dispatch cull knobs: whether to consult the Hi-Z pyramid (off for the
// mirror cull, whose pyramid is the wrong screen space), an optional pass-timing
// id, and the encoder debug label.
struct CullDispatchOptions<'a> {
    use_hiz: bool,
    timing: Option<super::pass_timing::PassId>,
    label: &'a str,
}

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,
    ) -> Result<Vec<Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>>, String> {
        self.rings
            .joint
            .write_all(&self.device, ring_slot, &self.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, and which
    // is also the only shape Metal accepts for a Slang-declared read-only
    // buffer (slangc lowers one to a mutable `device` pointer, and Metal API
    // validation rejects `setBytes` against that).
    pub(super) fn build_morph_weight_buffers(
        &mut self,
        ring_slot: usize,
    ) -> Result<Vec<Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>>, String> {
        self.rings
            .joint
            .write_weights(&self.device, ring_slot, &self.skinned.morph_weights)
    }

    // Build the previous-pose joint-palette buffers the velocity pre-pass
    // reprojects from, in a separate ring so they never alias the current
    // pose within the same frame.
    pub(super) fn build_prev_joint_buffers(
        &mut self,
        ring_slot: usize,
    ) -> Result<Vec<Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>>, String> {
        self.rings
            .prev_joint
            .write_all(&self.device, ring_slot, &self.skinned.prev_joint_matrices)
    }

    // Build the per-frame `GpuObjectData` buffer for the bindless static
    // pass: one record per `DrawObject`, indexed by the object id the draw
    // call passes as `[[base_instance]]`. Returns `None` when there is no
    // static geometry. 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,
    ) -> Result<Option<Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>>, String> {
        use crate::gfx::render_types::GpuObjectData;
        if self.draw.objects.is_empty() {
            return Ok(None);
        }
        let texture_count = self.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.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(GpuObjectData {
                model: obj.model,
                tint: obj.material.tint,
                roughness: obj.material.roughness,
                emissive: obj.material.emissive,
                metallic: obj.material.metallic,
                albedo_index: idx.albedo,
                normal_index: idx.normal,
                macro_variation: obj.material.macro_variation,
                terrain_blend: obj.material.terrain_blend,
                bb_min: obj.bb_min,
                cull_distance: obj.cull_distance,
                bb_max: obj.bb_max,
                secondary_blend_sharpness: obj.material.secondary_blend_sharpness,
                albedo_secondary_index: idx.albedo_secondary,
                normal_secondary_index: idx.normal_secondary,
                emissive_map_index: idx.emissive,
                orm_map_index: idx.orm,
                alpha_cutoff: obj.material.alpha_cutoff,
                _pad: [0.0; 3],
            });
        }
        // 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. `objects` was `mem::take`n, so this borrows only
        // `instanced.records`, leaving the other fields free.
        if self.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.draw.n_skinned > 0 {
            for obj in &self.skinned.draw_objects {
                objects.push(metal_skinned_record(obj, texture_count));
            }
        }
        let result = self.rings.object.write(
            &self.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 no static geometry. 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,
    ) -> Result<Option<Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>>, String> {
        use crate::gfx::render_types::{GpuDrawArgs, draw_args_flags};
        if self.draw.objects.is_empty() {
            return Ok(None);
        }
        // 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 obj in &self.draw.objects {
            // 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())
                    | crate::gfx::render_types::draw_args_bucket_bits(obj.shader_bucket),
            });
        }
        // 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. Static (base LOD only),
        // so a memcpy; per-instance LOD would move this build per-frame.
        if self.draw.n_instances > 0 {
            args.extend_from_slice(&self.instanced.draw_args);
        }
        // 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.draw.n_skinned > 0 {
            for obj in &self.skinned.draw_objects {
                let d = crate::gfx::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),
                });
            }
        }
        let result = self.rings.draw_args.write(
            &self.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: &crate::gfx::frustum::Frustum,
        cam_pos: [f32; 3],
        counts: crate::metal::context::DrawRecordCounts,
    ) -> Result<(), String> {
        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: Some(super::pass_timing::PassId::Cull),
                label: "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>,
        object_buffer: &ProtocolObject<dyn objc2_metal::MTLBuffer>,
        draw_args_buffer: &ProtocolObject<dyn objc2_metal::MTLBuffer>,
        frustum: &crate::gfx::frustum::Frustum,
        cam_pos: [f32; 3],
        slot: usize,
    ) -> Result<(), String> {
        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 },
            CullOutputTarget {
                icbs: std::slice::from_ref(&mirror.icb),
                arg_buf: &mirror.arg_buffer,
                status,
            },
            CullDispatchOptions {
                use_hiz: false,
                timing: None,
                label: "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,
    ) -> Result<(), String> {
        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 _, MTLComputePipelineState as _, MTLResourceUsage, MTLSize,
        };
        let Some(pipeline) = &self.cull.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-normalised 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,
            object_count: object_count as u32,
            prev_view_proj: self.cull.prev_view_proj,
            hiz_size,
            hiz_mip_count,
            hiz_enabled,
            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,
            _pad_skin: 0,
        };

        let cull_pass_desc = MTLComputePassDescriptor::new();
        if let (Some(t), Some(id)) = (&self.diagnostics.pass_timing, timing) {
            t.attach_compute(&cull_pass_desc, id);
        }
        let enc = ScopedEncoder::new(
            cmd_buf
                .computeCommandEncoderWithDescriptor(&cull_pass_desc)
                .ok_or("failed to get compute encoder")?,
            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.index_buffer, 0, 3);
        enc.set_buffer(arg_buf, 0, CULL_ICB_BUFFER_INDEX);
        // Per-object cull status at buffer(5). The kernel writes it
        // unconditionally; the main cull's status is read by phase 2 under
        // two-pass occlusion, the mirror cull's shared scratch is never read.
        enc.set_buffer(status, 0, 5);
        // Skinned index buffer at buffer(6): the 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 kernel writes draw commands into the ICBs through the argument
        // buffer, so each must be declared resident for the compute pass.
        for icb in icbs {
            enc.useResource_usage(ProtocolObject::from_ref(&**icb), MTLResourceUsage::Write);
        }

        // One thread per draw object, non-uniform grid: no remainder branch
        // needed beyond the kernel's own bounds guard.
        let tg = pipeline.maxTotalThreadsPerThreadgroup().clamp(1, 64);
        enc.dispatchThreads_threadsPerThreadgroup(
            MTLSize {
                width: object_count,
                height: 1,
                depth: 1,
            },
            MTLSize {
                width: tg,
                height: 1,
                depth: 1,
            },
        );
        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 `STATUS_HIZ_CANDIDATE` against the
    // fresh pyramid, projecting through *this* frame's view-projection
    // (`cull_cur_view_proj`), and encodes a draw into `cull_icb_2` for any
    // that turn out visible. `Main2` then issues `cull_icb_2`. 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: &crate::gfx::frustum::Frustum,
        cam_pos: [f32; 3],
    ) -> Result<u32, String> {
        use objc2_metal::{
            MTLComputeCommandEncoder as _, MTLComputePipelineState as _, MTLResourceUsage, MTLSize,
        };
        let (Some(pipeline), Some(arg_buf), Some(status), Some(hiz)) = (
            &self.cull.pipeline_phase2,
            &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 cull_uniforms = CullUniforms {
            planes,
            cam_pos,
            object_count: object_count as u32,
            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,
            skinned_base: self.skinned_record_base() as u32,
            cascade_base: 0,
            bucket_count: self.cull.icbs_2.len() as u32,
            _pad_skin: 0,
        };

        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("failed to get compute encoder")?,
            "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.index_buffer, 0, 3);
        enc.set_buffer(arg_buf, 0, CULL_ICB_BUFFER_INDEX);
        enc.set_buffer(status, 0, 5);
        // 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 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);
        }

        let tg = pipeline.maxTotalThreadsPerThreadgroup().clamp(1, 64);
        enc.dispatchThreads_threadsPerThreadgroup(
            MTLSize {
                width: object_count,
                height: 1,
                depth: 1,
            },
            MTLSize {
                width: tg,
                height: 1,
                depth: 1,
            },
        );
        Ok(0)
    }

    // Encode the GPU-driven cascaded-shadow cull: one
    // `cull_encode_shadow` dispatch per re-rendered cascade (gated by
    // `shadow.render_mask`), each frustum-testing every record against that
    // cascade's LIGHT frustum and encoding survivors into the cascade's slice of
    // the shared shadow ICB (`cascade_base = c * cull_count()`). 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 shadow ICB write lands in a command buffer committed before
    // the `Shadow` render pass's command buffer -- the exact cross-command-buffer
    // FIFO ordering the main cull -> main ICB already relies on. 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>,
    ) -> Result<(), String> {
        use crate::gfx::render_types::NUM_SHADOW_CASCADES;
        use objc2_metal::{
            MTLComputeCommandEncoder as _, MTLComputePipelineState as _, MTLResourceUsage, MTLSize,
        };
        let (Some(pipeline), Some(icb), Some(arg_buf)) = (
            &self.cull.shadow_pipeline,
            &self.cull.shadow_icb,
            &self.cull.shadow_icb_arg_buffer,
        ) else {
            return Ok(());
        };
        let object_count = self.cull_count();
        if object_count == 0 {
            return Ok(());
        }
        // 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 cull_pass_desc = MTLComputePassDescriptor::new();
        let enc = ScopedEncoder::new(
            cmd_buf
                .computeCommandEncoderWithDescriptor(&cull_pass_desc)
                .ok_or("failed to get shadow cull compute encoder")?,
            "shadow cull",
        );
        enc.set_pipeline(pipeline);
        enc.set_buffer(object_buffer, 0, 0);
        enc.set_buffer(draw_args_buffer, 0, 1);
        enc.set_buffer(&self.index_buffer, 0, 3);
        enc.set_buffer(arg_buf, 0, CULL_ICB_BUFFER_INDEX);
        // Skinned index buffer at buffer(6); the kernel bakes it into the
        // skinned-tail commands exactly like the main cull.
        enc.set_buffer(self.skinned_index_or_placeholder(), 0, 6);
        // The kernel writes draw commands into the shadow ICB through the
        // argument buffer, so it must be resident for the compute pass.
        enc.useResource_usage(ProtocolObject::from_ref(&**icb), MTLResourceUsage::Write);

        let tg = pipeline.maxTotalThreadsPerThreadgroup().clamp(1, 64);
        let skinned_base = self.skinned_record_base() as u32;
        for c in 0..NUM_SHADOW_CASCADES {
            if mask & (1u32 << c) == 0 {
                continue;
            }
            // Cascade light frustum: world-space planes from the cascade's light
            // view-projection (the caster-extent near push baked into light_vps
            // survives, so off-screen / tall casters are kept).
            let frustum = crate::gfx::frustum::Frustum::from_view_projection(
                self.shadow.uniforms.light_vps[c],
            );
            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; 3],
                object_count: object_count as u32,
                // 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,
                skinned_base,
                cascade_base: (c * object_count) as u32,
                // The shadow kernel writes one depth-only stream: icbs[0].
                bucket_count: 1,
                _pad_skin: 0,
            };
            enc.set_value(&cull_uniforms, 2);
            enc.dispatchThreads_threadsPerThreadgroup(
                MTLSize {
                    width: object_count,
                    height: 1,
                    depth: 1,
                },
                MTLSize {
                    width: tg,
                    height: 1,
                    depth: 1,
                },
            );
        }
        Ok(())
    }

    // Build the per-frame `BindlessTextures` argument buffer for the bindless
    // static pass: 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 at slot 0) followed by the
    // shadow map and the two IBL cubes. The bindless fragment shader can only
    // reach textures through this argument buffer because discrete texture
    // bindings make it incompatible with indirect command buffers. A fresh
    // buffer is allocated each frame (like the object / draw-args buffers)
    // so a streamed texture swap is picked up and the GPU never reads a buffer
    // the next frame's CPU encode is rewriting. `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,
    ) -> Result<Option<Retained<ProtocolObject<dyn objc2_metal::MTLBuffer>>>, String> {
        use objc2_metal::MTLArgumentEncoder as _;
        // Clone the encoder handle (cheap refcount bump) so no borrow of `self`
        // is held while the ring (a different field) is borrowed mutably below.
        let enc = match &self.bindless_tex_arg_encoder {
            Some(e) => e.clone(),
            None => return Ok(None),
        };
        let len = enc.encodedLength().max(16);
        // Ring slot, grown to the encoder's `encodedLength()` instead of a fresh
        // allocation each frame. The argument encoder rewrites it in place; the
        // fence guarantees the prior user of this slot has retired on the GPU.
        let buf = self.rings.bindless_tex.slot(&self.device, ring_slot, len)?;
        // SAFETY: `buf` was sized to the encoder's `encodedLength()`, and every
        // texture index below is within the `BindlessTextures` layout.
        unsafe {
            enc.setArgumentBuffer_offset(Some(&buf), 0);
        }
        let count = super::context::BINDLESS_TEXTURE_COUNT;
        // The shared pool: every real texture, then the reserved fallbacks --
        // flat-normal at `texture_count`, white at `texture_count + 1` -- and
        // white again across the unused tail, so an over-cap or clamped index
        // still samples a valid texture.
        let texture_count = self.textures.len();
        for i in 0..count {
            let tex = if i < texture_count {
                self.textures[i].as_ref()
            } else if i == texture_count {
                self.fallback_textures[0].as_ref()
            } else {
                self.fallback_textures[1].as_ref()
            };
            // SAFETY: every resource bound here is owned by `self` and outlives the encoder, at the
            // buffer/texture indices the shaders declare.
            unsafe {
                enc.setTexture_atIndex(Some(tex), i);
            }
        }
        // SAFETY: every texture bound here is owned by `self` and outlives the encoder, and the
        // argument ids match the layout the shaders declare: `count` shadow map, then irradiance,
        // prefilter, AO, and `MAX_PROBES` probe cubes.
        unsafe {
            enc.setTexture_atIndex(Some(self.shadow.map.as_ref()), count);
            enc.setTexture_atIndex(Some(self.env_map.irradiance.as_ref()), count + 1);
            enc.setTexture_atIndex(Some(self.env_map.prefilter.as_ref()), count + 2);
            // SSAO occlusion: the blurred AO when SSAO is on, else 1×1 white.
            enc.setTexture_atIndex(Some(self.ao_output_texture()), count + 3);
            // Local reflection probe cube array (specular only): one slice per
            // baked probe, the sky prefilter for unused slots. Occupies argument
            // ids `count + 4 ..= count + 4 + MAX_PROBES`.
            for i in 0..concinnity_render::uniforms::MAX_PROBES {
                enc.setTexture_atIndex(Some(self.probe_cube_or_sky(i)), count + 4 + i);
            }
            // Spot shadow map array (1x1 fallback when nothing casts), just past
            // the probe cubes.
            enc.setTexture_atIndex(
                Some(self.spot_shadow.map.as_ref()),
                count + 4 + concinnity_render::uniforms::MAX_PROBES,
            );
            // The two area-light LTC tables follow the spot shadow array.
            enc.setTexture_atIndex(
                Some(self.ltc_matrix_texture.as_ref()),
                count + 5 + concinnity_render::uniforms::MAX_PROBES,
            );
            enc.setTexture_atIndex(
                Some(self.ltc_magnitude_texture.as_ref()),
                count + 6 + concinnity_render::uniforms::MAX_PROBES,
            );
        }
        Ok(Some(buf))
    }
    // 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.
    pub(super) fn use_bindless_textures(
        &self,
        encoder: &ProtocolObject<dyn objc2_metal::MTLRenderCommandEncoder>,
    ) {
        use objc2_metal::{MTLRenderStages, MTLResourceUsage};
        for tex in self.textures.iter().chain(self.fallback_textures.iter()) {
            encoder.useResource_usage_stages(
                ProtocolObject::from_ref(&**tex),
                MTLResourceUsage::Read,
                MTLRenderStages::Fragment,
            );
        }
        for tex in [
            self.shadow.map.as_ref(),
            self.spot_shadow.map.as_ref(),
            self.ltc_matrix_texture.as_ref(),
            self.ltc_magnitude_texture.as_ref(),
            self.env_map.irradiance.as_ref(),
            self.env_map.prefilter.as_ref(),
        ] {
            encoder.useResource_usage_stages(
                ProtocolObject::from_ref(tex),
                MTLResourceUsage::Read,
                MTLRenderStages::Fragment,
            );
        }
        // SSAO occlusion travels in the BindlessTextures argument buffer too.
        encoder.useResource_usage_stages(
            ProtocolObject::from_ref(self.ao_output_texture()),
            MTLResourceUsage::Read,
            MTLRenderStages::Fragment,
        );
        // The reflection probe cube array (each slice, or its sky fallback) rides
        // the argument buffer, so every bound slice must be resident.
        for i in 0..concinnity_render::uniforms::MAX_PROBES {
            encoder.useResource_usage_stages(
                ProtocolObject::from_ref(self.probe_cube_or_sky(i)),
                MTLResourceUsage::Read,
                MTLRenderStages::Fragment,
            );
        }
    }
}

// The GPU-driven cull stage: a compute pipeline plus the argument encoder
// that wires an `MTLIndirectCommandBuffer` into the kernel. The kernel
// reaches the ICB only through an argument buffer, so the encoder must be
// kept to (re)encode that argument buffer whenever the ICB is recreated.
//
// The phase-2 pipeline + its argument encoder drive two-pass occlusion: the
// `cull_encode_phase2` kernel re-tests phase-1's Hi-Z-occluded objects against
// the rebuilt pyramid and encodes survivors into a second ICB. Built from the
// same library whenever the bindless path is active; used only when
// `occlusion_two_pass` is on.
pub(super) struct CullPipeline {
    pub state: Retained<ProtocolObject<dyn MTLComputePipelineState>>,
    pub icb_arg_encoder: Retained<ProtocolObject<dyn MTLArgumentEncoder>>,
    pub state_phase2: Retained<ProtocolObject<dyn MTLComputePipelineState>>,
    pub icb2_arg_encoder: Retained<ProtocolObject<dyn MTLArgumentEncoder>>,
}

// Build the GPU-driven cull pipeline. The `cull_encode` kernel runs
// one thread per `DrawObject`: it frustum/distance-tests the object against
// `CullUniforms` and, for survivors, encodes an indexed draw into the
// indirect command buffer; culled or disabled objects have their command
// reset to a no-op. The render pass then issues the whole buffer with one
// `executeCommandsInBuffer`, so the CPU never walks the draw list.
//
// The frustum and distance maths mirror `gfx::frustum` exactly (the six
// planes are extracted CPU-side and handed in already normalised), so the
// GPU path culls identically to the CPU BVH path it replaces.
pub(super) fn build_cull_pipeline(
    device: &ProtocolObject<dyn objc2_metal::MTLDevice>,
    hot_reload: bool,
) -> Result<CullPipeline, String> {
    let library = shader_library(device, hot_reload, "cull.metal")?;

    let cull_fn = library
        .newFunctionWithName(&ns_str("cull_encode"))
        .ok_or("cull_encode not found in cull library")?;

    let state = device
        .newComputePipelineStateWithFunction_error(&cull_fn)
        .map_err(|e| format!("failed to create cull pipeline state: {:?}", e))?;

    // SAFETY: CULL_ICB_BUFFER_INDEX is the static buffer index the kernel
    // declares its argument-buffer parameter at.
    let icb_arg_encoder =
        unsafe { cull_fn.newArgumentEncoderWithBufferIndex(CULL_ICB_BUFFER_INDEX) };

    // Second-pass cull (two-pass occlusion): same library, the
    // `cull_encode_phase2` kernel. It declares its ICB argument buffer at the
    // same buffer index, so the encoder is built the same way but tied to the
    // second-pass function.
    let cull_fn_phase2 = library
        .newFunctionWithName(&ns_str("cull_encode_phase2"))
        .ok_or("cull_encode_phase2 not found in cull library")?;
    let state_phase2 = device
        .newComputePipelineStateWithFunction_error(&cull_fn_phase2)
        .map_err(|e| format!("failed to create phase-2 cull pipeline state: {:?}", e))?;
    // SAFETY: same static buffer index: `cull_encode_phase2` declares its
    // ICBContainer argument buffer at CULL_ICB_BUFFER_INDEX too.
    let icb2_arg_encoder =
        unsafe { cull_fn_phase2.newArgumentEncoderWithBufferIndex(CULL_ICB_BUFFER_INDEX) };

    Ok(CullPipeline {
        state,
        icb_arg_encoder,
        state_phase2,
        icb2_arg_encoder,
    })
}

// The GPU-driven shadow cull pipeline + the argument encoder that wires its
// shadow ICB into the `cull_encode_shadow` kernel.
pub(super) type ShadowCullPipeline = (
    Retained<ProtocolObject<dyn MTLComputePipelineState>>,
    Retained<ProtocolObject<dyn MTLArgumentEncoder>>,
);

// Build the GPU-driven cascaded-shadow cull pipeline: the
// `cull_encode_shadow` kernel + the argument encoder that wires its shadow ICB
// into the kernel. Compiled from the same `cull.metal` source as the main cull
// (a separate compile keeps the call site shadow-gated rather than always
// building it on the bindless path). 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,
) -> Result<ShadowCullPipeline, String> {
    let library = shader_library(device, hot_reload, "cull.metal")?;
    let shadow_fn = library
        .newFunctionWithName(&ns_str("cull_encode_shadow"))
        .ok_or("cull_encode_shadow not found in cull library")?;
    let state = device
        .newComputePipelineStateWithFunction_error(&shadow_fn)
        .map_err(|e| format!("failed to create shadow cull pipeline state: {:?}", e))?;
    // SAFETY: same static buffer index as the main cull kernels:
    // `cull_encode_shadow` declares its ICBContainer argument buffer at
    // CULL_ICB_BUFFER_INDEX.
    let icb_arg_encoder =
        unsafe { shadow_fn.newArgumentEncoderWithBufferIndex(CULL_ICB_BUFFER_INDEX) };
    Ok((state, icb_arg_encoder))
}

#[cfg(test)]
mod tests {
    use super::metal_flat_pool_indices;
    use crate::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 {
            albedo_secondary_index: 1,
            normal_secondary_index: 2,
            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.albedo_secondary, 1);
        assert_eq!(idx.normal_secondary, 2);
        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);
    }
}