concinnity-device 0.19.119

//! Metal-side executor for the render graph. `MtlContext::execute_graph`
//! walks a `CompiledGraph` and dispatches each pass by `PassId` to the
//! existing `encode_*` method. Every Metal pass that ever ran inline is
//! now in the graph. Composite plus Shadow, Main, Cull, AutoExposure,
//! Bloom, Velocity, TaaResolve, SsrResolve, ParticlesSim, ParticlesDraw,
//! Fog, Decals, GBufferPrepass, and SsaoBlur are the dispatchable PassIds.
//! PassIds `SsaoDepth` and `SsaoKernel` are timing-only: their per-pass
//! timing slots fire from `diagnostics.pass_timing.attach_*` calls inside
//! the bundled `encode_ssao` Rust function, but they must never appear as
//! graph nodes (the executor rejects them with a clear error if mis-added).
//!
//! Per-pass command buffers, two queues. Each non-composite pass runs on its own
//! freshly-minted `MTLCommandBuffer`, encoded on a rayon worker and committed by
//! the main thread onto the queue the graph's schedule assigned it
//! (`CompiledPass::queue`). Each queue's buffers commit in ascending compiled
//! index, so a queue's own commit order is that queue's graph order; the two
//! queues are ordered against each other only by the `MTLEvent` signal / wait
//! pairs the schedule derived, laid out by `metal/graph_events.rs` and carried
//! by `metal/graph_queues.rs`.
//!
//! The `Composite` pass keeps using the outer cmd_buf that `draw_frame` owns (so
//! `presentDrawable` + the completion handler stay attached to the cmd buf that
//! actually writes to the drawable). It is the last graphics pass, so it also
//! carries that queue's frame terminal signal, and `draw_frame` records the
//! terminal only once it has committed the buffer.
//!
//! Ordering within a queue is still submission order, never events. An earlier
//! draft had workers commit their own cmd bufs in arbitrary thread-schedule
//! order with an `MTLEvent` chain enforcing GPU ordering, and the renderer drew
//! into a black drawable: Apple's command queue executes cmd bufs FIFO in commit
//! order regardless of events, so committing out of order broke the dependency
//! chain (later passes ran while earlier passes' writes were still queued behind
//! them). That is exactly why the asynchronous passes need a *second* queue
//! rather than out-of-order commits on one, and why events only ever cross
//! between the two.
//!
//! Per-pass command buffers also sidestep the `MTLParallelRenderCommandEncoder`
//! abort that reliably trips G14X (M2/M3 Pro/Max-class GPUs) on macOS 26.4 after
//! ~20-90 s of rendering, regardless of how few sub-encoders we minted.
//!
//! The executor is a `&mut self` method on `MtlContext` taking the
//! concrete per-frame params.
//!
//! Per-pass barriers (`pass.barriers_before` and `pass.barriers_after`) are not
//! applied on Metal: the DX/VK seam (`barrier_translate` + a per-resource
//! registry + `emit_graph_barriers`) emits explicit resource-state TRANSITIONS,
//! and Metal has none to emit.
//!
//!   1. Hazards are tracked automatically -- within a queue. Every resource here
//!      uses the default tracked hazard mode (no `MTLHeap` / untracked
//!      resources). Apple documents that mode as "delay write operations until
//!      all previous read operations finish" and "prevent subsequent commands
//!      from running until write operations finish", without claiming a scope
//!      wider than the queue; the resource-synchronization overview places the
//!      mechanisms in a ladder where a fence "synchronizes resource memory
//!      operations across different passes within a command queue" and only an
//!      `MTLEvent` "synchronizes resource memory operations in passes across all
//!      command queues". So automatic tracking is taken to cover encoders and
//!      command buffers on ONE queue only, and every cross-queue edge is carried
//!      by an event instead: the in-frame ones the schedule derived, plus the
//!      frame-start wait that closes the cross-frame hazard on the persistent
//!      resources both queues touch (see `metal/graph_events.rs`).
//!   2. Same-queue cross-pass ordering is free. A queue's passes commit in that
//!      queue's graph order, so the producer -> consumer ordering the two barrier
//!      lists encode is already guaranteed by submission order, whichever side of
//!      a pass the graph chose to record a read run's transition on.
//!   3. The only Metal "barrier-analogue" is `useResource` residency, and it is
//!      a DIFFERENT concern the graph cannot drive: it is per-encoder (every
//!      encoder reaching a resource INDIRECTLY -- through an ICB, an argument
//!      buffer, or an acceleration structure -- must declare it), not a one-shot
//!      producer -> consumer transition, and it covers resources the graph does
//!      not model (the bindless texture pool, the env maps, the accel BLAS). So
//!      it stays inline and comprehensive at each indirect-access encoder: the
//!      ICB write residency + the bindless main pass `use_bindless_textures`
//!      (cull.rs), and the RT trace (rt_reflections.rs / raytrace.rs).
//!
//! The Vulkan / DirectX executors consume the same `BarrierOp` list; Metal reads
//! the graph for ordering + resource lifetimes only. (If untracked / heap
//! resources are ever introduced, the point-1 assumption breaks even within a
//! queue, and explicit `MTLFence`s become necessary.)

