concinnity-device 0.19.119

//! Main pass (off-screen HDR + 4x MSAA). Renders the visible scene into the
//! HDR color + depth attachments, multisample-resolving to `hdr_resolve`.
//!
//! One geometry path: the GPU-driven bindless pass, issued through the cull's
//! indirect command buffers. Static objects, folded instances and the folded
//! skinned tail are all cull records, so the encode is one prefix range per
//! shader bucket plus the skinned tail. `object_buffer` / `bindless_tex_args`
//! are `Some` exactly when the world has something to draw.
//!
//! The pass is encoded on a single `MTLRenderCommandEncoder`. Two
//! `MTLParallelRenderCommandEncoder` attempts (one full-fat with parallel
//! encoders on every shadow cascade, one scoped to just this main pass with 3
//! sub-encoders + no pass-timing) both tripped G14X (M2/M3 Pro/Max class) into
//! an abort inside `IOGPUMetalCommandBufferStorageAllocResourceAtIndex`: the
//! crash window scaled with parallel-encoder usage rate (~20 s with shadow
//! split, ~90 s with main-only) but never went away. The mechanism appears
//! fundamentally incompatible with our usage on this hardware / macOS 26.4
//! combo.
#![deny(unsafe_op_in_unsafe_fn)]

use concinnity_core::gfx::render_types;
use concinnity_core::render::error::{RenderError, RenderResult};
use concinnity_core::render::planar_reflection::PixelRect;
use concinnity_core::render::post::reflection_path::ReflectionPath;
use concinnity_core::render::uniforms::ViewUniforms;
use objc2::rc::Retained;
use objc2::runtime::ProtocolObject;
use objc2_foundation::ns_string;
use objc2_metal::{
    MTLBuffer, MTLClearColor, MTLCommandBuffer as _, MTLCommandEncoder as _, MTLLoadAction,
    MTLRenderCommandEncoder as _, MTLRenderPassDescriptor, MTLScissorRect, MTLStoreAction,
};

use crate::metal::context::{BINDLESS_TEXTURE_ARG_BUFFER_INDEX, MtlContext};
use crate::metal::depth::CLEAR_DEPTH;
use crate::metal::encode::RenderEncode;
use crate::metal::pass_timing::PassTimer;
use crate::metal::scoped_encoder::ScopedEncoder;

// Camera state a main-pass encode builds its ViewUniforms from. `view` is
// `self.view_matrix` for the on-screen main pass (and its phase-2 sibling) but
// the caller's face view matrix for a reflection-probe face capture, so it
// travels with the rest of the camera state rather than being read off `self`.
#[derive(Clone, Copy)]
pub(in crate::metal) struct MainPassCamera {
    pub elapsed: f32,
    pub vp: [[f32; 4]; 4],
    pub view: [[f32; 4]; 4],
    pub cam_pos: [f32; 3],
}

// The GPU-driven buffers the main pass consumes. `object_buffer` and
// `bindless_tex_args` are `Some` together exactly when the cull path ran this
// frame; a world with nothing to draw leaves them `None` and the pass stops at
// its clear. `counts` is the record set those buffers were built for, which the
// indirect-draw ranges address; it is the live draw list for a frame pass and an
// earlier snapshot for the reflection-probe bake.
#[derive(Clone, Copy)]
pub(in crate::metal) struct GpuFrameBuffers<'a> {
    pub object_buffer: Option<&'a Retained<ProtocolObject<dyn MTLBuffer>>>,
    pub material_params: Option<&'a Retained<ProtocolObject<dyn MTLBuffer>>>,
    pub bindless_tex_args: Option<&'a Retained<ProtocolObject<dyn MTLBuffer>>>,
    pub deformed_skinned: Option<&'a Retained<ProtocolObject<dyn MTLBuffer>>>,
    pub counts: crate::metal::context::DrawRecordCounts,
}

