concinnity-device 0.19.119

#![deny(unsafe_op_in_unsafe_fn)]

use concinnity_core::gfx::frustum::Frustum;
use concinnity_core::gfx::render_types;
use concinnity_core::gfx::render_types::{
    ClusterParams, DrawIndex, InstancedCluster, LightUniforms, NUM_SHADOW_CASCADES, ShadowUniforms,
};
use concinnity_core::profile;
use concinnity_core::render::backend;
use concinnity_core::render::backend_init;
use concinnity_core::render::cluster_range::ClusterReach;
use concinnity_core::render::decal;
use concinnity_core::render::error;
use concinnity_core::render::hdr_output;
use concinnity_core::render::particles;
use concinnity_core::render::probe_book::ProbeBook;
use concinnity_core::render::render_graph;
use concinnity_core::render::scene_flow;
use concinnity_core::render::scene_state::SceneState;
use concinnity_core::render::shadow_schedule;
use concinnity_core::render::spot_shadow;
use concinnity_core::render::view_history::ViewHistory;
use objc2::rc::Retained;
use objc2::runtime::ProtocolObject;
use objc2_metal::{
    MTLArgumentEncoder as _, MTLBuffer, MTLCommandBuffer as _, MTLCommandQueue,
    MTLDepthStencilState, MTLDevice as _, MTLIndirectCommandBuffer,
    MTLIndirectCommandBufferDescriptor, MTLIndirectCommandType, MTLPixelFormat,
    MTLRenderPipelineState, MTLResourceOptions, MTLSamplerState, MTLTexture,
};
use objc2_metal_kit::MTKView;

use super::allocator::{DeviceAllocator, PooledBuffer, PooledTexture};
use super::auto_exposure::AutoExposureGpu;
use super::cull::CullState;
use super::decal::DecalState;
use super::fog::FogState;
use super::line::LineState;
use super::particle::ParticleState;
use super::post::{
    GBufferState, MtlBloomPass, SsaoState, SsgiState, SsrState, TaaState, UpscaleState,
};
use super::raytrace::RtState;
use super::resources::skinning::SkinnedState;
use super::texture::{EnvironmentMapTextures, HdrTargets};
use super::transient_pool::TransientTexturePool;

// Size of the bindless texture pool the static main pass samples. The pool
// holds every albedo texture followed by every normal map; `GpuObjectData`
// carries pool indices into it. Worlds with more than this many textures fall
// back to clamped indices (logged once at init).
pub(super) const BINDLESS_TEXTURE_COUNT: usize =
    concinnity_core::render::uniforms::BINDLESS_POOL_SIZE;

// Fragment buffer index the bindless static pass binds its `BindlessTextures`
// argument buffer at. Discrete `[[texture(n)]]` bindings make a fragment
// shader unusable from an indirect command buffer on Apple GPUs, so the
// texture pool + shadow/IBL maps travel in an argument buffer instead. Must
// match the buffer(7) slot of the engine fragment in `main_bindless.hlsl`
// (locked by the build script's ABI assertion), which every world Shader
// compiles from.
pub(super) const BINDLESS_TEXTURE_ARG_BUFFER_INDEX: usize = 7;

// Fragment buffer index of the engine sampler block: indirect-command
// execution cannot see encoder-bound sampler state, so the engine's
// single-source fragment reads its three static samplers from this argument
// buffer.
pub(super) const BINDLESS_SAMPLER_ARG_BUFFER_INDEX: usize = 10;

// InstancedProp clusters and the per-instance records they draw through.
pub(super) struct InstancedState {
    // One entry per cluster. Each issues one drawIndexedInstanced call with all
    // of its per-instance transforms uploaded to a transient GPU buffer per
    // frame. Empty when there are no clusters in the scene.
    pub clusters: Vec<InstancedCluster>,
    // The per-instance `GpuObjectData` / `GpuDrawArgs` records, built once at
    // init (instances are placed at world load and never move).
    // `build_object_buffer` / `build_draw_args_buffer` append these after their
    // per-frame static fill, so the transient object / draw-args rings carry
    // both. `draw_args` carries each cluster's base LOD slice; the per-frame
    // build patches the instances of clusters that declare alternates.
    pub records: Vec<render_types::GpuObjectData>,
    pub draw_args: Vec<render_types::GpuDrawArgs>,
    // Whether any cluster declares LOD alternates. False skips the per-frame
    // per-instance LOD pass entirely: without alternates the base slice above
    // is the right answer for every instance, for the life of the world.
    pub any_lod: bool,
}

// Scene-captured reflection probes: each surface's specular reflection samples
// the nearest probe whose box contains it, while the background + diffuse keep
// the sky. See metal/probe.rs and metal/probe_set.rs.
pub(super) struct ProbeState {
    // The placements (declared `ReflectionProbe` assets or `auto_seed_probes`),
    // the record of every installed probe (the live count the shaders read) and
    // the staggered bake's queue, which spreads the load cost over several
    // frames instead of one.
    pub book: ProbeBook,
    // One cube per placement; the first `book.count()` hold baked probes.
    pub cubes: super::probe_set::ProbeCubeArray,
    // This frame's copy of the book's records, written by `build_probe_records`.
    // `None` before the first frame builds one.
    pub records_buf: Option<Retained<ProtocolObject<dyn MTLBuffer>>>,
    // At most one probe rendering its six cube faces on the GPU (owning the
    // reserved-ring-slot buffers + capture targets; the faces are submitted without
    // `waitUntilCompleted` and a completion handler flags GPU completion) and one
    // whose capture is convolving into its cube, one destination mip per frame. The
    // convolution overlaps the next probe's render, which shortens the bake
    // warm-up.
    pub bake: super::probe::MtlProbeBake,
    // The three convolution kernels, built at init under the same gate the bake
    // needs (the bindless cull pipeline). `None` disables baking.
    pub prefilter: Option<super::probe_prefilter::ProbePrefilterPipelines>,
    // Deferred-free pool for an in-flight bake's GPU resources when a
    // re-placement (`set_reflection_probes`) interrupts it: the capture command
    // buffers may still be reading those buffers/textures, so they are parked
    // here and freed once the frames-in-flight fence guarantees the bake has
    // retired.
    pub retire_pool: concinnity_core::render::retire_pool::RetirePool<super::probe::RetiredBake>,
}

// Cascaded shadow map resources + the cascade schedule. `enabled` is false when
// `shadow_map_size == 0`, in which case the shadow pass is skipped; `map` and
// `sampler` are always present (a 1x1 fallback every compare reads as fully
// lit when disabled) so fragment shaders can always sample texture(2) /
// sampler(1).
// Mirrors `DxContext::shadow`.
pub(super) struct ShadowState {
    pub enabled: bool,
    // Depth32Float texture array, one slice per cascade.
    pub map: Retained<ProtocolObject<dyn MTLTexture>>,
    // Per-cascade resolution, stored so the shadow pass can size the viewport to
    // the texture array's per-slice dimensions.
    pub map_size: u32,
    pub cadence: backend_init::ShadowCadence,
    // Round-robin clock + primed-set for the cascade schedule; advanced once per
    // frame by `update_shadow_schedule`.
    pub scheduler: shadow_schedule::ShadowCascadeScheduler,
    // Cascades re-rendered this frame (bit `i` = cascade `i`). Computed in
    // draw_frame and read by encode_shadow_pass so the two agree on which slices
    // to refresh and which to leave intact.
    pub render_mask: u32,
    pub sampler: Retained<ProtocolObject<dyn MTLSamplerState>>,
    // Cascaded light VPs + split depths. The near cascade's VP refreshes every
    // frame; far cascades' VPs persist between their round-robin refreshes
    // (Hybrid mode) so each slice is sampled with the VP it was rendered with.
    pub uniforms: ShadowUniforms,
    // World-space unit vector pointing TOWARD the first directional light.
    // Captured at init from the light uniforms; used by per-frame CSM updates.
    pub light_dir: [f32; 3],
}

// Spot-light shadow slices + their refresh schedule. Mirrors
// `DxContext::spot_shadow`.
pub(super) struct SpotShadowState {
    // One Depth32Float slice per shadow-casting spot light, indexed by
    // `GpuLight.shadow_index`. Always present so the fragment shader's
    // depth-array binding is valid; a 1x1 fallback (depth 1.0 = lit) stands in
    // when no spot casts shadows.
    pub map: Retained<ProtocolObject<dyn MTLTexture>>,
    // Per-slice light-space projections, uploaded once (local lights are
    // static). Empty when nothing casts, in which case the pass never runs.
    pub buffer: PooledBuffer,
    pub count: u32,
    // Each slice's light frustum, which its GPU cull keeps casters inside.
    pub frusta: Vec<Frustum>,
    // Prime-then-round-robin refresh schedule over the spot slices, the spot
    // analogue of `ShadowState::scheduler`.
    pub scheduler: spot_shadow::SpotShadowScheduler,
    // Which spot slices re-render this frame; set once per frame in draw_frame
    // and read by encode_spot_shadow_pass.
    pub render_mask: u32,
}

impl SpotShadowState {
    // Slices that re-render this frame.
    pub(super) fn refreshed_slices(&self) -> impl Iterator<Item = u32> {
        spot_shadow::refreshed_slices(self.render_mask, self.count)
    }
}