use concinnity_core::gfx::frustum::Frustum;
use concinnity_core::gfx::render_types::{
    ClusterParams, FogFroxelParams, FogParams, RtParams, SsaoParams, SsgiParams, SsrParams,
    TextDrawCall,
};
use concinnity_core::render::error::{RenderError, RenderResult};
use concinnity_core::render::planar_reflection::PlanarFramePlan;
use concinnity_core::render::reactive_mask::ReactiveMaskPlan;
#[cfg(debug_assertions)]
use concinnity_core::render::render_graph;
use concinnity_core::render::render_graph::{CompiledGraph, PassId, PassQueue};
use concinnity_core::render::uniforms::{GBufferView, PassCamera};
use concinnity_host::thread::jobs;
use objc2::rc::Retained;
use objc2::runtime::ProtocolObject;
use objc2_metal::{MTLBuffer, MTLCommandBuffer, MTLCommandQueue as _, MTLTexture};
use std::sync::atomic::Ordering;

use super::context::MtlContext;
use super::draw::main::ClusterGrid;
use super::frame_pacing::FrameJoin;
use super::graph_events;
use super::graph_events::PassSync;
use super::graph_queues::GraphQueues;
use super::parallel_encoder::{ParallelCtxRef, SendableCmdBuf};

// What `execute_graph` leaves for `draw_frame` to finish. The composite pass
// rides the command buffer `draw_frame` owns, so the graphics queue's frame
// terminal is signaled on a buffer this executor never commits.
pub(in crate::metal) struct GraphSubmission {
    // The graphics terminal value the composite command buffer signals, to be
    // handed to `MtlContext::record_graph_terminal` once it is committed.
    pub(in crate::metal) pending_terminal: Option<u64>,
}