// The reflection-probe face attachments `encode_main_into_face` renders into
// instead of the HDR targets: a square MSAA color + depth, resolving color
// into `resolve`. A face reuses the main pipelines, so it carries whatever
// sample count they were built at: at one sample `color_msaa` is `None` and the
// pass draws straight into `resolve`.
#[derive(Clone, Copy)]
pub(in crate::metal) struct FaceTargets<'a> {
    pub color_msaa: Option<&'a ProtocolObject<dyn objc2_metal::MTLTexture>>,
    pub depth: &'a ProtocolObject<dyn objc2_metal::MTLTexture>,
    pub resolve: &'a ProtocolObject<dyn objc2_metal::MTLTexture>,
    // Array slice of `resolve` this face resolves into. The probe capture passes
    // the cube face it is rendering; the planar mirror's target is a plain 2D
    // texture, so it passes 0.
    pub resolve_slice: usize,
}

// A light-cluster grid: the params a fragment places itself in the grid with
// and the per-cluster light lists and probe masks binned for them.
#[derive(Clone, Copy)]
pub(in crate::metal) struct ClusterGrid<'a> {
    pub params: &'a render_types::ClusterParams,
    pub lists: &'a ProtocolObject<dyn MTLBuffer>,
}

// How one off-camera render into a `FaceTargets` is drawn: the bindless ICB to
// execute instead of the main cull's (the planar mirror passes its slot's mirror
// ICB, culled against the reflected frustum; the probe capture reuses the main
// cull ICB), the texel rectangle the render is clipped to (`None` draws the
// whole target), the cluster grid binned for its viewpoint (`None` shades every
// local light and probe), and where its encoder sits in a pass's timing span.
#[derive(Clone, Copy)]
pub(in crate::metal) struct FacePass<'a> {
    pub icb_override: Option<&'a ProtocolObject<dyn objc2_metal::MTLIndirectCommandBuffer>>,
    pub scissor: Option<PixelRect>,
    pub clusters: Option<ClusterGrid<'a>>,
    pub timer: PassTimer,
}

impl FacePass<'_> {
    // The whole target through the main cull ICB, unclustered and untimed: the
    // probe capture.
    pub(in crate::metal) const PROBE: FacePass<'static> = FacePass {
        icb_override: None,
        scissor: None,
        clusters: None,
        timer: PassTimer::None,
    };
}

impl MtlContext {
    // 1.0 when a reflection-resolve pass (SSR resolve or RT reflections) will
    // composite over this frame's HDR target, else 0.0. Mirrors the
    // `scene_input` gate in draw/pass_uniforms.rs (both resolves feed the
    // reflection composite's output and the graph picks RT over SSR). The forward shader reads it from
    // `ViewUniforms.reflections_enabled` to hand glossy specular to that resolve.
    fn reflection_resolve_active(&self) -> f32 {
        if self.reflection_path().resolves() {
            1.0
        } else {
            0.0
        }
    }

    // Which reflection stages run this frame, from the authored SSR and the
    // live RT pass and BVH. This backend reads the composite's output only while
    // a resolve writes it, so the RT node never runs without a BVH.
    pub(in crate::metal) fn reflection_path(&self) -> ReflectionPath {
        ReflectionPath::new(
            self.ssr.settings.is_some(),
            self.rt.settings.is_some(),
            self.rt.accel.is_some(),
        )
    }

    // The frame's unlit flag for ViewUniforms, from the viewport view mode.
    fn shade_mode(&self) -> f32 {
        if self.state.view.mode == concinnity_core::gfx::view_modes::ViewMode::Unlit {
            1.0
        } else {
            0.0
        }
    }