// Per-frame-in-flight transient buffer rings, plus the CPU scratch that fills
// them. Each ring hands out this frame's slot so a build reuses a buffer instead
// of allocating one per frame; the scratch `Vec`s are `mem::take`n during a build
// and returned after, so the per-frame `collect` reuses one heap allocation.
pub(super) struct FrameRings {
    // Ring of per-frame `GpuObjectData` buffers. Written by `build_object_buffer`.
    pub object: super::frame_rings::TransientRing,
    // Ring of per-frame `GpuDrawArgs` buffers for the GPU-cull pass. Written by
    // `build_draw_args_buffer`.
    pub draw_args: super::frame_rings::TransientRing,
    // Ring of per-frame model-history buffers for the GPU-driven G-buffer /
    // velocity pre-pass: one column-major `float4x4` per cull record, indexed
    // identically to the object buffer. Filled on the GPU by
    // `encode_model_history`; frame `R` reads the slot frame `R - 1` wrote.
    pub model_history: super::frame_rings::TransientRing,
    // Ring of per-frame `BindlessTextures` argument buffers. The argument
    // encoder fills the slot in place each frame; see
    // `build_bindless_texture_args`.
    pub bindless_tex: super::frame_rings::TransientRing,
    // Ring of per-frame probe record buffers, written by
    // `build_probe_records`.
    pub probe_records: super::frame_rings::TransientRing,
    // Ring of per-skinned-object joint-palette buffers, one inner buffer per
    // object. Written by `build_joint_buffers`.
    pub joint: super::frame_rings::JointRing,
    pub object_scratch: Vec<render_types::GpuObjectData>,
    pub draw_args_scratch: Vec<render_types::GpuDrawArgs>,
    // The material parameter table, over the same slots as `object`.
    pub material_params: super::material_params::MaterialParamRing,
}

// Transparent water surfaces and the pipelines that draw them. The RT variants
// trace a sharp reflection ray against the scene BVH instead of sampling the
// probe cube; they are `Some` only when the world has >=1 `WaterSurface` AND the
// device supports ray tracing, and are selected per-frame only while `rt.accel`
// is live. The flat variant uses the per-object material tint as albedo (the
// non-bindless RT fallback); the textured variant samples the reflected hit's
// albedo / normal / emissive maps from the bindless pool (bound at buffer 10 for
// water, since the main pass's index 7 is the ProbeSet here) and is selected
// only in a bindless world.
pub(super) struct WaterState {
    // `Some` only when the world declared >=1 `WaterSurface`; the transparent
    // pass executor short-circuits otherwise.
    pub pipeline: Option<Retained<ProtocolObject<dyn MTLRenderPipelineState>>>,
    pub pipeline_rt: Option<Retained<ProtocolObject<dyn MTLRenderPipelineState>>>,
    pub pipeline_rt_textured: Option<Retained<ProtocolObject<dyn MTLRenderPipelineState>>>,
    // One GPU record per `WaterSurface` asset: tessellated VB+IB plus the
    // per-surface fragment / vertex uniforms, built once at init from the
    // asset.
    pub surfaces: Vec<super::water::WaterSurfaceRecord>,
}

// Transparent glass: the `GlassPanel` producer and the see-through mesh path.
// The panel RT variants follow the same flat / textured split as [`WaterState`].
pub(super) struct GlassState {
    // Shared pipeline for the `GlassPanel` transparent producer. `Some` only
    // when the world declared >=1 `GlassPanel`.
    pub pipeline: Option<Retained<ProtocolObject<dyn MTLRenderPipelineState>>>,
    pub pipeline_rt: Option<super::glass::TracedGlassPipelines>,
    pub pipeline_rt_textured: Option<super::glass::TracedGlassPipelines>,
    // Ray-traced see-through glass MESH pipelines (`glass_mesh.hlsl`): an
    // imported `Material` with `transparent: true` routed through the transparent
    // pass with a per-pixel RT trace off the interpolated mesh normal, instead of
    // the Layer-1 opaque-reflective fallback. Built only on RT-capable devices
    // (regardless of whether the world declares a transparent material -- a live
    // RT toggle then has the pipeline ready). `mesh_pipeline_rt.is_some()` gates
    // the whole transparent-mesh path: when live (RT on) transparent meshes are
    // skipped in the opaque pass + the RT BLAS and drawn here; otherwise they
    // render opaque (Layer 1).
    pub mesh_pipeline_rt: Option<super::glass::TracedGlassPipelines>,
    pub mesh_pipeline_rt_textured: Option<super::glass::TracedGlassPipelines>,
    // Indices into `draw.objects` of every see-through glass mesh (its material
    // has both `transparent` and `see_through` set), precomputed at init so the
    // per-frame Layer 2 producer does not rescan all objects. Empty on non-RT
    // devices or when no material opts into see-through (those transparent meshes
    // then render opaque, Layer 1). The objects stay IN `draw.objects` (a
    // DrawObject's position is a key into the cull / prev-model / RT parallel
    // arrays); this list only marks which to reroute. A non-empty list (plus a
    // built mesh pipeline) is what enables the Layer 2 path -- see
    // `seethrough_meshes_enabled`. See-through is opt-in per `Material` because
    // it only looks right when the space behind the glass is modeled; without
    // it, Layer 1's tinted reflective glass hides the interior.
    pub seethrough_mesh_indices: Vec<usize>,
    // One GPU record per `GlassPanel` asset: the static world-space quad VB+IB
    // plus the per-panel uniforms. Contributes to the transparent pass.
    pub panels: Vec<super::glass::GlassPanelRecord>,
    // The reduced targets the traced glass pre-pass writes, sized at render /
    // the RT trace divisor. `None` while glass traces in place.
    pub reflection_targets: Option<super::glass::GlassReflectionTargets>,
}

// Raymarched SDF volumes and the unit-cube proxy geometry the pass rasterizes.
pub(super) struct RaymarchState {
    // One GPU record per `SdfVolume` asset: the per-volume render pipeline
    // (compiled lazily at init from the user's fragment shader source + the
    // engine-shipped helpers/template) plus the static per-volume uniforms
    // (center, extent, params, ...). Drives the pass at `PassId::Raymarch`.
    pub volumes: Vec<super::raymarch::RaymarchVolumeRecord>,
    // Shared unit-cube proxy geometry (8 vertices, 36 indices). `Some` whenever
    // any `SdfVolume` exists in the world; the encoder reads them per-frame and
    // the asset cost is fixed (96 + 72 bytes).
    pub cube_vertex_buffer: Option<Retained<ProtocolObject<dyn MTLBuffer>>>,
    pub cube_index_buffer: Option<Retained<ProtocolObject<dyn MTLBuffer>>>,
}

// HUD text pass. `pipeline_state` is `None` when no Font assets were declared
// and `atlas_textures` is empty in the same case. The pipeline targets the
// single-sample drawable in the composite pass (after HDR tonemap), so text is
// rendered in display-referred LDR.
pub(super) struct TextState {
    pub pipeline_state: Option<Retained<ProtocolObject<dyn MTLRenderPipelineState>>>,
    pub atlas_textures: Vec<PooledTexture>,
    // Linear-clamp sampler for glyph atlas lookups.
    pub sampler: Retained<ProtocolObject<dyn MTLSamplerState>>,
    // Ring of per-frame HUD text geometry buffers. `draw_frame` writes the whole
    // frame's labels into this frame's slot before the graph runs; the composite
    // pass binds sub-ranges of it.
    pub upload: super::text_upload::TextUploadRing,
}

// Built-in shader hot reload. `enabled` is true only under `cn debug`: it
// switches the built-in shader source loader to a disk-first read with
// embedded fallback, so an edited shader is picked up by
// [`MtlContext::reload_shaders`]. Under `cn run` production keeps the embedded
// path. `reload_pending` is the atomic flag the `reload-shaders` debug tool
// call sets, polled at the top of `draw_frame`; the debug server reads its
// `Arc` clone via `GraphicsSystem`. `Some` only when `enabled`.
pub(super) struct HotReloadState {
    pub enabled: bool,
    pub reload_pending: Option<std::sync::Arc<std::sync::atomic::AtomicBool>>,
    // Bumped by every engine-template reload, so a world Shader pipeline a
    // worker built from the templates before it is never installed after it.
    pub generation: u64,
}

impl HotReloadState {
    pub(super) fn new(enabled: bool) -> Self {
        Self {
            enabled,
            reload_pending: enabled
                .then(|| std::sync::Arc::new(std::sync::atomic::AtomicBool::new(false))),
            generation: 0,
        }
    }
}

// The window this context draws into.
pub(super) struct WindowState {
    // The shared AppKit window + input layer (crate::appkit), which also owns the
    // NSWindow, the fullscreen tracking, and the display-mode hold. Holds the
    // same view as `view` below, upcast to `NSView`.
    pub appkit: crate::appkit::AppKitWindow,
    // MTKView with isPaused=true and enableSetNeedsDisplay=false so its internal
    // display link never fires. draw() is called manually from draw_frame().
    // Metal-only: the drawable and its render-pass descriptor come from here,
    // which is why the shared window layer keeps only the NSView upcast.
    pub view: Retained<MTKView>,
}