// Per-frame params the executor threads into each pass's `encode_*`
// method. The union of what every graph pass reads; a pass ignores the fields
// it does not consume. `scene_color` is `Option` because the pre-graph runs before TAA /
// SSR resolve has produced it; the post-graph (which contains
// Composite) supplies it.
//
// `Send + Sync` is asserted so the struct can be shared by reference into
// the rayon::scope fan-out inside `execute_graph`. Workers only read from
// it; the only non-Sync field is `cmd_buf`, which workers do not use
// (each worker mints its own `MTLCommandBuffer`). `cmd_buf` is consumed
// solely on the main thread for the Composite pass and the present /
// completion handler attached by `draw_frame`.
pub(in crate::metal) struct GraphFrameParams<'a> {
    pub cmd_buf: &'a ProtocolObject<dyn MTLCommandBuffer>,
    pub cam_pos: [f32; 3],
    pub skinned_joint_bufs: &'a [Retained<ProtocolObject<dyn MTLBuffer>>],
    // Per-skinned-object morph weights for this frame, from the same ring slot
    // as `skinned_joint_bufs`. Read only by the Cull pass's skin fold.
    pub skinned_morph_weight_bufs: &'a [Retained<ProtocolObject<dyn MTLBuffer>>],
    pub scene_color: Option<&'a Retained<ProtocolObject<dyn MTLTexture>>>,
    pub text_calls: &'a [TextDrawCall],
    // This frame's transient-buffer ring slot (`frame_ring_index % frames_in_flight`).
    // The auto-exposure pass writes the readback buffer at this slot.
    pub ring_slot: usize,
    // An opaque menu backdrop hides the scene: the Main pass runs as a bare
    // clear (it is the only surviving world pass in the masked graph), skipping
    // every geometry sub-path so nothing of the world draws behind the menu.
    pub world_hidden: bool,
    // Main-pass params, computed by draw_frame before the graph dispatch.
    pub elapsed: f32,
    pub vp: [[f32; 4]; 4],
    // Inverse of `vp`, computed once in `draw_frame` and shared by every pass
    // that reconstructs world-space position from depth (fog / decals /
    // raymarch / transparent) instead of each re-inverting `vp`.
    pub inv_vp: [[f32; 4]; 4],
    pub frustum: &'a Frustum,
    // Which planar mirrors this frame renders, and the screen rectangle each
    // covers; computed once from `vp` so the mirror pass and the transparent
    // pass that samples it agree.
    pub planar: PlanarFramePlan,
    pub object_buffer: Option<&'a Retained<ProtocolObject<dyn MTLBuffer>>>,
    // This frame's copy of the material parameter table; `Some` with
    // `object_buffer`.
    pub material_params: Option<&'a Retained<ProtocolObject<dyn MTLBuffer>>>,
    pub bindless_tex_args: Option<&'a Retained<ProtocolObject<dyn MTLBuffer>>>,
    // This frame's skinned deformed-vertex buffer. `Some` only when a
    // SkinnedMesh has uploaded. The Cull pass's `encode_main_skin` writes it; the Main /
    // Main2 skinned ICB tail binds it as the vertex buffer for that draw range.
    pub deformed_skinned: Option<&'a Retained<ProtocolObject<dyn MTLBuffer>>>,
    // The previous-frame skinned deformed-vertex buffer: the slot one
    // frame behind `deformed_skinned` in the per-frame deformed ring. The
    // GPU-driven G-buffer skinned tail reads it as the previous vertex stream to
    // emit a per-vertex skin motion vector. `Some` only when the fold is active;
    // the priming gate (`deformed_primed`) handles the unposed first frame.
    pub deformed_prev: Option<&'a Retained<ProtocolObject<dyn MTLBuffer>>>,
    // Per-frame parallel `prev_model` buffer for the GPU-driven G-buffer pass:
    // one `float4x4` per cull record, read by the bindless G-buffer
    // VS at `[[base_instance]]`. `Some` only when the GPU-driven pre-pass runs.
    pub prev_model_buffer: Option<&'a Retained<ProtocolObject<dyn MTLBuffer>>>,
    // Model-history ring slots this frame's snapshot fills, after the pre-pass
    // has read `prev_model_buffer`.
    pub history_targets: &'a [Retained<ProtocolObject<dyn MTLBuffer>>],
    // GPU-cull output the bindless Main pass consumes via
    // executeCommandsInBuffer. `Some` only when bindless cull ran this
    // frame (i.e. matches `FrameGraphInputs::bindless_cull_enabled`).
    pub draw_args_buffer: Option<&'a Retained<ProtocolObject<dyn MTLBuffer>>>,
    // The G-buffer pre-pass's view block: motion matrices and the previous
    // clock, collapsed onto this frame's own when `velocity_active` is unset.
    pub gbuffer_view: &'a GBufferView,
    // A consumer reads the pre-pass's motion this frame.
    pub velocity_active: bool,
    // Pre-TAA scene texture that `TaaResolve` reads (the SSR resolve
    // output when SSR is on, otherwise the raw `hdr_resolve`). `Some`
    // only when the `TaaResolve` pass is in the graph this frame.
    pub scene_pre_taa: Option<&'a Retained<ProtocolObject<dyn MTLTexture>>>,
    // SSR ray-march params (96 B, copy-able). `Some` only when the
    // `SsrResolve` pass is in the graph this frame (matches
    // `FrameGraphInputs::ssr_enabled`).
    pub ssr_params: Option<&'a SsrParams>,
    // Volumetric-fog pass uniforms. `Some` only when the `Fog` pass is
    // in the graph this frame (matches `FrameGraphInputs::fog_enabled`).
    pub fog_params: Option<&'a FogParams>,
    // Metal-only froxel-volume extras (view matrix + volume dims +
    // near/far). `Some` only when `Fog` is in the graph this frame; the
    // FogFroxel compute pass + the Fog fragment shader sample path both
    // consume it.
    pub fog_froxel_params: Option<&'a FogFroxelParams>,
    // Clustered light-binning params. `Some` only when the `LightCull` pass is
    // in the graph this frame (matches `FrameGraphInputs::clustering_enabled`).
    pub cluster_params: Option<&'a ClusterParams>,
    // SSAO kernel + blur params. `Some` only when the `SsaoBlur` pass
    // is in the graph this frame (matches `FrameGraphInputs::ssao_enabled`).
    pub ssao_params: Option<&'a SsaoParams>,
    // SSGI pyramid, trace and composite params. `Some` only when the `Ssgi` pass is
    // in the graph this frame (matches `FrameGraphInputs::ssgi_enabled`).
    pub ssgi_params: Option<&'a SsgiParams>,
    // RT-reflection params (camera + sun + tunables). `Some` only when the
    // `RtReflections` pass is in the graph this frame (matches
    // `FrameGraphInputs::rt_reflections_enabled`).
    pub rt_reflection_params: Option<&'a RtParams>,
    // How the particle and transparent passes treat the reactive mask, and
    // whether TAA or the upscaler reads it.
    pub reactive: ReactiveMaskPlan,
}

// SAFETY: see the type-level docs. The non-Sync `cmd_buf` field is used
// exclusively on the main thread (Composite + present + completion
// handler); every worker spawned by `execute_graph` mints its own cmd buf.
// All other fields are Sync (POD or thread-safe Apple Metal handles).
unsafe impl<'a> Send for GraphFrameParams<'a> {}
// SAFETY: as for `Send` above.
unsafe impl<'a> Sync for GraphFrameParams<'a> {}

fn pass_input<T>(value: Option<T>, pass: PassId, input: &str) -> RenderResult<T> {
    value.ok_or_else(|| {
        RenderError::Other(format!(
            "graph executor: {pass:?} pass requires {input} but none was supplied"
        ))
    })
}

fn encode_waits(
    cmd_buf: &ProtocolObject<dyn MTLCommandBuffer>,
    queues: &GraphQueues,
    sync: &PassSync,
) {
    for &(event_queue, value) in &sync.waits {
        cmd_buf.encodeWaitForEvent_value(queues.event(event_queue), value);
    }
}

fn encode_signals(
    cmd_buf: &ProtocolObject<dyn MTLCommandBuffer>,
    queues: &GraphQueues,
    sync: &PassSync,
    queue: PassQueue,
) {
    for &value in &sync.signals {
        cmd_buf.encodeSignalEvent_value(queues.event(queue), value);
    }
}