    // 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_main_pass(
        &self,
        cmd_buf: &ProtocolObject<dyn objc2_metal::MTLCommandBuffer>,
        camera: MainPassCamera,
        gpu: GpuFrameBuffers,
        world_hidden: bool,
    ) -> RenderResult<u32> {
        // Only `object_buffer` gates the descriptor's store action below; the
        // rest of `gpu` travels intact into `encode_main_static_into`.
        let object_buffer = gpu.object_buffer;
        // Build the HDR render pass descriptor. Color writes into the MSAA
        // attachment and resolves into the single-sample target at end-of-pass;
        // depth lives entirely on the MSAA attachment and is discarded unless
        // a later pass (projected decals, volumetric fog) needs to sample it.
        // Under two-pass occlusion the phase-2 main pass (`Main2`) loads this
        // pass's MSAA color to draw the disoccluded geometry on top, so the
        // MSAA samples must be stored, not just resolved away. `Main2`
        // performs the final resolve. Without two-pass we resolve-and-discard
        // the MSAA samples as before. The decision mirrors the graph's two-pass
        // gating (`two_pass_occlusion` config AND the bindless cull path active
        // this frame, i.e. an `object_buffer` exists).
        let store_msaa_color = self.cull.two_pass_occlusion
            && object_buffer.is_some()
            && self.cull.pipeline_phase2.is_some();
        let main_pass_desc = MTLRenderPassDescriptor::new();
        let [r, g, b, a] = self.state.view.clear_color;
        // SAFETY: plain descriptor property setters; the subscripted slots are ones this descriptor
        // declares.
        unsafe {
            let ca = main_pass_desc
                .colorAttachments()
                .objectAtIndexedSubscript(0);
            ca.setTexture(Some(self.targets.hdr.color_attachment()));
            ca.setLoadAction(MTLLoadAction::Clear);
            if self.targets.hdr.multisampled() {
                ca.setResolveTexture(Some(self.targets.hdr.hdr_resolve.as_ref()));
                ca.setStoreAction(if store_msaa_color {
                    MTLStoreAction::StoreAndMultisampleResolve
                } else {
                    MTLStoreAction::MultisampleResolve
                });
            } else {
                // The color attachment *is* the spine, so there is nothing to
                // resolve and the pass just stores what it drew.
                ca.setStoreAction(MTLStoreAction::Store);
            }
            ca.setClearColor(MTLClearColor {
                red: r as f64,
                green: g as f64,
                blue: b as f64,
                alpha: a as f64,
            });

            let da = main_pass_desc.depthAttachment();
            da.setTexture(Some(self.targets.hdr.depth_attachment()));
            da.setLoadAction(MTLLoadAction::Clear);
            da.setClearDepth(CLEAR_DEPTH);
            // `depth_resolve` is the canonical post-rasterize scene depth the
            // post chain consumes: raymarch writes hit depth into it, and
            // water / decal / fog sample it. Multisampled, the Main pass
            // produces it by resolving with the `Sample0` filter (single-sample
            // is enough -- those passes only ever read sample 0) and keeps the
            // MSAA attachment alive for the Hi-Z build; without MSAA it is the
            // attachment, and the store is the whole of it.
            if self.targets.hdr.multisampled() {
                da.setResolveTexture(Some(self.targets.hdr.depth_resolve.as_ref()));
                da.setDepthResolveFilter(objc2_metal::MTLMultisampleDepthResolveFilter::Sample0);
                da.setStoreAction(MTLStoreAction::StoreAndMultisampleResolve);
            } else {
                da.setStoreAction(MTLStoreAction::Store);
            }
        }

        if let Some(t) = &self.diagnostics.pass_timing {
            t.attach_render(&main_pass_desc, super::super::pass_timing::PassId::Main);
        }
        // ScopedEncoder ends the main HDR pass (MSAA resolves into hdr_resolve)
        // and pops the debug group when it drops at end of scope.
        let encoder = ScopedEncoder::new(
            cmd_buf
                .renderCommandEncoderWithDescriptor(&main_pass_desc)
                .ok_or_else(|| RenderError::Other("failed to get render encoder".to_string()))?,
            ns_string!("main pass"),
        );
        // Wireframe view: fill mode is encoder state that indirect commands
        // inherit, so the one call covers the ICB sub-paths too.
        if self.state.view.mode == concinnity_core::gfx::view_modes::ViewMode::Wireframe {
            encoder.setTriangleFillMode(objc2_metal::MTLTriangleFillMode::Lines);
        }

        let view_uniforms = self.main_view_uniforms(&camera);

        // While the world is hidden behind an opaque menu, the pass stops at the
        // descriptor's Clear load action. A scene-less world (no main pipeline)
        // takes the same bare-clear shape every frame.
        if world_hidden || self.cull.main_pipeline.is_none() {
            return Ok(0);
        }

        // Main camera: bind the per-cluster light lists once for the pass.
        self.bind_clusters(&encoder, Some(self.main_cluster_grid()));
        let draw_calls = self.encode_main_geometry_into(
            &encoder,
            &view_uniforms,
            gpu,
            // Main pass: the main cull ICB (no override).
            None,
        );
        // The sky lands behind phase 1's geometry. Under two-pass occlusion
        // `Main2` then draws the disoccluded rest over it, which wins the depth
        // test because the sky writes none.
        if self.draws_sky(self.state.view.mode) {
            self.encode_sky(&encoder, &view_uniforms);
        }
        Ok(draw_calls)
    }

    // Render the main pass into one reflection-probe cube face instead of the
    // HDR targets. A thin sibling of `encode_main_pass`: same three geometry
    // sub-paths and shared bindings, but the render-pass descriptor points at a
    // square MSAA color + depth (resolving color into `face_resolve`), and the
    // view + view-projection are the caller's face matrices (not `self.*`), so
    // the capture never disturbs the frame's camera state. Depth is cleared and
    // discarded -- the probe consumes only the resolved color. Driven by
    // `capture_reflection_probe` (metal/probe.rs), once at first frame.
    pub(in crate::metal) fn encode_main_into_face(
        &self,
        cmd_buf: &ProtocolObject<dyn objc2_metal::MTLCommandBuffer>,
        face_targets: FaceTargets,
        camera: MainPassCamera,
        gpu: GpuFrameBuffers,
        pass: FacePass<'_>,
    ) -> RenderResult<u32> {
        let FaceTargets {
            color_msaa: face_color_msaa,
            depth: face_depth,
            resolve: face_resolve,
            resolve_slice,
        } = face_targets;
        let MainPassCamera {
            elapsed,
            vp,
            view,
            cam_pos,
        } = camera;
        let desc = MTLRenderPassDescriptor::new();
        let [r, g, b, a] = self.state.view.clear_color;
        // SAFETY: plain descriptor property setters; the subscripted slots are ones this descriptor
        // declares.
        unsafe {
            let ca = desc.colorAttachments().objectAtIndexedSubscript(0);
            match face_color_msaa {
                Some(msaa) => {
                    ca.setTexture(Some(msaa));
                    ca.setResolveTexture(Some(face_resolve));
                    ca.setResolveSlice(resolve_slice);
                    ca.setStoreAction(MTLStoreAction::MultisampleResolve);
                }
                // Single-sampled: the face draws into the destination slice
                // itself, so the slice moves from the resolve to the attachment.
                None => {
                    ca.setTexture(Some(face_resolve));
                    ca.setSlice(resolve_slice);
                    ca.setStoreAction(MTLStoreAction::Store);
                }
            }
            ca.setLoadAction(MTLLoadAction::Clear);
            ca.setClearColor(MTLClearColor {
                red: r as f64,
                green: g as f64,
                blue: b as f64,
                alpha: a as f64,
            });

            let da = desc.depthAttachment();
            da.setTexture(Some(face_depth));
            da.setLoadAction(MTLLoadAction::Clear);
            da.setClearDepth(CLEAR_DEPTH);
            da.setStoreAction(MTLStoreAction::DontCare);
        }
        if let Some(t) = &self.diagnostics.pass_timing {
            t.attach_render_timer(&desc, pass.timer);
        }

        let encoder = ScopedEncoder::new(
            cmd_buf
                .renderCommandEncoderWithDescriptor(&desc)
                .ok_or_else(|| {
                    RenderError::Other("failed to get probe render encoder".to_string())
                })?,
            ns_string!("probe face"),
        );
        if let Some(r) = pass.scissor {
            encoder.setScissorRect(MTLScissorRect {
                x: r.x as usize,
                y: r.y as usize,
                width: r.width as usize,
                height: r.height as usize,
            });
        }

        let view_uniforms = ViewUniforms {
            vp,
            view,
            elapsed,
            // Probe-face bake: no reflection resolve runs over the probe cube, so
            // the forward probe specular is the only reflection source here. Keep
            // it (0.0) so captured glossy surfaces are not flattened.
            reflections_enabled: 0.0,
            cam_pos,
            prefilter_mip_count: self.scene.env_map.prefilter_mip_count as f32,
            // A probe capture is always lit, whatever the viewport shows.
            shade_mode: 0.0,
            ambient_occlusion: 0.0,
            sky_rot: self.state.view.sky_rot,
        };

        self.bind_clusters(&encoder, pass.clusters);
        let draw_calls =
            self.encode_main_geometry_into(&encoder, &view_uniforms, gpu, pass.icb_override);
        // A face is always rendered lit, whatever the viewport shows.
        if self.draws_sky(concinnity_core::gfx::view_modes::ViewMode::Lit) {
            self.encode_sky(&encoder, &view_uniforms);
        }
        Ok(draw_calls)
    }

    // Phase-2 main pass for two-pass occlusion (`Main2`). Loads (does not
    // clear) the HDR color + depth that `encode_main_pass` (phase 1) wrote
    // and re-runs the bindless indirect draw through `cull_icb_2` (the phase-2
    // cull's output), depth-compositing the disoccluded geometry with phase 1.
    // Folded instances AND folded skinned objects ride the unified cull buffers
    // as cullable records, so both are Hi-Z-culled through both phases: drawn in
    // phase 1 when visible, and redrawn here only when phase 1 Hi-Z-occluded
    // them and the rebuilt pyramid (Cull2) disoccludes them. The shared
    // `execute_bindless_static_icb` issues the static+instance range then the
    // skinned tail of `cull_icb_2` (the phase-2 cull resets every already-drawn
    // slot, so nothing double-draws). Resolves color + depth at end-of-pass so
    // the post-decoration stack reads the combined result. A no-op (returns 0)
    // when there is nothing to redraw: two-pass off, no bindless geometry, or
    // the phase-2 ICB was not built.
    pub(in crate::metal) fn encode_main_pass_phase2(
        &self,
        cmd_buf: &ProtocolObject<dyn objc2_metal::MTLCommandBuffer>,
        camera: MainPassCamera,
        gpu: GpuFrameBuffers,
    ) -> RenderResult<u32> {
        let GpuFrameBuffers {
            object_buffer,
            bindless_tex_args,
            ..
        } = gpu;
        let (Some(obj_buf), Some(tex_args), false) = (
            object_buffer,
            bindless_tex_args,
            self.cull.icbs_2.is_empty(),
        ) else {
            return Ok(0);
        };

        // Load the phase-1 MSAA color + depth (phase 1 stored them under
        // two-pass), draw the disoccluded geometry on top, and resolve both at
        // end-of-pass: this resolve is the one the post stack consumes.
        let main_pass_desc = MTLRenderPassDescriptor::new();
        // SAFETY: plain descriptor property setters; the subscripted slots are ones this descriptor
        // declares.
        unsafe {
            let ca = main_pass_desc
                .colorAttachments()
                .objectAtIndexedSubscript(0);
            ca.setTexture(Some(self.targets.hdr.color_attachment()));
            ca.setLoadAction(MTLLoadAction::Load);

            let da = main_pass_desc.depthAttachment();
            da.setTexture(Some(self.targets.hdr.depth_attachment()));
            da.setLoadAction(MTLLoadAction::Load);

            if self.targets.hdr.multisampled() {
                ca.setResolveTexture(Some(self.targets.hdr.hdr_resolve.as_ref()));
                ca.setStoreAction(MTLStoreAction::StoreAndMultisampleResolve);
                da.setResolveTexture(Some(self.targets.hdr.depth_resolve.as_ref()));
                da.setDepthResolveFilter(objc2_metal::MTLMultisampleDepthResolveFilter::Sample0);
                da.setStoreAction(MTLStoreAction::StoreAndMultisampleResolve);
            } else {
                ca.setStoreAction(MTLStoreAction::Store);
                da.setStoreAction(MTLStoreAction::Store);
            }
        }

        if let Some(t) = &self.diagnostics.pass_timing {
            t.attach_render(&main_pass_desc, super::super::pass_timing::PassId::Main2);
        }
        let encoder = ScopedEncoder::new(
            cmd_buf
                .renderCommandEncoderWithDescriptor(&main_pass_desc)
                .ok_or_else(|| RenderError::Other("failed to get render encoder".to_string()))?,
            ns_string!("main2 pass"),
        );
        if self.state.view.mode == concinnity_core::gfx::view_modes::ViewMode::Wireframe {
            encoder.setTriangleFillMode(objc2_metal::MTLTriangleFillMode::Lines);
        }

        let view_uniforms = self.main_view_uniforms(&camera);
        self.bind_main_pass_shared(&encoder, &view_uniforms);
        // Main2 is the same main camera as phase 1, so it reads the clusters too.
        self.bind_clusters(&encoder, Some(self.main_cluster_grid()));
        let draw_calls =
            self.execute_bindless_static_icb(&encoder, obj_buf, tex_args, &self.cull.icbs_2, gpu);

        Ok(draw_calls)
    }

    // Issue the GPU-driven bindless pass: bind the per-object data +
    // bindless-texture argument buffer, declare the index buffer + sampled
    // textures resident (they are reached only through indirect commands /
    // the argument buffer, never bound on the encoder), then execute the `icb`.
    // Shared by the phase-1 main pass (`icb = cull_icb`) and the phase-2 main
    // pass (`icb = cull_icb_2`) so both issue the bindless geometry identically.
    //
    // The ICB is split into two `executeCommandsInBuffer` ranges because Metal
    // bakes the index buffer into each indirect command: the prefix records are
    // static + instances (static u32 index buffer, static vertex buffer already
    // bound at binding 1 by `bind_main_pass_shared`); the tail records are the
    // folded skinned ones, which the cull kernel encoded against the skinned
    // u16 index buffer and which read the compute-deformed vertices, so this
    // rebinds the deformed buffer at binding 1 for that range (inherited by the
    // ICB commands). Both ranges come from `counts`, the record set the caller's
    // buffers were built and culled for, so a snapshot capture never addresses
    // slots the live draw list grew after its ICB was sized. Returns the number
    // of indirect draws issued (1 or 2).
    fn execute_bindless_static_icb(
        &self,
        enc: &ProtocolObject<dyn objc2_metal::MTLRenderCommandEncoder>,
        obj_buf: &Retained<ProtocolObject<dyn MTLBuffer>>,
        tex_args: &Retained<ProtocolObject<dyn MTLBuffer>>,
        icbs: &[Retained<ProtocolObject<dyn objc2_metal::MTLIndirectCommandBuffer>>],
        gpu: GpuFrameBuffers,
    ) -> u32 {
        use objc2_metal::{MTLRenderStages, MTLResourceUsage};
        let GpuFrameBuffers {
            material_params,
            deformed_skinned,
            counts,
            ..
        } = gpu;
        // Object records (binding 9) + bindless textures (binding 7) are shared
        // by both ranges: the skinned records live in the same object buffer and
        // sample the same flat pool. The object id reaches the shader via each
        // command's [[base_instance]].
        enc.set_vertex_buffer(obj_buf, 0, 9);
        enc.set_fragment_buffer(obj_buf, 0, 9);
        // The material parameter table either hook may read.
        if let Some(params) = material_params {
            let slot = crate::metal::material_params::MATERIAL_PARAMS_BUFFER_INDEX;
            enc.set_vertex_buffer(params, 0, slot);
            enc.set_fragment_buffer(params, 0, slot);
        }
        enc.set_fragment_buffer(tex_args, 0, BINDLESS_TEXTURE_ARG_BUFFER_INDEX);
        // The engine sampler block (single-source main program only;
        // world-authored fragments use inline samplers and ignore the
        // slot).
        if let Some(sampler_args) = &self.arg_buffers.bindless_sampler_args {
            enc.set_fragment_buffer(
                sampler_args,
                0,
                crate::metal::context::BINDLESS_SAMPLER_ARG_BUFFER_INDEX,
            );
        }
        self.use_bindless_textures(enc);

        let mut draw_calls = 0u32;

        // Static + instance prefix, once per shader bucket: bucket 0 executes
        // under the pipeline `bind_main_pass_shared` bound, each later bucket
        // under its material shader's pipeline. The cull kernel wrote every
        // record's command into exactly one bucket's ICB, so the ranges never
        // double-draw. The static u32 index buffer is referenced only inside the
        // indirect commands (not bound), so make it resident; the static vertex
        // buffer is already bound at binding 1.
        if let Some(prefix) = counts.prefix(0) {
            enc.useResource_usage_stages(
                ProtocolObject::from_ref(&*self.scene.index_buffer),
                MTLResourceUsage::Read,
                MTLRenderStages::Vertex,
            );
            let range = crate::metal::context::ns_range(prefix);
            for (b, icb) in icbs.iter().enumerate() {
                if b > 0 {
                    // A bucket whose Shader is not resident yet (its scene has
                    // not pinned) has no pipeline: skip it until warmup builds
                    // one rather than drawing it with the wrong program.
                    let Some(pipelines) = self.cull.world_pipelines.get(b) else {
                        continue;
                    };
                    enc.set_pipeline(&pipelines.main);
                }
                // SAFETY: the prefix spans the static + instance command slots
                // (`ensure_icb_capacity` sized every ICB for `counts.total`).
                unsafe {
                    enc.executeCommandsInBuffer_withRange(icb, range);
                }
                draw_calls += 1;
            }
            // Restore the default pipeline for the skinned tail below (and for
            // the caller's subsequent sub-paths, which re-bind anyway).
            if icbs.len() > 1
                && let Some(pipelines) = &self.cull.main_pipeline
            {
                enc.set_pipeline(&pipelines.main);
            }
        }

        // Skinned tail: bind the deformed vertex buffer (inherited by the ICB
        // commands) and make the skinned index buffer resident (the cull
        // kernel baked it into these commands). `deformed_skinned` is `Some`
        // exactly when the fold is active, which is also when the tail is
        // non-empty. Skinned draws always render under the world default shader,
        // so only bucket 0's ICB executes here.
        if let (Some(deformed), Some(icb0), Some(tail)) =
            (deformed_skinned, icbs.first(), counts.skinned_tail(0))
        {
            enc.set_vertex_buffer(deformed, 0, 1);
            if let Some(skinned_ib) = self.skinned.index_buffer.as_ref() {
                enc.useResource_usage_stages(
                    ProtocolObject::from_ref(&**skinned_ib),
                    MTLResourceUsage::Read,
                    MTLRenderStages::Vertex,
                );
            }
            // SAFETY: the tail spans the folded skinned command slots.
            unsafe {
                enc.executeCommandsInBuffer_withRange(icb0, crate::metal::context::ns_range(tail));
            }
            draw_calls += 1;
        }
        draw_calls
    }

    // The view block the on-screen main pass shades through, which the G-buffer
    // pre-pass binds too so a vertex hook places a surface identically in both.
    pub(in crate::metal) fn main_view_uniforms(&self, camera: &MainPassCamera) -> ViewUniforms {
        ViewUniforms {
            vp: camera.vp,
            view: self.state.view.matrix,
            elapsed: camera.elapsed,
            reflections_enabled: self.reflection_resolve_active(),
            cam_pos: camera.cam_pos,
            prefilter_mip_count: self.scene.env_map.prefilter_mip_count as f32,
            shade_mode: self.shade_mode(),
            ambient_occlusion: 1.0,
            sky_rot: self.state.view.sky_rot,
        }
    }

    // The main camera's cluster grid, binned by the frame's `LightCull` node.
    pub(in crate::metal) fn main_cluster_grid(&self) -> ClusterGrid<'_> {
        ClusterGrid {
            params: &self.cluster_params,
            lists: &self.light_cull.cluster_buffer,
        }
    }

    // Bind the clustered-lighting inputs the forward pass reads: the params at
    // fragment buffer(11) + the per-cluster light lists and probe masks at
    // buffer(12). Bound once per pass on the shared encoder. Without a grid
    // (a probe capture) `use_clusters` is cleared, so every local light and
    // probe is iterated.
    fn bind_clusters(
        &self,
        enc: &ProtocolObject<dyn objc2_metal::MTLRenderCommandEncoder>,
        grid: Option<ClusterGrid<'_>>,
    ) {
        let (params, lists) = match grid {
            Some(grid) => (*grid.params, grid.lists),
            None => (
                render_types::ClusterParams {
                    use_clusters: 0,
                    ..self.cluster_params
                },
                &*self.light_cull.cluster_buffer,
            ),
        };
        enc.set_fragment_value(&params, 11);
        enc.set_fragment_buffer(lists, 0, 12);
    }

    // Apply the encoder state every main-pass encode needs: the pipeline and
    // depth state, the view uniforms, the shared vertex buffer at binding 1,
    // and the light / shadow / probe records the fragment stage reads. All of
    // it is buffers, because the ICB-executed draws inherit encoder buffer
    // bindings but reach every texture and sampler through the argument
    // buffers `execute_bindless_static_icb` binds instead.
    // Returns false without touching the encoder when the world has no main
    // pipeline (no 3D scene content); the caller then skips its draws.
    fn bind_main_pass_shared(
        &self,
        enc: &ProtocolObject<dyn objc2_metal::MTLRenderCommandEncoder>,
        view_uniforms: &ViewUniforms,
    ) -> bool {
        let Some(pipelines) = &self.cull.main_pipeline else {
            return false;
        };
        enc.set_pipeline(&pipelines.main);
        enc.set_depth_stencil(&self.targets.depth_state);

        enc.set_vertex_value(view_uniforms, 0);
        enc.set_fragment_value(view_uniforms, 0);
        enc.set_vertex_buffer(&self.scene.vertex_buffer, 0, 1);
        enc.set_fragment_value(&self.light_uniforms, 4);
        // Local-light storage buffer at fragment buffer(8). Encoder-bound
        // buffers are inherited by the ICB-executed bindless draws (the same
        // way the object buffer at binding 9 is), so this single bind covers
        // the main pass and the planar / probe re-renders.
        enc.set_fragment_buffer(&self.scene.local_light_buffer, 0, 8);
        // Per-slice spot shadow projections at fragment buffer(13), inherited
        // by the ICB draws like the local lights above. The matching depth
        // array is not bound here: every texture the pass samples travels in
        // the BindlessTextures argument buffer, because discrete texture
        // bindings break ICB compatibility.
        enc.set_fragment_buffer(&self.spot_shadow.buffer, 0, 13);
        // Rect area-light extents at fragment buffer(14), indexed by
        // GpuLight.data_index. Inherited by the ICB draws like the buffers
        // above, so this one bind covers every main-pass variant.
        enc.set_fragment_buffer(&self.scene.area_light_buffer, 0, 14);
        enc.set_fragment_value(&self.shadow.uniforms, 5);
        // Reflection-probe count at fragment buffer(6) and one parallax record
        // per baked probe at buffer(15); the cube array rides the texture
        // argument buffer. The shader weights every box covering the surface.
        self.probe_bindings().bind(
            enc,
            crate::metal::probe_set::ProbeSlots {
                set: 6,
                records: 15,
                cubes: None,
                cluster: None,
            },
        );
        true
    }

    // Encode the frame's geometry: the GPU-driven bindless pass, one indirect
    // range per shader bucket plus the folded skinned tail. Returns 0 without
    // touching the encoder when the world has nothing to draw this frame.
    fn encode_main_geometry_into(
        &self,
        enc: &ProtocolObject<dyn objc2_metal::MTLRenderCommandEncoder>,
        view_uniforms: &ViewUniforms,
        gpu: GpuFrameBuffers,
        // ICB to execute instead of the main cull's `self.cull.icb`: the planar
        // mirror render passes its slot's mirror ICB (culled against the
        // reflected frustum) here; the main + probe paths pass `None` to use the
        // main ICB.
        icb_override: Option<&ProtocolObject<dyn objc2_metal::MTLIndirectCommandBuffer>>,
    ) -> u32 {
        let GpuFrameBuffers {
            object_buffer,
            bindless_tex_args,
            ..
        } = gpu;
        enc.pushDebugGroup(&objc2_foundation::NSString::from_str("main geometry"));
        if !self.bind_main_pass_shared(enc, view_uniforms) {
            enc.popDebugGroup();
            return 0;
        }

        // The planar mirror override is a single command stream executed under
        // the encoder's default pipeline; the main + probe paths execute the
        // per-bucket main cull set.
        let override_holder;
        let bucket_icbs: &[Retained<_>] = match icb_override {
            Some(icb) => {
                override_holder = [objc2::Message::retain(icb)];
                &override_holder
            }
            None => &self.cull.icbs,
        };
        let mut draw_calls: u32 = 0;
        if let (Some(obj_buf), Some(tex_args), false) =
            (object_buffer, bindless_tex_args, bucket_icbs.is_empty())
        {
            // Shared with the phase-2 main pass under two-pass occlusion: see
            // `execute_bindless_static_icb`.
            draw_calls +=
                self.execute_bindless_static_icb(enc, obj_buf, tex_args, bucket_icbs, gpu);
        }
        enc.popDebugGroup();
        draw_calls
    }
}