impl Drop for WindowState {
    fn drop(&mut self) {
        // Always release the cursor on teardown so the OS mouse association and
        // cursor visibility are restored even if the caller didn't do it.
        self.appkit.release_cursor();
        if let Some(window) = self.appkit.window() {
            // Close the window so it doesn't linger after the run loop exits.
            window.close();
        } else {
            // In embedded mode (no NSWindow), the MTKView was added as a subview.
            // Explicitly remove it so it doesn't outlive the preview session.
            self.view.removeFromSuperview();
        }
    }
}

// Per-frame counters, GPU timings, and the fault reporting that crosses the
// backend boundary. Most of these are written from a GPU completion handler on a
// callback thread, so they are shared behind atomics.
pub(super) struct Diagnostics {
    // Per-frame draw-call / VRAM / GPU-time counters surfaced to the profiler
    // overlay via `render_stats`.
    pub frame_stats: profile::RenderStats,
    // GPU execution time of the last completed frame, in microseconds. Written
    // by each command buffer's completion handler.
    pub gpu_time_us: std::sync::Arc<std::sync::atomic::AtomicU32>,
    // First classified GPU failure observed on a completed frame command buffer,
    // parked here by the completion handler until the next draw_frame reports it
    // across the backend boundary.
    pub device_error: std::sync::Arc<std::sync::Mutex<Option<error::RenderError>>>,
    // Per-pass GPU sample buffers, when the active device supports the
    // `MTLCommonCounterSetTimestamp` counter set. Each `draw_frame` rotates to
    // the next slot in the ring; the completion handler resolves that slot into
    // `diagnostics.pass_times_us` for the profiler overlay.
    pub pass_timing: Option<super::pass_timing::PassTimingResources>,
    // Per-pass GPU microseconds from the most recently resolved frame. One
    // atomic per pass slot, shared between the GPU completion handler that
    // writes it and `render_stats()` that reads it.
    pub pass_times_us:
        std::sync::Arc<[std::sync::atomic::AtomicU32; super::pass_timing::PASS_COUNT]>,
    // Atomic accumulator the parallel-dispatched workers fetch_add their draw
    // counts into. Drained into `diagnostics.frame_stats.draw_calls` at the end of
    // `execute_graph`. AtomicU32 because workers may run concurrently.
    pub draw_calls_accum: std::sync::atomic::AtomicU32,
}

// The device layer every per-world resource is built on: device, queues,
// allocator and window, plus the display settings negotiated with them. The
// window is declared last so everything that presents to it releases first.
pub(super) struct MtlHardware {
    pub device: Retained<ProtocolObject<dyn objc2_metal::MTLDevice>>,
    // Block pool every persistent CPU-written, GPU-read-only buffer and texture
    // is placed in, so the world's resource count costs a handful of heaps
    // rather than one device allocation each. See `metal/allocator.rs`.
    pub allocator: DeviceAllocator,
    pub command_queue: Retained<ProtocolObject<dyn MTLCommandQueue>>,
    // The async-compute queue plus the per-queue `MTLEvent`s the render-graph
    // executor submits the schedule's two queues over. `None` when the device
    // would not create them, which drops the executor back to recording every
    // pass onto `command_queue` in compiled order.
    pub graph_queues: Option<super::graph_queues::GraphQueues>,
    // Pixel format the MTKView's CAMetalLayer is currently presenting at:
    // `BGRA8Unorm` for SDR, `RGBA16Float` for HDR EDR. The post + text
    // pipelines bake this format into their color attachment descriptors,
    // so it has to be stable between `MtlContext::new` and any subsequent
    // hot-reload rebuild. The per-frame `PostProcessParams.hdr_output` flag
    // carries the EDR signal into the shader.
    pub swap_pixel_format: MTLPixelFormat,
    // The resolved swapchain color-output mode. EDR negotiation can fall the
    // world's HDR request back to SDR, so this can disagree with
    // `swapchain_config`.
    pub hdr_mode: hdr_output::HdrOutputMode,
    // The swapchain config (ring depth + HDR request) this context was built
    // with, reported by `hot_swap_config` so a live editor reload reuses this
    // hardware only when the new world's config still matches. Holding the
    // request rather than the negotiated mode keeps a display without EDR
    // headroom from forcing a rebuild every save.
    pub swapchain_config: backend_init::SwapchainConfig,
    // `Option` so a live `cn editor` world reload can MOVE the window (and its
    // live cursor / keymap / fullscreen state) into the rebuilt context; a
    // normal context always holds `Some`. Access via `window`/`window_mut`.
    pub window: Option<WindowState>,
}

impl MtlHardware {
    // The hardware an outgoing context hands its successor on a live editor
    // `reload_world`: clones of the device and command queue, with the window
    // and the graph queues moved out. The successor places into a fresh
    // allocator; the outgoing world releases into its own.
    pub(super) fn hand_over(&mut self) -> error::RenderResult<Self> {
        Ok(Self {
            window: Some(self.window.take().ok_or_else(|| {
                error::RenderError::Other("apply_world_reload: window already taken".to_string())
            })?),
            graph_queues: self.graph_queues.take(),
            allocator: DeviceAllocator::new(&self.device, self.swapchain_config.frames_in_flight),
            device: self.device.clone(),
            command_queue: self.command_queue.clone(),
            swap_pixel_format: self.swap_pixel_format,
            hdr_mode: self.hdr_mode,
            swapchain_config: self.swapchain_config,
        })
    }

    // Maximum extended-range multiplier on the HDR path, `None` on SDR.
    pub(super) fn max_edr(&self) -> Option<f32> {
        match self.hdr_mode {
            hdr_output::HdrOutputMode::Hdr { max_edr, .. } => Some(max_edr),
            hdr_output::HdrOutputMode::Sdr => None,
        }
    }

    // HDR encoding (scRGB-linear vs PQ) on the HDR path, `None` on SDR.
    pub(super) fn hdr_encoding(&self) -> Option<hdr_output::HdrEncoding> {
        match self.hdr_mode {
            hdr_output::HdrOutputMode::Hdr { encoding, .. } => Some(encoding),
            hdr_output::HdrOutputMode::Sdr => None,
        }
    }
}

// Render-resolution scene targets, rebuilt on resize, plus the depth tests the
// scene passes run against them.
pub(super) struct MtlTargets {
    // Off-screen HDR render targets (MSAA RGBA16Float + resolve + MSAA
    // Depth32Float). Re-created lazily in `draw_frame` whenever the
    // drawable size changes. The main + instanced pipelines render into
    // these, and the post-process pass samples `hdr_resolve` to write
    // the tonemapped + FXAA-filtered output into the drawable.
    pub hdr: HdrTargets,
    // The drawable (output) size the bloom chain and the pool's half-output
    // `bloom_top` were built for. Bloom follows it rather than the render
    // resolution, so the glow stays on the panel's pixel grid under upscaling.
    pub output: (u32, u32),
    // Pool backing the render graph's transient textures
    // (`gfx::render_graph::alias`). Owns `bloom_top` (the bloom chain's top
    // octave) and, when SSAO is on, `ao_output`; their disjoint lifetimes put
    // them on one aliased `MTLHeap` slot. Rebuilt on resize. See
    // [`TransientTexturePool`].
    pub transient_pool: TransientTexturePool,
    // The opaque-geometry test: nearer fragments pass and write. The shadow
    // casters draw with it too: the cascades, the spot slices and the
    // raymarched casters drawn into them.
    pub depth_state: Retained<ProtocolObject<dyn MTLDepthStencilState>>,
    // The test that also passes at equal depth, with write. Opaque
    // raymarched surfaces bind it: their shader writes a depth no farther than
    // the proxy's rasterized one.
    pub depth_state_inclusive: Retained<ProtocolObject<dyn MTLDepthStencilState>>,
    // Read-only depth state: the inclusive test, no write. Used by translucent
    // draws that must be occluded by nearer opaque geometry but must not
    // update the depth buffer (volumetric raymarch volumes). Metal forbids
    // `setDepthStencilState(nil)` under the validation layer, so translucent
    // passes bind this instead of clearing the state.
    pub depth_state_read_only: Retained<ProtocolObject<dyn MTLDepthStencilState>>,
    // True when the world has no 3D geometry (e.g. a text-only world). The
    // off-screen HDR / bloom / effect targets are then allocated at 1x1 since
    // nothing is rendered into them; the composite pass still runs at the full
    // drawable size, so text stays crisp.
    pub geometry_less: bool,
}