impl MtlContext {
    // Walk a compiled render graph and dispatch each pass to its
    // existing per-backend encoder. `params` carries the per-frame state
    // that cannot live on `&mut self` (the command buffer is per-frame;
    // the scene-color texture is computed each frame after SSR / TAA
    // resolve).
    //
    // Any pass not matched by the match arm below returns an error so
    // a caller that drops an unhandled PassId into the graph by
    // mistake fails loudly rather than silently no-op'ing.
    pub(in crate::metal) fn execute_graph(
        &mut self,
        graph: &CompiledGraph,
        params: &GraphFrameParams<'_>,
        join: &std::sync::Arc<FrameJoin>,
    ) -> RenderResult<GraphSubmission> {
        #[cfg(debug_assertions)]
        render_graph::assert_slot_aliasing_sound(
            graph,
            self.targets.transient_pool.slot_labels(),
            "metal",
        );
        // Both submission paths need the compiled order to be a topological
        // order for each queue at once, with every wait naming a producer
        // recorded earlier: the two-queue path because it commits each queue's
        // buffers in ascending compiled index and encodes a wait for a value
        // that queue signals no later, the single-queue fallback because it
        // flattens the schedule back into one stream. This asserts both.
        #[cfg(debug_assertions)]
        render_graph::assert_serial_order_honors_schedule(graph, "metal");

        // Per-frame particle-state mutations live on `&mut self` and have
        // to happen before the read-only `encode_particles` path runs. We
        // build the (dt, frame_index, per-emitter spawn budgets) tuple
        // once here and stash it for the match arm below.
        let particle_frame = self.prepare_particle_pass(params.elapsed);
        self.diagnostics
            .draw_calls_accum
            .store(0, Ordering::Relaxed);

        let composite_idx = graph
            .passes
            .iter()
            .position(|p| matches!(p.id, PassId::Composite));
        // The graphics queue's terminal signal rides whichever pass the plan put
        // it on. That is the composite pass in every graph that presents, and
        // its command buffer is committed by `draw_frame` rather than here, so
        // the terminal is handed back and recorded after that commit.
        let last_graphics = (0..graph.passes.len())
            .rev()
            .find(|&i| graph.passes[i].queue == PassQueue::Graphics);
        let deferred_terminal = (composite_idx.is_some() && composite_idx == last_graphics)
            .then_some(PassQueue::Graphics);

        // This frame's event values. `None` when the second queue could not be
        // created: the fallback records every pass onto the graphics queue in
        // compiled order and encodes no events at all.
        let plan = self.hw.graph_queues.as_ref().map(|queues| {
            let (events, previous) = queues.begin_frame(graph.passes.len());
            graph_events::plan_frame(graph, events, previous)
        });

        // Pre-allocate per-pass slots. Workers encode into their own
        // freshly-minted `MTLCommandBuffer` in parallel and hand the
        // encoded-but-uncommitted buffer back through the matching slot.
        // The main thread then commits each slot onto its own queue in
        // compiled order, so each queue's commit order is that queue's graph
        // order, which is that queue's GPU execution order.
        let worker_slots: std::sync::Mutex<Vec<Option<SendableCmdBuf>>> =
            std::sync::Mutex::new((0..graph.passes.len()).map(|_| None).collect());
        let first_error: std::sync::Mutex<Option<RenderError>> = std::sync::Mutex::new(None);

        let plan_ref = plan.as_ref();
        let ctx_ref = ParallelCtxRef::new(self);
        jobs::pool().install(|| {
            rayon::scope(|scope| {
                for (idx, pass) in graph.passes.iter().enumerate() {
                    if Some(idx) == composite_idx {
                        continue;
                    }
                    let pass_id = pass.id;
                    let pass_queue = pass.queue;
                    let particle_ref = particle_frame.as_ref();
                    let first_error_ref = &first_error;
                    let worker_slots_ref = &worker_slots;
                    // A rayon worker drains no autorelease pool of its own, so
                    // without this one every autoreleased object the pass
                    // encode creates (the command buffer and its encoders
                    // among them) would be retained for the life of the
                    // process.
                    scope.spawn(move |_| {
                        objc2::rc::autoreleasepool(|_| {
                            let ctx = ctx_ref.as_ctx();
                            let queue = match ctx.hw.graph_queues.as_ref() {
                                Some(queues) => queues.queue(pass_queue, &ctx.hw.command_queue),
                                None => &ctx.hw.command_queue,
                            };
                            let cmd_buf = match queue.commandBuffer() {
                                Some(cb) => cb,
                                None => {
                                    let mut e = first_error_ref.lock().unwrap();
                                    if e.is_none() {
                                        *e = Some(RenderError::Other(
                                            "graph executor: failed to mint per-pass cmd buf"
                                                .into(),
                                        ));
                                    }
                                    return;
                                }
                            };
                            // Waits before the pass's own encoders, signals
                            // after them: Metal accepts an event command only
                            // while the command buffer has no open encoder.
                            let sync = ctx.hw.graph_queues.as_ref().zip(plan_ref);
                            if let Some((queues, plan)) = sync {
                                encode_waits(&cmd_buf, queues, plan.pass(idx));
                            }
                            match ctx.encode_pass_into(pass_id, &cmd_buf, params, particle_ref) {
                                Ok(count) => {
                                    if let Some((queues, plan)) = sync {
                                        encode_signals(
                                            &cmd_buf,
                                            queues,
                                            plan.pass(idx),
                                            pass_queue,
                                        );
                                    }
                                    ctx.diagnostics
                                        .draw_calls_accum
                                        .fetch_add(count, Ordering::Relaxed);
                                    let mut lock = worker_slots_ref.lock().unwrap();
                                    lock[idx] = Some(SendableCmdBuf(cmd_buf));
                                }
                                Err(e) => {
                                    let mut lock = first_error_ref.lock().unwrap();
                                    if lock.is_none() {
                                        *lock = Some(e);
                                    }
                                }
                            }
                        })
                    });
                }
            });
        });

        // Nothing has been committed yet, so an encode failure leaves the
        // frame's event values unsignaled and unrecorded: the next frame
        // reuses the slice rather than waiting on a value nothing reaches.
        if let Some(err) = first_error.into_inner().unwrap_or(None) {
            return Err(err);
        }

        // Commit every worker-encoded cmd buf onto its own queue, walking the
        // compiled order so each queue sees its passes in graph order.
        let slots = worker_slots.into_inner().map_err(|_| {
            RenderError::Other("graph executor: worker slot mutex poisoned".to_string())
        })?;
        for (idx, slot) in slots.into_iter().enumerate() {
            if let Some(cb) = slot {
                // Each pass commits its own command buffer, so its handler names
                // the pass when the GPU faults it. The same handler reports the
                // buffer to the frame's completion join, so the frame-in-flight
                // slot outlives every queue's work rather than just the
                // presenting buffer's.
                let pass_id = graph.passes.get(idx).map(|p| p.id);
                let part = std::sync::Arc::clone(join);
                join.add_part();
                let handler = block2::RcBlock::new(
                    move |cbh: std::ptr::NonNull<ProtocolObject<dyn MTLCommandBuffer>>| {
                        // SAFETY: Metal hands the completion handler a live command buffer, and the
                        // borrow does not escape the block.
                        let cbh = unsafe { cbh.as_ref() };
                        match pass_id {
                            Some(id) => super::fault_log::report_fault(
                                cbh,
                                format_args!("render-graph pass {id:?}"),
                            ),
                            None => super::fault_log::report_fault(cbh, "render-graph pass"),
                        }
                        part.arrive();
                    },
                );
                // SAFETY: addCompletedHandler copies the block, so the RcBlock
                // may drop at end of iteration.
                unsafe {
                    cb.0.addCompletedHandler(block2::RcBlock::as_ptr(&handler));
                }
                cb.0.commit();
            }
        }

        // Every terminal but the deferred one is now committed, so the next
        // frame may wait on it. Also advances the value slice.
        let mut submission = GraphSubmission {
            pending_terminal: None,
        };
        if let (Some(queues), Some(plan)) = (self.hw.graph_queues.as_mut(), plan.as_ref()) {
            queues.end_submission(plan, deferred_terminal);
            submission.pending_terminal = deferred_terminal.and_then(|q| plan.terminal(q));
        }

        // Composite stays on the outer cmd buf that `draw_frame` owns, so
        // `presentDrawable` + the completion handler attach to the same cmd buf
        // that writes to the drawable. It is committed by `draw_frame` after
        // this returns; every other graphics-queue cmd buf has already
        // committed, so the queue order places it strictly after them.
        if let Some(idx) = composite_idx {
            let sync = self.hw.graph_queues.as_ref().zip(plan.as_ref());
            if let Some((queues, plan)) = sync {
                encode_waits(params.cmd_buf, queues, plan.pass(idx));
            }
            let count = self.encode_pass_into(
                PassId::Composite,
                params.cmd_buf,
                params,
                particle_frame.as_ref(),
            )?;
            if let Some((queues, plan)) = sync {
                encode_signals(params.cmd_buf, queues, plan.pass(idx), PassQueue::Graphics);
            }
            self.diagnostics
                .draw_calls_accum
                .fetch_add(count, Ordering::Relaxed);
        }

        self.diagnostics.frame_stats.draw_calls +=
            self.diagnostics.draw_calls_accum.load(Ordering::Relaxed);
        Ok(submission)
    }