// The world's scene assets: the shared geometry buffers, the texture pool, the
// light tables, the IBL cubes and color-grading LUT, and the samplers they are
// read through. None of it depends on the drawable.
pub(super) struct MtlSceneAssets {
    // Single shared vertex buffer containing geometry for all draw objects.
    pub vertex_buffer: PooledBuffer,
    // Single shared index buffer containing indices for all draw objects.
    pub index_buffer: PooledBuffer,
    // Shared texture pool (slot == handle): every texture (albedo, normal map,
    // emissive/ORM, terrain secondary) lives here once, matching DX/VK. A 1x1
    // opaque-white fallback is always present at slot 0 so shaders that sample
    // texture(0) produce correct output even when no Texture asset was declared.
    pub textures: Vec<PooledTexture>,
    // Holds only the 1x1 flat-normal fallback (RGBA 128,128,255,255) a
    // normal-less draw samples; its pool slot is one past the last real texture.
    // Real normal maps are entries in `textures`.
    pub fallback_textures: Vec<PooledTexture>,
    // Per-scene local lights (point + spot + area) for the clustered forward
    // pass, bound at fragment buffer(8). Always holds at least one element (a
    // neutral placeholder when the scene declares no local lights) so the
    // binding is valid; `light_uniforms.num_local_lights` bounds iteration.
    pub local_light_buffer: PooledBuffer,
    // The local lights' reach, which places the far end of the cluster grid.
    pub cluster_reach: ClusterReach,
    // Per-scene rect area-light extents, indexed by `GpuLight.data_index`.
    // Uploaded once; a one-element placeholder when the world declares none.
    pub area_light_buffer: PooledBuffer,
    // The two linearly-transformed-cosine lookup tables the area-light shading
    // path samples: the inverse transforms (RGBA32Float) and the matching
    // (albedo, Fresnel) pairs (RG32Float). Generated at build time, so these are
    // scene-independent and created once.
    pub ltc_matrix_texture: PooledTexture,
    pub ltc_magnitude_texture: PooledTexture,
    // IBL cubemaps + mip count. Always Some: the runtime synthesizes a 1x1
    // gray fallback for both cubes when no EnvironmentMap was supplied, so
    // the fragment shader's texture(3) / texture(4) bindings are always
    // valid. `prefilter_mip_count == 0` means no EnvironmentMap is declared,
    // and the fragment shader uses the flat albedo ambient term instead of IBL.
    pub env_map: EnvironmentMapTextures,
    // 3D color-grading LUT sampled in the composite pass. Holds the declared
    // `ColorLut` payload, or a 2x2x2 identity LUT when the world declares
    // none, so the composite pass binds a valid 3D texture either way.
    pub color_lut: PooledTexture,
    // Linear-filter, repeat-wrap sampler the texture pool is read through.
    pub sampler: Retained<ProtocolObject<dyn MTLSamplerState>>,
    // Linear-clamp sampler bound at sampler(2) for cubemap sampling.
    pub cube_sampler: Retained<ProtocolObject<dyn MTLSamplerState>>,
}

// The bindless main pass's argument buffers, plus what decides when a ring
// slot is rewritten: the slot gates, the residency set and the texture epoch.
// See `metal/bindless_args.rs`. The texture block is written as plain resource
// ids (argument-buffer tier 2), so it needs no encoder; only the sampler block,
// written once, goes through one.
pub(super) struct MtlArgumentBuffers {
    // Per-ring-slot change gates for the texture block: the block's contents
    // and its unused tail, gated apart because a stream-in moves the former
    // every time and the latter almost never. See `build_bindless_texture_args`.
    pub bindless_tex_gates: super::bindless_args::SlotGates,
    pub bindless_tail_gates: super::bindless_args::SlotGates,
    // Every texture the `BindlessTextures` block names, declared resident in
    // one batched call per encoder. Refreshed once per frame by
    // `refresh_bindless_residency`.
    pub bindless_residency: super::bindless_args::ResidencySet,
    // The engine sampler block bound at fragment buffer(10) for the
    // single-source main program: three static samplers written once at init
    // (samplers never stream, so no per-frame ring is needed). `None` when the
    // bindless cull pipeline is off.
    pub bindless_sampler_args: Option<Retained<ProtocolObject<dyn MTLBuffer>>>,
    // Bumped by every announced change to the textures those argument buffers
    // name (a pool slot streamed in or evicted, an env-map or probe swap), so a
    // handle that happens to reuse a freed one's address still moves the
    // signature the gates compare.
    pub texture_epoch: u64,
}

// The composite pass: ACES tonemap + gamma 2.2 + FXAA (SDR drawable) or linear
// extended-range output (HDR drawable) from the resolved HDR target.
pub(super) struct CompositeState {
    pub pipeline: Retained<ProtocolObject<dyn MTLRenderPipelineState>>,
    // Linear-clamp sampler bound at sampler(0) when sampling the HDR resolve
    // target. Also reused by the bloom passes and the post-process chain.
    pub sampler: Retained<ProtocolObject<dyn MTLSamplerState>>,
}

// Metal rendering context. Owns all GPU resources and the window.
// Only ever accessed from the main thread.
pub(crate) struct MtlContext {
    // Color texture of the most recently presented drawable, retained so the
    // `cn debug` `screenshot` command can blit it back to a host buffer and
    // PNG-encode it (see metal/screenshot.rs). Set each frame only under
    // `hot_reload` (the path that runs the debug endpoint able to request a
    // capture); `None` in production and before the first present, so a
    // capture then returns a clean error. The MTKView has `framebufferOnly`
    // switched off under the same gate so the drawable is blit-readable.
    // Mirrors the DX/VK `last_present_index`.
    pub(super) last_present_texture: Option<Retained<ProtocolObject<dyn MTLTexture>>>,
    // GPU-driven main pass + cull feature state: the main and world-shader
    // pipelines, 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 pyramid + view-projection
    // snapshots the occlusion test reprojects through. All `Some`/active only
    // when the world has 3D scene content. See [`CullState`].
    pub(super) cull: CullState,
    // Argument encoders and their re-encode gates. See [`MtlArgumentBuffers`].
    pub(super) arg_buffers: MtlArgumentBuffers,
    // The CPU-side scene: draw list, view, skinned slots, model history and
    // streamed-geometry placement. See [`SceneState`].
    pub(super) state: SceneState,
    // The frame graph last compiled, keyed by the `FrameGraphInputs` it was
    // built from. `build_frame_graph` is a pure function of those inputs (which
    // change only when a feature toggles or a target resizes), so a frame whose
    // inputs match the cached key reuses it instead of rebuilding. Taken out
    // during `execute_graph` (which needs `&mut self`) and put back after, so a
    // steady scene compiles once.
    pub(super) graph_cache: Option<(render_graph::FrameGraphInputs, render_graph::CompiledGraph)>,
    // InstancedProp clusters. See [`InstancedState`].
    pub(super) instanced: InstancedState,
    // Scene assets. See [`MtlSceneAssets`].
    pub(super) scene: MtlSceneAssets,
    // All scene lights packed and pushed to the fragment shader at buffer(4).
    pub(super) light_uniforms: LightUniforms,
    // Cascaded shadow map + its schedule. See [`ShadowState`].
    pub(super) shadow: ShadowState,
    // Spot-light shadow slices + their schedule. See [`SpotShadowState`].
    pub(super) spot_shadow: SpotShadowState,
    // Local reflection probes: the scene captured into one cube per placement
    // (metal/probe.rs). Distinct from `scene.env_map` (which stays the sky -- it
    // drives the background + diffuse irradiance) so the bake never corrupts the
    // visible sky. See [`ProbeState`].
    pub(super) probe: ProbeState,
    pub(super) text: TextState,
    // Render-resolution scene targets. See [`MtlTargets`].
    pub(super) targets: MtlTargets,
    // Composite pass. See [`CompositeState`].
    pub(super) composite: CompositeState,
    // The bloom chain. None for a world with no 3D scene content: the graph
    // never inserts the Bloom pass, and the composite's unconditional bind of
    // the pool's `bloom_top` reads it unwritten.
    pub(super) bloom: Option<MtlBloomPass>,
    // Post-process tunables (bloom intensity / threshold / knee). Pushed to
    // the bloom prefilter and composite fragment shaders. `bloom_intensity`
    // of 0 skips the bloom passes entirely.
    pub(super) post_process: render_types::PostProcessParams,
    // Temporal-anti-aliasing feature state: the toggle, resolve pipeline,
    // ping-pong history buffers, and per-frame bookkeeping. See [`TaaState`].
    pub(super) taa: TaaState,
    // Previous frame's un-jittered view-projection, fed to the velocity
    // pre-pass to reproject motion. (Shared by both TAA and the upscaler when
    // either drives the velocity pre-pass, so it is kept flat rather than
    // under `taa`.)
    pub(super) view_history: ViewHistory,
    // MetalFX-temporal-upscaling feature state: the scaler, the input/output
    // scale ratio, the per-frame projection jitter, and the history-reset
    // flag. When the scaler is `Some`, the 3D scene renders at
    // `(output * upscale.scale)` and the scaler reconstructs a
    // drawable-resolution image the bloom + composite stack reads as
    // `scene_color`; the TAA pass is bypassed (the scaler accumulates
    // temporally), though the velocity pre-pass + projection jitter still
    // run. See [`UpscaleState`].
    pub(super) upscale: UpscaleState,
    // SSAO (GTAO) feature state: resolved tunables, occlusion targets, the
    // kernel + blur pipelines, and the 1×1 white fallback. See [`SsaoState`].
    pub(super) ssao: SsaoState,
    // Screen-space-reflection feature state: resolved tunables, the resolve
    // output target (shared with SSGI/RT), and the resolve pipeline. See
    // [`SsrState`].
    pub(super) ssr: SsrState,
    // Unified G-buffer pre-pass feature state: the shared normal+depth /
    // roughness / velocity / sampleable-depth targets plus the static /
    // instanced / skinned pipelines. See [`GBufferState`].
    pub(super) gbuffer: GBufferState,
    // Screen-space-GI feature state: resolved tunables and the shared pass
    // with its pyramids and accumulation. See [`SsgiState`].
    pub(super) ssgi: SsgiState,
    // Hardware-ray-traced-reflection feature state: resolved tunables, the
    // scene acceleration structure, the dynamic-update mode + failure flag, and
    // the resolve / textured-resolve / skinning pipelines. See [`RtState`].
    pub(super) rt: RtState,
    // World-space line state (editor origin axes, collider previews):
    // the pipeline, built on the first frame that publishes lines. See
    // [`LineState`].
    pub(super) lines: LineState,
    // Projected-decal feature state: the decal slot table, the pipeline, the
    // shared unit-cube geometry, and the sampler. See [`DecalState`]. The pipeline / cube buffers / sampler are
    // built lazily at init (≥1 declared decal) or on the first runtime
    // [`MtlContext::add_decal`].
    pub(super) decal: DecalState,
    // Volumetric-fog feature state: resolved tunables, the ray-march
    // pipeline, and the froxel-volume compute pipeline + 3D output volume.
    // See [`FogState`].
    pub(super) fog: FogState,
    // Clustered-lighting state: the binning compute pipeline + the per-cluster
    // light-index buffer. See [`LightCullState`].
    pub(super) light_cull: super::light_cull::LightCullState,
    // Per-frame clustered-lighting params (main camera), rebuilt each frame in
    // draw_frame and bound to the light-cull pass + the forward pass. Mirrors
    // shadow.uniforms' per-frame update + shared bind.
    pub(super) cluster_params: ClusterParams,
    // Particle-system feature state: the per-emitter records (+ tombstone
    // free-list), the parallel per-emitter GPU pools, the shared compute +
    // render pipelines, and the per-frame timing bookkeeping. See
    // [`ParticleState`].
    pub(super) particle: ParticleState,
    // Auto-exposure feature state: resolved tunables, the EMA-tracked adapted
    // EV + authored bias, the histogram/average compute pipelines + buffers,
    // and the per-frame timing bookkeeping. See [`AutoExposureGpu`].
    pub(super) auto_exposure: AutoExposureGpu,
    pub(super) hot_reload: HotReloadState,
    // The world default Shader's compiled programs, `None` for the engine's
    // own. Kept past init for the skinned pipeline built at upload and for the
    // hot-reload rebuild.
    pub(super) world_shader: Option<concinnity_core::components::ShaderPrograms>,
    // Keep each presented drawable's texture retained so `screenshot` can
    // blit it back (the view is blit-readable under the same flag). On under
    // the dev loop and `cn run --screenshot`; false in plain production.
    pub(super) capture: bool,
    // Skinned-mesh rendering feature state: the shared skinned vertex / index
    // buffers, the pre-skin pipeline and its deformed buffers, and the morph
    // bindings. See [`SkinnedState`].
    pub(super) skinned: SkinnedState,
    pub(super) diagnostics: Diagnostics,
    // Frames-in-flight pacing. `draw_frame` acquires a slot before encoding
    // and the frame command buffer's completion handler releases it once the
    // GPU retires the frame, bounding how far the CPU may queue ahead of the
    // GPU (and thus how many sets of per-frame transient buffers can pile up).
    pub(super) frame_pacing: super::frame_pacing::FrameInFlight,
    // Ring depth for the per-frame transient buffers below: equal to the
    // frames-in-flight count (≥1). The fence guarantees frame `R − depth` has
    // retired before frame `R` reuses ring slot `R % depth`, so overwriting
    // that slot's buffer never races an in-flight GPU read.
    pub(super) frames_in_flight: usize,
    // Monotonic counter over frames that build per-frame buffers; `% depth`
    // selects this frame's ring slot. Advanced once per such frame.
    pub(super) frame_ring_index: u64,
    pub(super) rings: FrameRings,
    pub(super) water: WaterState,
    // Planar reflection targets, one set per distinct reflector plane (water
    // surfaces + glass panes, grouped by `assign_planar_slots`). `Some` only when
    // the world declared >=1 such reflector; the scene is re-rendered mirrored
    // across each plane into these each frame (RT off) and the reflective shader
    // samples the resolve of its slot. Rebuilt on resize alongside `targets.hdr`.
    pub(super) planar_reflection: Option<super::planar::PlanarReflectionSet>,
    pub(super) glass: GlassState,
    pub(super) raymarch: RaymarchState,
    // The environment drawn behind the opaque scene. See `metal/sky.rs`.
    pub(super) sky: super::sky::SkyState,
    // Device, queues, allocator and window. Declared last so every resource
    // above releases before the device layer and the window it presents to.
    pub(super) hw: MtlHardware,
}

// SAFETY: MtlContext is only ever accessed from the main thread (as documented
// on the struct). The Retained<ProtocolObject<...>> Metal handles aren't Send by
// default, but `RenderBackend: Send` requires it so GraphicsSystem can box
// the backend behind a trait object. Mirrors DxContext / VkContext.
unsafe impl Send for MtlContext {}

// Debug-only guard that the caller is on the main thread.
//
// The `unsafe impl Send for MtlContext` above is sound only because the
// context is touched from the main thread alone: AppKit and Metal command
// submission are both main-thread-affine. `draw_frame` proves this with a
// `MainThreadMarker`, but the `RenderBackend` mutation entry points (reached
// through the boxed trait object) did not, so scheduling `GraphicsSystem`
// off the main thread would silently race AppKit/Metal instead of failing.
// This makes that mistake panic loudly in debug builds and compiles to
// nothing in release. `entry` is the offending method name, for the message.
#[inline]
#[track_caller]
pub(super) fn debug_assert_main_thread(entry: &str) {
    debug_assert!(
        objc2::MainThreadMarker::new().is_some(),
        "{entry} must be called from the main thread: MtlContext is main-thread-only \
         (see `unsafe impl Send for MtlContext`); driving GraphicsSystem off the main \
         thread races AppKit/Metal",
    );
}

// The cull-record set one GPU-driven encode covers: `total` command slots, of
// which the folded skinned tail starts at `skinned_base`. Cull dispatches and
// indirect-draw ranges take this instead of reading the live draw list, so an
// encode whose object + draw-args buffers were snapshotted on an earlier frame
// keeps addressing exactly the records those buffers hold. The reflection-probe
// bake is the one such consumer: it builds its buffers once and renders the six
// cube faces one per frame, while runtime spawns keep growing the live list.
#[derive(Clone, Copy, Debug, PartialEq, Eq)]
pub(super) struct DrawRecordCounts {
    pub total: usize,
    pub skinned_base: usize,
}

impl DrawRecordCounts {
    // Command slots holding the static objects + folded instances, which draw
    // through the static index buffer. `block` is where the record set starts in
    // the target ICB: 0 for the main ICB, the cascade's block for the shadow ICB.
    // `None` when the set has no such records.
    pub(super) fn prefix(&self, block: usize) -> Option<std::ops::Range<usize>> {
        (self.skinned_base > 0).then(|| block..block + self.skinned_base)
    }

    // Command slots holding the folded skinned tail, which draws compute-deformed
    // vertices through the skinned index buffer. `None` when nothing is folded.
    pub(super) fn skinned_tail(&self, block: usize) -> Option<std::ops::Range<usize>> {
        (self.total > self.skinned_base).then(|| block + self.skinned_base..block + self.total)
    }
}

// The `NSRange` an indirect-draw `executeCommandsInBuffer` takes for `r`.
pub(super) fn ns_range(r: std::ops::Range<usize>) -> objc2_foundation::NSRange {
    objc2_foundation::NSRange {
        location: r.start,
        length: r.end - r.start,
    }
}

impl MtlContext {
    // The live window. `None` only on the outgoing context of a `reload_world`
    // (its window was moved into the successor), which is dropped without any
    // further window access.
    #[inline]
    pub(super) fn window(&self) -> &WindowState {
        self.hw
            .window
            .as_ref()
            .expect("MtlContext window taken by reload_world")
    }

    #[inline]
    pub(super) fn window_mut(&mut self) -> &mut WindowState {
        self.hw
            .window
            .as_mut()
            .expect("MtlContext window taken by reload_world")
    }

    // The record set the live draw list covers. Every frame-driven pass takes
    // this; only the reflection-probe bake substitutes its own snapshot.
    pub(super) fn draw_record_counts(&self) -> DrawRecordCounts {
        DrawRecordCounts {
            total: self.cull_count(),
            skinned_base: self.skinned_record_base(),
        }
    }

    // Number of records the GPU-driven cull processes this frame: the static
    // draw objects, every folded instance, then every folded skinned object.
    // Metal has no separate `n_objects` field; `draw.objects.len()` is the
    // static count. Drives the cull dispatch width + `object_count` uniform, the
    // shared ICB capacity, and the indirect-draw `NSRange`. Equals
    // `draw.objects.len()` for static-only worlds, so those paths are untouched.
    pub(super) fn cull_count(&self) -> usize {
        self.state.draw.objects.len() + self.state.draw.n_instances + self.state.draw.n_skinned
    }

    // Index in the unified cull list where the folded skinned records begin
    // (static objects + instances precede them). The cull kernel draws records
    // at or past this through the skinned index buffer (the skinned tail),
    // and the main pass binds the deformed vertex buffer for that range. Equals
    // `cull_count()` when no skinned mesh is folded.
    pub(super) fn skinned_record_base(&self) -> usize {
        self.state.draw.objects.len() + self.state.draw.n_instances
    }