    // Record the frame terminal whose command buffer `draw_frame` commits
    // itself. Called after that commit, so a frame abandoned between recording
    // and committing never leaves the next frame waiting on a value the GPU is
    // not going to reach.
    pub(in crate::metal) fn record_graph_terminal(&mut self, value: u64) {
        if let Some(queues) = self.hw.graph_queues.as_mut() {
            queues.record_terminal(PassQueue::Graphics, value);
        }
    }

    // Dispatch a single pass into a freshly-minted `MTLCommandBuffer`. Takes
    // Build the per-frame `RaymarchView` from the graph params. Shared by the
    // `Raymarch` pass (live SDF surface) and the `Shadow` pass (SDF shadow
    // casters) so both agree on the camera VP / time / viewport: the shadow
    // fragment only reads `time`, but constructing the full view keeps the
    // binding identical to the main pass.
    fn build_raymarch_view(&self, params: &GraphFrameParams<'_>) -> super::raymarch::RaymarchView {
        super::raymarch::RaymarchView::new(&self.pass_camera(params))
    }

    // The camera the raymarch and transparent passes reconstruct from: the
    // frame's jittered VP and its inverse at the HDR target size, with this
    // frame's IBL and sky rotation.
    fn pass_camera(&self, params: &GraphFrameParams<'_>) -> PassCamera {
        PassCamera {
            vp: params.vp,
            inv_vp: params.inv_vp,
            cam_pos: params.cam_pos,
            viewport: [
                self.targets.hdr.width as f32,
                self.targets.hdr.height as f32,
            ],
            time: params.elapsed,
            prefilter_mip_count: self.scene.env_map.prefilter_mip_count as f32,
            sky_rot: self.state.view.sky_rot,
        }
    }