    // The buffer to bind at the cull kernel's skinned-index slot (buffer 6):
    // the skinned index buffer when a SkinnedMesh has uploaded, else the
    // static index buffer as a harmless placeholder. The kernel only reads it
    // for records at/after `skinned_base`, which equals `cull_count()` (so the
    // skinned branch never fires) whenever the skinned buffer is absent; Metal
    // still requires a referenced buffer to be bound, hence the placeholder.
    pub(super) fn skinned_index_or_placeholder(
        &self,
    ) -> &ProtocolObject<dyn objc2_metal::MTLBuffer> {
        match self.skinned.index_buffer.as_ref() {
            Some(b) => b.as_ref(),
            None => self.scene.index_buffer.as_ref(),
        }
    }

    // Ensure the GPU-driven cull indirect command buffer has a command slot
    // for every one of `count` draw objects, rebuilding it (and re-encoding
    // its argument buffer) when the draw list has outgrown it. A no-op for
    // non-bindless contexts, which have no cull pipeline. New capacity is
    // rounded up to the next power of two so streamed chunks growing
    // `draw.objects` do not rebuild the ICB every frame.
    //
    // The per-object `cull_status_buffer` (always, for the phase-1 kernel's
    // buffer(5) binding) and (under two-pass occlusion) the phase-2 ICB
    // `cull_icb_2` + its argument buffer are grown in lockstep on the same
    // trigger, so all three stay sized to the live draw-object count.
    pub(super) fn ensure_icb_capacity(&mut self, count: usize) -> error::RenderResult<()> {
        // Retained is reference-counted; cloning the handle lets the rest of
        // the method mutate `self` without holding a borrow on the encoder.
        let arg_encoder = match &self.cull.icb_arg_encoder {
            Some(e) => e.clone(),
            None => return Ok(()),
        };
        if !self.cull.icbs.is_empty() && count <= self.cull.icb_capacity {
            return Ok(());
        }
        let new_cap = count.next_power_of_two().max(64);
        let bucket_count = self.cull.bucket_count.max(1);

        let mut icbs = Vec::with_capacity(bucket_count);
        for _ in 0..bucket_count {
            icbs.push(self.build_cull_icb(new_cap)?);
        }

        // The argument buffer is a fixed-size handle to the ICB array; create
        // it once, then re-encode it to point at each freshly built set.
        if self.cull.icb_arg_buffer.is_none() {
            let len = arg_encoder.encodedLength().max(16);
            let buf = self
                .hw
                .device
                .newBufferWithLength_options(len, MTLResourceOptions::StorageModeShared)
                .ok_or_else(|| super::error::allocation_failed("ICB argument buffer"))?;
            self.cull.icb_arg_buffer = Some(buf);
        }
        let arg_buf = self
            .cull
            .icb_arg_buffer
            .as_ref()
            .expect("ICB argument buffer was just ensured");
        // SAFETY: the argument buffer was sized to `encodedLength()`, and each
        // bucket's ICB is encoded at its bucket index within the `icbs`
        // `[[id(0)]]` array of the kernel's `ICBContainer` struct. Entries
        // past `bucket_count` stay null; the kernel never touches them.
        unsafe {
            arg_encoder.setArgumentBuffer_offset(Some(arg_buf), 0);
            for (b, icb) in icbs.iter().enumerate() {
                arg_encoder.setIndirectCommandBuffer_atIndex(Some(icb), b);
            }
        }

        // Per-object status buffer: one u32 per command slot, private storage
        // (GPU-written by phase-1 cull, GPU-read by phase-2 cull). Always
        // allocated so the phase-1 kernel's buffer(5) binding always resolves.
        let status = self
            .hw
            .device
            .newBufferWithLength_options(
                new_cap * std::mem::size_of::<u32>(),
                MTLResourceOptions::StorageModePrivate,
            )
            .ok_or_else(|| super::error::allocation_failed("cull status buffer"))?;
        self.cull.status_buffer = Some(status);

        // Second-pass ICB + argument buffer, only when two-pass occlusion is on.
        if self.cull.two_pass_occlusion {
            let mut icbs_2 = Vec::with_capacity(bucket_count);
            for _ in 0..bucket_count {
                icbs_2.push(self.build_cull_icb(new_cap)?);
            }
            if self.cull.icb_2_arg_buffer.is_none() {
                let len = arg_encoder.encodedLength().max(16);
                let buf = self
                    .hw
                    .device
                    .newBufferWithLength_options(len, MTLResourceOptions::StorageModeShared)
                    .ok_or_else(|| {
                        super::error::allocation_failed("phase-2 ICB argument buffer")
                    })?;
                self.cull.icb_2_arg_buffer = Some(buf);
            }
            let arg_buf2 = self
                .cull
                .icb_2_arg_buffer
                .as_ref()
                .expect("phase-2 ICB argument buffer was just ensured");
            // SAFETY: same layout as the phase-1 set: each bucket's ICB at its
            // index within the `icbs` `[[id(0)]]` array, and the argument
            // buffer was sized to `encodedLength()`.
            unsafe {
                arg_encoder.setArgumentBuffer_offset(Some(arg_buf2), 0);
                for (b, icb) in icbs_2.iter().enumerate() {
                    arg_encoder.setIndirectCommandBuffer_atIndex(Some(icb), b);
                }
            }
            self.cull.icbs_2 = icbs_2;
        }

        self.cull.icbs = icbs;
        self.cull.icb_capacity = new_cap;
        Ok(())
    }

    // Ensure the GPU-driven shadow view ICBs and their status buffers have a
    // slot for every view of every record: `views * count` per set (view `v`'s
    // live at `[v*count, (v+1)*count)`, the same stride `encode_shadow_culls`
    // writes at and the shadow render passes execute), for the cascades and
    // for the spot slices. A no-op for non-bindless / no-shadow contexts (no
    // shadow decision pipeline). Rounded to the next power of two so a streamed
    // chunk growing `cull_count()` does not rebuild an ICB every frame. Called
    // from `draw_frame` (where `&mut self` is available) right after
    // `ensure_icb_capacity`, so the encode pass only ever reads sized ICBs.
    pub(super) fn ensure_shadow_icb_capacity(&mut self, count: usize) -> error::RenderResult<()> {
        if self.cull.shadow_pipeline.is_none() {
            return Ok(());
        }
        let mut cascades = std::mem::take(&mut self.cull.shadow_views);
        let grown = self.grow_shadow_view_icb(&mut cascades, NUM_SHADOW_CASCADES, count);
        self.cull.shadow_views = cascades;
        grown?;
        let mut spots = std::mem::take(&mut self.cull.spot_views);
        let grown = self.grow_shadow_view_icb(&mut spots, self.spot_shadow.count as usize, count);
        self.cull.spot_views = spots;
        grown
    }

    // Grow one view set's ICB, argument buffer and status buffer to hold
    // `views * count` command slots. A set with no view stays empty.
    fn grow_shadow_view_icb(
        &self,
        set: &mut crate::metal::cull::ShadowViewIcb,
        views: usize,
        count: usize,
    ) -> error::RenderResult<()> {
        let Some(arg_encoder) = &self.cull.icb_arg_encoder else {
            return Ok(());
        };
        let needed = count.saturating_mul(views);
        if needed == 0 || (set.icb.is_some() && needed <= set.capacity) {
            return Ok(());
        }
        let new_cap = needed.next_power_of_two().max(64);
        let icb = self.build_cull_icb(new_cap)?;
        let arg_buf = match set.arg_buffer.take() {
            Some(buf) => buf,
            None => {
                let len = arg_encoder.encodedLength().max(16);
                self.hw
                    .device
                    .newBufferWithLength_options(len, MTLResourceOptions::StorageModeShared)
                    .ok_or_else(|| {
                        super::error::allocation_failed("shadow view ICB argument buffer")
                    })?
            }
        };
        // SAFETY: the argument buffer was sized to `encodedLength()`, and the ICB
        // is encoded at slot 0 (the single `[[id(0)]]` member of the kernel's
        // `ICBContainer` argument-buffer struct), exactly like the main ICB.
        unsafe {
            arg_encoder.setArgumentBuffer_offset(Some(&arg_buf), 0);
            arg_encoder.setIndirectCommandBuffer_atIndex(Some(&icb), 0);
        }
        // One status word per command slot: each view's decision dispatch
        // writes its region, the encode dispatch reads them all back.
        let status = self
            .hw
            .device
            .newBufferWithLength_options(
                new_cap * std::mem::size_of::<u32>(),
                MTLResourceOptions::StorageModePrivate,
            )
            .ok_or_else(|| super::error::allocation_failed("shadow view cull status buffer"))?;
        set.arg_buffer = Some(arg_buf);
        set.status = Some(status);
        set.icb = Some(icb);
        set.capacity = new_cap;
        Ok(())
    }

    // Ensure `slot_count` per-planar-slot mirror cull ICBs exist, each with a
    // command slot for every one of `count` draw objects (static + folded
    // instances + skinned, exactly like the main ICB). The slots' ICBs + argument
    // buffers are rebuilt when the slot count changes or the draw list outgrows
    // the capacity; the shared single-pass status scratch grows in lockstep. A
    // no-op for non-bindless contexts (no main ICB argument encoder to reuse) and
    // when `slot_count` is 0 (no planar set -> the slots are cleared). Called from
    // `draw_frame` right after `ensure_icb_capacity`, so the planar pass only ever
    // reads a sized mirror ICB.
    pub(super) fn ensure_mirror_icb_capacity(
        &mut self,
        slot_count: usize,
        count: usize,
    ) -> error::RenderResult<()> {
        if slot_count == 0 {
            self.cull.mirror_slots.clear();
            self.cull.mirror_status = None;
            self.cull.mirror_icb_capacity = 0;
            return Ok(());
        }
        // Reuse the main phase-1 ICB argument encoder: a mirror ICB has the
        // identical `ICBContainer` layout (one `[[id(0)]]` member). Absent on
        // non-bindless contexts, where there is no GPU cull to mirror.
        let arg_encoder = match &self.cull.icb_arg_encoder {
            Some(e) => e.clone(),
            None => return Ok(()),
        };
        if self.cull.mirror_slots.len() == slot_count && count <= self.cull.mirror_icb_capacity {
            return Ok(());
        }
        let new_cap = count.next_power_of_two().max(64);
        let mut slots = Vec::with_capacity(slot_count);
        for _ in 0..slot_count {
            let icb = self.build_cull_icb(new_cap)?;
            let len = arg_encoder.encodedLength().max(16);
            let arg_buffer = self
                .hw
                .device
                .newBufferWithLength_options(len, MTLResourceOptions::StorageModeShared)
                .ok_or_else(|| super::error::allocation_failed("mirror ICB argument buffer"))?;
            // SAFETY: the argument buffer is sized to `encodedLength()` and the
            // ICB is encoded at slot 0, exactly like the main + shadow ICBs. Each
            // slot re-points the shared encoder at its own (arg buffer, ICB) pair;
            // the encoding is fully written before the next slot re-points it.
            unsafe {
                arg_encoder.setArgumentBuffer_offset(Some(&arg_buffer), 0);
                arg_encoder.setIndirectCommandBuffer_atIndex(Some(&icb), 0);
            }
            slots.push(super::cull::MirrorCullSlot { icb, arg_buffer });
        }
        // Shared single-pass status scratch: the mirror cull writes per-object
        // status the same as phase 1, but nothing reads it (no phase-2 over the
        // mirror), so one buffer serves every slot. One u32 per command slot.
        let status = self
            .hw
            .device
            .newBufferWithLength_options(
                new_cap * std::mem::size_of::<u32>(),
                MTLResourceOptions::StorageModePrivate,
            )
            .ok_or_else(|| super::error::allocation_failed("mirror cull status buffer"))?;
        self.cull.mirror_status = Some(status);
        self.cull.mirror_slots = slots;
        self.cull.mirror_icb_capacity = new_cap;
        Ok(())
    }

    // Create one `DrawIndexed` indirect command buffer with `cap` command
    // slots. Both the phase-1 (`cull_icb`) and phase-2 (`cull_icb_2`) ICBs
    // share this shape: each command inherits the render encoder's buffer
    // bindings + pipeline state, so the cull kernels only encode the
    // indexed-draw arguments. Private storage keeps it GPU-resident
    // (kernel-written, render-pass-consumed, never CPU-touched).
    fn build_cull_icb(
        &self,
        cap: usize,
    ) -> error::RenderResult<Retained<ProtocolObject<dyn MTLIndirectCommandBuffer>>> {
        let desc = MTLIndirectCommandBufferDescriptor::new();
        desc.setCommandTypes(MTLIndirectCommandType::DrawIndexed);
        desc.setInheritBuffers(true);
        desc.setInheritPipelineState(true);
        desc.setMaxVertexBufferBindCount(0);
        desc.setMaxFragmentBufferBindCount(0);
        // SAFETY: `cap` is a valid command count; private storage as documented.
        unsafe {
            self.hw
                .device
                .newIndirectCommandBufferWithDescriptor_maxCommandCount_options(
                    &desc,
                    cap,
                    MTLResourceOptions::StorageModePrivate,
                )
        }
        .ok_or_else(|| super::error::allocation_failed("indirect command buffer"))
    }

    // Device capability flags for the settings menu. Ray tracing is queried
    // from the live MTLDevice (cheap; the same check the RT pass gates on).
    pub(crate) fn capabilities(&self) -> backend::DeviceCapabilities {
        backend::DeviceCapabilities {
            ray_tracing: super::raytrace::raytracing_supported(&self.hw.device),
            // Upscaling always goes through MetalFX; there is no selector.
            selectable_upscaler: false,
            // The per-frame RT topology refresh re-admits recycled build-time
            // slots, so every retired slot may be reused.
            reuses_build_slots: true,
            // The per-object buffer is rebuilt from `draw.objects` every frame,
            // so a slot's rewritten material draws on the next one.
            rewrites_draws: true,
        }
    }

    // Coarse GPU performance profile for default-quality selection, read live
    // from the MTLDevice (cheap; the same kind of device query as capabilities).
    pub(crate) fn gpu_profile(&self) -> backend::GpuProfile {
        super::gpu_profile::device_profile(&self.hw.device)
    }

    // Render statistics for the most recent `draw_frame`, for the profiler
    // overlay. The GPU frame time is the last value reported by a completed
    // command buffer, so it may lag the draw counts by a frame or two.
    pub(crate) fn render_stats(&self) -> profile::RenderStats {
        let mut stats = self.diagnostics.frame_stats;
        // `gpu_wait_us` was written by `draw_frame` itself and rides along in
        // `frame_stats`.
        stats.gpu_frame_us = self
            .diagnostics
            .gpu_time_us
            .load(std::sync::atomic::Ordering::Relaxed);
        // Per-pass timings are filled in pass-index order with their stable
        // names, leaving slots past PASS_COUNT at the default ("", 0).
        // Reports zero for any pass that did not write its sample slot this
        // frame (e.g. SSR when disabled, or any pass not yet wired up).
        for (i, name) in super::pass_timing::PASS_NAMES.iter().enumerate() {
            let micros =
                self.diagnostics.pass_times_us[i].load(std::sync::atomic::Ordering::Relaxed);
            stats.pass_times_us[i] = (*name, micros);
        }
        // Surface the auto-exposure EMA state. `None` when the world did not
        // opt in to auto-exposure (the static-exposure path leaves the field
        // empty so the StatHud chip stays blank).
        stats.auto_exposure_ev = self
            .auto_exposure
            .adaptation
            .as_ref()
            .map(|a| a.current_ev());
        // Surface the active panel's EDR headroom. `None` on SDR: both the
        // world-opt-out case and the request-on-an-SDR-display fallback case
        // map to the same blank chip.
        stats.max_edr = self.hw.max_edr();
        stats
    }

    // Replace one row of the material parameter table; every ring slot is
    // rewritten as the frames come round to it.
    pub(crate) fn set_material_params(
        &mut self,
        row: u32,
        params: [f32; render_types::MATERIAL_PARAM_COUNT],
    ) {
        self.rings.material_params.set(row, params);
    }

    // Append a projected-decal record at runtime, returning a stable slot
    // index the caller hands to [`Self::remove_decal`] later. Builds the
    // decal pipeline + unit-cube buffers on first use so a world that never
    // declared a decal still pays zero pipeline cost until the first add.
    // A vacated slot from [`Self::remove_decal`] is reused before growing
    // the slot table so a steady-state spawn/despawn pattern (bullet holes,
    // footprints) stays bounded.
    pub(crate) fn add_decal(&mut self, record: decal::DecalRecord) -> error::RenderResult<usize> {
        if self.decal.pipeline.is_none() {
            let (ps, vbuf, ibuf, samp) = super::init::world_fx::build_decal_resources_for_runtime(
                &self.hw.device,
                self.hot_reload.enabled,
            )?;
            self.decal.pipeline = Some(ps);
            self.decal.cube_vertex_buffer = Some(vbuf);
            self.decal.cube_index_buffer = Some(ibuf);
            self.decal.sampler = Some(samp);
        }
        self.decal
            .set
            .insert(record)
            .map_err(|_| error::RenderError::Other("add_decal: decal set is full".to_string()))
    }

    // Tombstone a runtime decal slot. The slot index returned by
    // [`Self::add_decal`] becomes invalid; the next add may reuse it.
    // Returns an error when the index is out of range or already tombstoned.
    // The decal pipeline + unit-cube buffers are kept around so a later add
    // does not pay the rebuild cost.
    pub(crate) fn remove_decal(&mut self, decal_id: usize) -> error::RenderResult<()> {
        self.decal
            .set
            .remove(decal_id)
            .map_err(|e| error::RenderError::Other(format!("remove_decal: id {decal_id} {e}")))
    }

    // Append a particle-emitter record at runtime, returning a stable slot
    // index. Allocates the per-emitter GPU pool + atomic counter buffer
    // (matching the init-time path) and builds the compute + render
    // pipelines on first use so a world that never declared an emitter
    // pays zero pipeline cost until the first add. Tombstoned slots from
    // [`Self::remove_emitter`] are reused before growing the vec.
    pub(crate) fn add_emitter(
        &mut self,
        record: particles::ParticleEmitterRecord,
    ) -> error::RenderResult<usize> {
        if self.particle.pipelines.is_none() {
            let pipelines = super::particle::build_particle_pipelines(
                &self.hw.device,
                self.hot_reload.enabled,
            )?;
            self.particle.pipelines = Some(pipelines);
        }
        let gpu_state = super::particle::build_emitter_gpu_state(&self.hw.device, &record)?;
        let idx = if let Some(slot) = self.particle.free_slots.pop() {
            self.particle.records[slot] = Some(record);
            self.particle.emitter_state[slot] = Some(gpu_state);
            slot
        } else {
            self.particle.records.push(Some(record));
            self.particle.emitter_state.push(Some(gpu_state));
            self.particle.records.len() - 1
        };
        Ok(idx)
    }