    // `&self` so it can be called from both the main-thread Composite path
    // and the rayon-spawned per-pass workers.
    fn encode_pass_into(
        &self,
        pass_id: PassId,
        cmd_buf: &ProtocolObject<dyn MTLCommandBuffer>,
        params: &GraphFrameParams<'_>,
        particle_frame: Option<&super::particle::ParticleFrame>,
    ) -> RenderResult<u32> {
        Ok(match pass_id {
            PassId::Cull => {
                let object_buffer =
                    pass_input(params.object_buffer, PassId::Cull, "object_buffer")?;
                let draw_args_buffer =
                    pass_input(params.draw_args_buffer, PassId::Cull, "draw_args_buffer")?;
                // Skinned fold: pre-skin into this frame's deformed
                // buffer in the Cull command buffer (before the cull dispatch).
                // Committed before Main, so Metal hazard-tracks the deformed
                // write ahead of the main pass's vertex read. Only when the fold
                // is active (deformed buffer supplied).
                if let Some(deformed) = params.deformed_skinned {
                    self.encode_main_skin(
                        cmd_buf,
                        deformed,
                        super::raytrace::MainSkinBuffers {
                            joints: params.skinned_joint_bufs,
                            morph_weights: params.skinned_morph_weight_bufs,
                        },
                    )?;
                }
                self.encode_cull(
                    cmd_buf,
                    object_buffer,
                    draw_args_buffer,
                    params.frustum,
                    params.cam_pos,
                    self.draw_record_counts(),
                )?;
                // GPU-driven shadow views: fill the cascade and spot-slice ICBs
                // in this same Cull command buffer (committed before the Shadow
                // and SpotShadow render passes' command buffers), so the
                // cross-command-buffer FIFO order makes them ready when those
                // passes read them -- the same ordering the main cull -> Main ICB
                // relies on. A no-op when the shadow-bindless path is inactive.
                self.encode_shadow_culls(cmd_buf, object_buffer, draw_args_buffer)?;
                0
            }
            PassId::HizBuild | PassId::HizFinal => {
                // Two Hi-Z builds share one encoder. `HizBuild` rebuilds the
                // pyramid mid-frame from phase-1 depth so Cull2 re-tests
                // against up-to-date occluders; `HizFinal` reduces the frame's
                // final depth for the next frame's phase-1 cull.
                self.encode_hiz_build(cmd_buf);
                0
            }
            PassId::Cull2 => {
                let object_buffer =
                    pass_input(params.object_buffer, PassId::Cull2, "object_buffer")?;
                let draw_args_buffer =
                    pass_input(params.draw_args_buffer, PassId::Cull2, "draw_args_buffer")?;
                self.encode_cull_phase2(
                    cmd_buf,
                    object_buffer,
                    draw_args_buffer,
                    params.frustum,
                    params.cam_pos,
                )?
            }
            PassId::Main2 => self.encode_main_pass_phase2(
                cmd_buf,
                crate::metal::draw::main::MainPassCamera {
                    elapsed: params.elapsed,
                    vp: params.vp,
                    view: self.state.view.matrix,
                    cam_pos: params.cam_pos,
                },
                crate::metal::draw::main::GpuFrameBuffers {
                    object_buffer: params.object_buffer,
                    material_params: params.material_params,
                    bindless_tex_args: params.bindless_tex_args,
                    deformed_skinned: params.deformed_skinned,
                    counts: self.draw_record_counts(),
                },
            )?,
            PassId::Shadow => {
                // Build the raymarch view only when a volume opts into
                // shadow casting; otherwise pass `None` so the shadow encoder
                // skips the SDF caster sub-pass with zero overhead. The view
                // matches the matching `PassId::Raymarch` build later this
                // frame (same camera VP / time / viewport). Mirrors DirectX.
                let raymarch_view = if self.any_raymarch_shadow_casters() {
                    Some(self.build_raymarch_view(params))
                } else {
                    None
                };
                self.encode_shadow_pass(
                    cmd_buf,
                    params.object_buffer,
                    params.deformed_skinned,
                    raymarch_view.as_ref(),
                )?
            }
            PassId::SpotShadow => self.encode_spot_shadow_pass(
                cmd_buf,
                params.object_buffer,
                params.deformed_skinned,
            )?,
            PassId::Main => self.encode_main_pass(
                cmd_buf,
                crate::metal::draw::main::MainPassCamera {
                    elapsed: params.elapsed,
                    vp: params.vp,
                    view: self.state.view.matrix,
                    cam_pos: params.cam_pos,
                },
                crate::metal::draw::main::GpuFrameBuffers {
                    object_buffer: params.object_buffer,
                    material_params: params.material_params,
                    bindless_tex_args: params.bindless_tex_args,
                    deformed_skinned: params.deformed_skinned,
                    counts: self.draw_record_counts(),
                },
                params.world_hidden,
            )?,
            PassId::AutoExposure => self.encode_auto_exposure(cmd_buf, params.ring_slot)?,
            PassId::Bloom => {
                let scene_color = pass_input(params.scene_color, PassId::Bloom, "scene_color")?;
                self.encode_bloom(cmd_buf, scene_color)?
            }
            PassId::GBufferPrepass => {
                // The surfaces rasterize through the main pass's view block, so
                // the vertex hook places them as the main pass will; the
                // un-jittered cur/prev VPs drive the motion vector.
                let main_view =
                    self.main_view_uniforms(&crate::metal::draw::main::MainPassCamera {
                        elapsed: params.elapsed,
                        vp: params.vp,
                        view: self.state.view.matrix,
                        cam_pos: params.cam_pos,
                    });
                let raymarch_view = super::raymarch::RaymarchView::for_gbuffer(
                    &self.pass_camera(params),
                    params.gbuffer_view,
                );
                self.encode_gbuffer_prepass(
                    cmd_buf,
                    crate::metal::post::gbuffer::GbufferPrepassViews {
                        gbuffer: params.gbuffer_view,
                        main: &main_view,
                        raymarch: &raymarch_view,
                        frustum: params.frustum,
                    },
                    crate::metal::post::gbuffer::GbufferGpuBuffers {
                        object_buffer: params.object_buffer,
                        material_params: params.material_params,
                        bindless_tex_args: params.bindless_tex_args,
                        prev_model_buffer: params.prev_model_buffer,
                        draw_args_buffer: params.draw_args_buffer,
                        history_targets: params.history_targets,
                        deformed_current: params.deformed_skinned,
                        deformed_prev: params.deformed_prev,
                    },
                    params.velocity_active,
                )?
            }
            PassId::TaaResolve => {
                let scene_pre_taa =
                    pass_input(params.scene_pre_taa, PassId::TaaResolve, "scene_pre_taa")?;
                self.encode_taa(cmd_buf, scene_pre_taa, params.reactive.readable)?
            }
            PassId::SsrResolve => {
                let ssr_params = pass_input(params.ssr_params, PassId::SsrResolve, "ssr_params")?;
                self.encode_ssr_resolve(cmd_buf, ssr_params)?
            }
            PassId::Ssgi => {
                let ssgi_params = pass_input(params.ssgi_params, PassId::Ssgi, "ssgi_params")?;
                self.encode_ssgi(cmd_buf, ssgi_params)?
            }
            PassId::RtReflections => {
                let rt_params = pass_input(
                    params.rt_reflection_params,
                    PassId::RtReflections,
                    "rt_reflection_params",
                )?;
                self.encode_rt_reflections(cmd_buf, rt_params, params.bindless_tex_args)?
            }
            PassId::SsaoBlur => {
                // PassId::SsaoBlur dispatches the bundled `encode_ssao` (GTAO
                // depth copy, kernel + depth-aware blur). It reads the unified
                // G-buffer pre-pass output, so SSAO runs no geometry redraw of
                // its own; per-pass timing for the sub-passes is wired inline
                // inside `encode_ssao`.
                let ssao_params = pass_input(params.ssao_params, PassId::SsaoBlur, "ssao_params")?;
                self.encode_ssao(cmd_buf, ssao_params)?
            }
            PassId::SsaoDepth | PassId::SsaoKernel => {
                // Bundled inside `encode_ssao` (dispatched via
                // PassId::SsaoBlur). These PassIds keep their
                // per-pass timing slots via inline
                // `diagnostics.pass_timing.attach_render` calls inside
                // encode_ssao, but they must not appear as their
                // own graph nodes.
                return Err(RenderError::Other(format!(
                    "graph executor: pass {} is bundled inside SsaoBlur \
                         (encode_ssao encodes all three SSAO sub-passes); it \
                         should not appear as its own graph node",
                    pass_id.name()
                )));
            }
            PassId::Sky => {
                // Drawn inline at the tail of Main and of every probe face and
                // mirror render; it only names a timing slot.
                return Err(RenderError::Other(format!(
                    "graph executor: pass {} is drawn inline by the opaque scene \
                         passes; it should not appear as its own graph node",
                    pass_id.name()
                )));
            }
            PassId::ReflectionComposite => {
                // Encoded inline at the tail of SsrResolve / RtReflections (it
                // blurs + composites the reflection target they wrote). Keeps a
                // timing slot via an inline `attach_render`, but is never a graph
                // node of its own -- same pattern as the bundled SSAO sub-passes.
                return Err(RenderError::Other(format!(
                    "graph executor: pass {} is encoded inline by SsrResolve / \
                         RtReflections; it should not appear as its own graph node",
                    pass_id.name()
                )));
            }
            PassId::Decals => {
                self.encode_decals(cmd_buf, params.vp, params.inv_vp, params.frustum)?
            }
            PassId::Fog => {
                let fog_params = pass_input(params.fog_params, PassId::Fog, "fog_params")?;
                let fog_froxel_params =
                    pass_input(params.fog_froxel_params, PassId::Fog, "fog_froxel_params")?;
                self.encode_fog(cmd_buf, fog_params, fog_froxel_params)?
            }
            PassId::FogFroxel => {
                let fog_params = pass_input(params.fog_params, PassId::FogFroxel, "fog_params")?;
                let fog_froxel_params = pass_input(
                    params.fog_froxel_params,
                    PassId::FogFroxel,
                    "fog_froxel_params",
                )?;
                self.encode_fog_froxel(cmd_buf, fog_params, fog_froxel_params)?
            }
            PassId::LightCull => {
                let cluster_params =
                    pass_input(params.cluster_params, PassId::LightCull, "cluster_params")?;
                self.encode_light_cull(
                    cmd_buf,
                    ClusterGrid {
                        params: cluster_params,
                        lists: &self.light_cull.cluster_buffer,
                    },
                    Some(PassId::LightCull),
                )?
            }
            PassId::ParticlesSim => {
                // Integrates every live emitter's persistent pool in place. The
                // graph's only edge out of it is the draw's vertex-stage read of
                // those pools, so the schedule is free to put it on the async
                // queue; this executor still records it in the compiled order.
                // Per-frame particle-state mutations (`particle.last_elapsed`,
                // `particle.frame_index`, per-emitter spawn budget) ran on
                // `&mut self` before the fan-out via `prepare_particle_pass`;
                // both read-only halves consume the precomputed frame.
                if let Some(frame) = particle_frame {
                    self.encode_particles_sim(cmd_buf, frame)?;
                }
                0
            }
            PassId::ParticlesDraw => {
                let write = params.reactive.particles;
                let draws = if let Some(frame) = particle_frame {
                    self.encode_particles_draw(cmd_buf, frame, params.vp, params.frustum, write)?
                } else {
                    0
                };
                if draws == 0 {
                    self.clear_reactive_mask(cmd_buf, write)?;
                }
                draws
            }
            PassId::Lines => self.encode_lines(cmd_buf, params.vp)?,
            PassId::Composite => {
                let scene_color = pass_input(params.scene_color, PassId::Composite, "scene_color")?;
                self.encode_composite_and_text(
                    cmd_buf,
                    scene_color,
                    params.text_calls,
                    params.reactive.readable,
                )?
            }
            PassId::Upscale => {
                let scene_pre_taa =
                    pass_input(params.scene_pre_taa, PassId::Upscale, "scene_pre_taa")?;
                self.encode_upscale(cmd_buf, scene_pre_taa, params.reactive.readable)?
            }
            PassId::Transparent => {
                let scene_pre_taa =
                    pass_input(params.scene_pre_taa, PassId::Transparent, "scene_pre_taa")?;
                let view = concinnity_core::render::uniforms::TransparentView::new(
                    &self.pass_camera(params),
                    &self.light_uniforms,
                );
                // The mirrors `PlanarReflection` rendered this frame, if the pass
                // samples mirrors at all (the same gate that put the node in the
                // graph). A reflector whose slot was not rendered keeps its probe
                // path.
                let mirrors = if self.planar_mirrors_needed() {
                    params.planar
                } else {
                    PlanarFramePlan::default()
                };

                // Gather every translucent producer's draws, then let the
                // shared encoder sort them back-to-front and issue them.
                let mut draws = Vec::new();
                self.collect_water_transparent_draws(
                    &view,
                    params.bindless_tex_args.is_some(),
                    &mirrors,
                    &mut draws,
                );
                self.collect_glass_transparent_draws(
                    &view,
                    params.bindless_tex_args.is_some(),
                    &mirrors,
                    &mut draws,
                );
                // Transparent glass MESHES (Layer 2): imported `transparent`
                // materials traced per-pixel when RT is live. Inert otherwise
                // (those meshes render opaque in the main pass).
                self.collect_mesh_transparent_draws(
                    &view,
                    params.bindless_tex_args.is_some(),
                    &mut draws,
                );
                self.encode_transparent(
                    cmd_buf,
                    scene_pre_taa,
                    &draws,
                    super::transparent::TransparentFrame {
                        view: &view,
                        rt_params: params.rt_reflection_params,
                        bindless_tex_args: params.bindless_tex_args,
                        reactive: params.reactive.transparent,
                    },
                )?
            }
            PassId::PlanarReflection => {
                // Mirror renders for the flat reflectors in view, each cropped to
                // the screen rectangle its reflectors cover (see `planar.rs`).
                // `Transparent` samples them later on the same queue.
                self.encode_planar_reflections(cmd_buf, params)?;
                0
            }
            PassId::Raymarch => {
                let view = self.build_raymarch_view(params);
                self.encode_raymarch(cmd_buf, &view, params.frustum)?
            }
        })
    }
}

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

    #[test]
    fn pass_input_names_the_pass_and_the_missing_input() {
        assert_eq!(pass_input(Some(7), PassId::Fog, "fog_params"), Ok(7));
        assert_eq!(
            pass_input::<u32>(None, PassId::Fog, "fog_params"),
            Err(RenderError::Other(
                "graph executor: Fog pass requires fog_params but none was supplied".into()
            ))
        );
    }
}