    // Tombstone a runtime emitter slot. Drops the `ParticleEmitterGpuState`:
    // Metal keeps the underlying pool + counter buffers alive until any
    // in-flight command buffer referencing them completes, so this is safe
    // to call mid-frame between encode passes (a debug tool call runs in
    // the per-frame hook window before the world step). Returns an error
    // when the index is out of range or already tombstoned.
    pub(crate) fn remove_emitter(&mut self, emitter_id: usize) -> error::RenderResult<()> {
        let rec_slot = self.particle.records.get_mut(emitter_id).ok_or_else(|| {
            error::RenderError::Other(format!("remove_emitter: id {emitter_id} out of range"))
        })?;
        if rec_slot.is_none() {
            return Err(error::RenderError::Other(format!(
                "remove_emitter: id {emitter_id} already removed"
            )));
        }
        *rec_slot = None;
        if let Some(gpu_slot) = self.particle.emitter_state.get_mut(emitter_id) {
            *gpu_slot = None;
        }
        self.particle.free_slots.push(emitter_id);
        Ok(())
    }

    // Returns true if the window has been closed by the user.
    pub(crate) fn window_closed(&self) -> bool {
        self.window().appkit.closed()
    }

    // Block until the GPU has finished all in-flight work.
    pub(crate) fn wait_idle(&self) {
        if let Some(cmd_buf) = self.hw.command_queue.commandBuffer() {
            cmd_buf.commit();
            cmd_buf.waitUntilCompleted();
        }
    }

    // Drain the classified GPU failure a completion handler parked, if any.
    pub(super) fn take_device_error(&self) -> Option<error::RenderError> {
        self.diagnostics
            .device_error
            .lock()
            .ok()
            .and_then(|mut slot| slot.take())
    }
}

impl scene_flow::SceneControl for MtlContext {
    fn update_visibility(&mut self, draw_idx: DrawIndex, visible: bool) {
        self.state.update_visibility(draw_idx, visible);
    }

    fn set_fade(&mut self, fade: f32) {
        self.state.set_fade(fade);
    }
}

impl Drop for MtlContext {
    fn drop(&mut self) {
        // Wait for all in-flight GPU work to finish before releasing Metal
        // objects. Without this, releasing the command queue while the GPU is
        // still executing a committed command buffer can corrupt ObjC retain
        // counts and produce EXC_BAD_ACCESS in objc_release.
        self.wait_idle();
        // Write whatever metallibs were compiled lazily since init's own
        // checkpoint. The window, when this context still holds one, closes as
        // `hw` drops after every GPU resource.
        crate::shader::runtime_cache::checkpoint();
    }
}

// Reinterpret a POD slice as raw bytes for a GPU buffer copy.
// Copy the first `len` bytes of `src` into `dst`. Both must be
// shared-storage buffers at least `len` bytes long; used by
// `setup_chunk_streaming` to carry build-time geometry into a grown buffer.
pub(super) fn copy_buffer_prefix(
    src: &ProtocolObject<dyn MTLBuffer>,
    dst: &ProtocolObject<dyn MTLBuffer>,
    len: usize,
) {
    if len == 0 {
        return;
    }
    assert!(
        len <= src.length() && len <= dst.length(),
        "copy_buffer_prefix: {len} bytes exceeds src {} / dst {}",
        src.length(),
        dst.length()
    );
    // SAFETY: both buffers are shared storage, so `contents()` is a live CPU
    // mapping, and the assert above proved each holds at least `len` bytes.
    // They are distinct allocations, so the ranges cannot overlap.
    unsafe {
        let s = src.contents().as_ptr() as *const u8;
        let d = dst.contents().as_ptr() as *mut u8;
        std::ptr::copy_nonoverlapping(s, d, len);
    }
}

pub(super) fn bytes_of_slice<T: bytemuck::NoUninit>(slice: &[T]) -> &[u8] {
    bytemuck::cast_slice(slice)
}

// Copy a `#[repr(C)]` slice into a shared-storage buffer at offset 0,
// bounds-checked against the buffer length.
pub(super) fn write_buffer_slice<T: Copy>(
    buffer: &ProtocolObject<dyn MTLBuffer>,
    data: &[T],
) -> error::RenderResult<()> {
    let bytes = std::mem::size_of_val(data);
    if bytes == 0 {
        return Ok(());
    }
    let len = buffer.length();
    if bytes > len {
        return Err(error::RenderError::Other(format!(
            "buffer write of {bytes} bytes exceeds buffer length {len}"
        )));
    }
    // SAFETY: `buffer` is shared storage so `contents()` is a live CPU mapping, and the bounds
    // check above proved it holds `bytes`. `data` is a separate live borrow of exactly that
    // many bytes (padding is copied, never inspected), so the ranges cannot overlap.
    unsafe {
        std::ptr::copy_nonoverlapping(
            data.as_ptr().cast::<u8>(),
            buffer.contents().as_ptr().cast::<u8>(),
            bytes,
        );
    }
    Ok(())
}

// Copy `src` into a shared-storage buffer at `offset` bytes, bounds-checked
// against the buffer length.
pub(super) fn write_buffer_region(
    buffer: &ProtocolObject<dyn MTLBuffer>,
    offset: usize,
    src: &[u8],
) -> error::RenderResult<()> {
    let len = buffer.length();
    if offset.checked_add(src.len()).is_none_or(|end| end > len) {
        return Err(error::RenderError::Other(format!(
            "buffer write [{}, {}) exceeds buffer length {}",
            offset,
            offset.saturating_add(src.len()),
            len
        )));
    }
    if src.is_empty() {
        return Ok(());
    }
    let dst = buffer.contents().as_ptr() as *mut u8;
    // SAFETY: `buffer` is shared storage so `contents()` is a live CPU mapping, and the bounds
    // check above proved `offset + src.len()` is within its length. `src` is a separate borrow, so
    // the ranges cannot overlap.
    unsafe {
        std::ptr::copy_nonoverlapping(src.as_ptr(), dst.add(offset), src.len());
    }
    Ok(())
}

// Zero a `len`-byte region of a shared-storage buffer at `offset` bytes,
// bounds-checked against the buffer length.
pub(super) fn zero_buffer_region(
    buffer: &ProtocolObject<dyn MTLBuffer>,
    offset: usize,
    len: usize,
) -> error::RenderResult<()> {
    let buf_len = buffer.length();
    if offset.checked_add(len).is_none_or(|end| end > buf_len) {
        return Err(error::RenderError::Other(format!(
            "buffer zero [{}, {}) exceeds buffer length {}",
            offset,
            offset.saturating_add(len),
            buf_len
        )));
    }
    if len == 0 {
        return Ok(());
    }
    let dst = buffer.contents().as_ptr() as *mut u8;
    // SAFETY: `buffer` is shared storage so `contents()` is a live CPU mapping, and the bounds
    // check above proved `offset + len` is within its length.
    unsafe {
        std::ptr::write_bytes(dst.add(offset), 0, len);
    }
    Ok(())
}

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

    fn counts(total: usize, skinned_base: usize) -> DrawRecordCounts {
        DrawRecordCounts {
            total,
            skinned_base,
        }
    }

    #[test]
    fn static_only_set_is_all_prefix() {
        let c = counts(12, 12);
        assert_eq!(c.prefix(0), Some(0..12));
        assert_eq!(c.skinned_tail(0), None);
    }

    #[test]
    fn folded_skinned_set_splits_at_the_base() {
        let c = counts(12, 9);
        assert_eq!(c.prefix(0), Some(0..9));
        assert_eq!(c.skinned_tail(0), Some(9..12));
    }

    #[test]
    fn skinned_only_set_has_no_prefix() {
        let c = counts(4, 0);
        assert_eq!(c.prefix(0), None);
        assert_eq!(c.skinned_tail(0), Some(0..4));
    }

    #[test]
    fn empty_set_yields_no_ranges() {
        let c = counts(0, 0);
        assert_eq!(c.prefix(0), None);
        assert_eq!(c.skinned_tail(0), None);
    }

    #[test]
    fn block_offsets_both_ranges_into_a_cascade() {
        let c = counts(12, 9);
        assert_eq!(c.prefix(24), Some(24..33));
        assert_eq!(c.skinned_tail(24), Some(33..36));
    }

    #[test]
    fn ranges_stay_inside_the_block_they_start() {
        let c = counts(12, 9);
        for cascade in 0..4 {
            let block = cascade * c.total;
            let tail = c.skinned_tail(block).expect("skinned tail");
            assert_eq!(c.prefix(block).expect("prefix").start, block);
            assert_eq!(tail.end, block + c.total);
        }
    }

    #[test]
    fn ns_range_carries_start_and_length() {
        let r = ns_range(9..12);
        assert_eq!(r.location, 9);
        assert_eq!(r.length, 3);
    }
}