concinnity-device 0.19.119

//! Hi-Z (depth-mip pyramid) build pass used by the GPU-cull compute kernel for
//! occlusion culling. Each frame, after the main depth buffer has been written
//! by the graph, we reduce it into an `R32_SFLOAT` mip chain (MAX reduction:
//! standard depth, so larger = farther). The *next* frame's `Cull` kernel
//! projects each `DrawObject` AABB through the previous frame's un-jittered
//! view-projection, picks the Hi-Z mip whose texels are ~the size of the
//! projected rect, 4-tap-samples the max occluder depth, and culls the AABB when
//! its nearest projected NDC depth is strictly behind. Mirrors the DirectX
//! implementation in `directx/hiz.rs`; Metal keeps the older per-mip chain (see
//! `metal/hiz.rs`). Every backend's kernels ship from the single-source
//! `src/render/shaders/hiz_build.hlsl` (one variant compile per kernel):
//!
//!   * `hiz_spd_single`: reduce a single-sample main depth into mips 0..6.
//!   * `hiz_spd_msaa`  : the same for an MSAA main depth, taking the MAX over
//!     every sample so the result is conservative.
//!   * `hiz_spd_tail`  : continue from mip 6 into mips 7..12.
//!
//! Each workgroup reduces a 64x64 tile through seven levels, so the whole
//! pyramid is two dispatches with one barrier between them rather than one
//! dispatch and one barrier per mip. `core::render::hiz_spd::Plan` decides the
//! dispatch geometry; DirectX builds its pyramid from the same plan.
//!
//! The build is a graph node: `HizFinal` (terminal, reducing the frame's last
//! depth version for the next frame's cull) and, under two-pass occlusion,
//! `HizBuild` (mid-frame, reducing phase-1 depth for `Cull2`). Both dispatch this
//! encoder, so the graph owns the pyramid's lifetime and the main depth's layout
//! chain reaches the end of the frame.
//!
//! Each mip is written through its own single-level R32F storage-image view; the
//! whole Hi-Z image stays in GENERAL during the build, with a compute
//! write -> read memory barrier between each step. That per-mip chain is finer
//! than the graph's one-state-per-resource granularity, so it stays inline here;
//! the open and close around it are graph-derived. Between frames the image rests
//! in `SHADER_READ_ONLY_OPTIMAL` so the cull kernel reads it as a sampled image
//! (set 1). A single shared image read one frame and written the next is
//! hazard-free on a single queue: the executor's end-of-frame restore (GENERAL ->
//! SHADER_READ_ONLY) orders the write before the next frame's cull read, and the
//! producer barrier derived from `Cull`'s declared pyramid read orders that read
//! before this frame's write.

use ash::vk;
use concinnity_core::render::error::RenderResult;

use super::allocator::{DeviceAllocator, PooledBuffer, PooledImage};
use super::descriptor_layout::{Binding, PoolSizes};
use super::pipeline_desc::compute_pipeline;
use super::resources::{alloc_descriptor_sets, create_descriptor_set_layout};
use super::set_writes::SetWrites;
use super::texture::{
    LayoutTransition, SubresourceRange, one_shot_submit, transition_image_layout_range,
};
use crate::vulkan::owned::{
    OwnedDescriptorPool, OwnedPipeline, OwnedPipelineLayout, OwnedSetLayout, VkDevice,
};
use crate::vulkan::record::Recorder;

// Upper bound on the Hi-Z mip count, used to size the dedicated descriptor pool
// for the tail's sets. `hiz_mip_count` caps at 32 - leading_zeros,
// so 16 covers any render target up to 32768 px on its longer edge.
const MAX_HIZ_MIPS: usize = 16;

use crate::vulkan::builtin_shaders::CompileProgram;
use concinnity_core::render::hiz_spd::{self, Plan};
use concinnity_core::render::uniforms::HizSpdParams;
use concinnity_core::render::uniforms::vulkan::CullHizParams;

// Mip count for a Hi-Z of size (w, h): `floor(log2(max(w, h))) + 1`. Power-of-
// two sources end exactly at 1x1; non-power-of-two sources stop one mip short
// of true 1x1 in the smaller dimension, which is fine: the cull kernel clamps
// to the actual mip dims. Mirrors `directx::hiz::hiz_mip_count` /
// `metal::hiz::hiz_mip_count`.
pub(super) fn hiz_mip_count(width: u32, height: u32) -> u32 {
    let m = width.max(height).max(1);
    32 - m.leading_zeros()
}

// Compute pipelines + image + per-mip views + descriptor sets for the Hi-Z
// build, plus the cull-read set (set 1 of the cull pipeline) and its per-frame
// uniform buffers. `Some` on the context exactly when the GPU-cull pipeline is
// active (same gating as `cull_kernels`).
pub(super) struct HiZResources {
    // Build pipelines + their layouts (the phase-1 SPD and tail kernels bind
    // different set layouts, so each needs its own pipeline layout).
    spd_pipeline: OwnedPipeline,
    spd_tail_pipeline: OwnedPipeline,
    spd_pipeline_layout: OwnedPipelineLayout,
    spd_tail_pipeline_layout: OwnedPipelineLayout,
    spd_set_layout: OwnedSetLayout,
    spd_tail_set_layout: OwnedSetLayout,

    // Cull-read set layout (set 1 of the cull pipeline): the Hi-Z sampled image
    // + CullHizParams UBO. Held here because `init.rs` threads it into the cull
    // pipeline layout, and the layout survives a resize.
    pub(super) read_set_layout: OwnedSetLayout,

    // Dedicated descriptor pool for every Hi-Z set (build + cull-read).
    descriptor_pool: OwnedDescriptorPool,

    // R32F mip-chain image. Written mip-by-mip during the build (GENERAL),
    // sampled by the cull kernel between frames (SHADER_READ_ONLY). Its views
    // below are attached to the lease, so replacing it on a resize retires the
    // whole set together. The graph executor resolves it as the `hiz_pyramid`
    // barrier target.
    pub(super) pyramid: PooledImage,
    // All-mips sampled view bound in the cull-read set.
    sampled_view: vk::ImageView,
    // One single-level storage view per mip, bound as the mips phase 1 writes
    // and the level the tail reduces from. Length = `mip_count`.
    mip_views: Vec<vk::ImageView>,
    // Build sets: one phase-1 set per frame (depth differs per frame slot) and
    // one tail set (frame-independent, Hi-Z mips only).
    spd_sets: Vec<vk::DescriptorSet>,
    spd_tail_sets: Vec<vk::DescriptorSet>,
    // Cull-read sets, one per frame (the UBO differs per frame slot).
    pub(super) read_sets: Vec<vk::DescriptorSet>,
    // Per-frame CullHizParams uniform buffers (host-mapped), bound in
    // `read_sets[i]` binding 1 and written by `encode_cull`.
    pub(super) cull_ubos: Vec<PooledBuffer>,

    // Two-pass occlusion phase-2 cull-read sets + their own per-frame UBOs.
    // Empty unless two-pass occlusion is active. The phase-2 `Cull2` dispatch
    // needs a separate UBO from phase 1 because both consume their UBO at
    // different GPU times within one frame (phase 1's `prev_view_proj`, phase
    // 2's current-frame VP); sharing one host-mapped buffer would clobber
    // phase 1's value before the GPU reads it. The image binding points at
    // the same pyramid `sampled_view`, re-pointed alongside `read_sets` on a
    // resize. Uses the shared `read_set_layout`.
    pub(super) read_sets2: Vec<vk::DescriptorSet>,
    pub(super) cull_ubos2: Vec<PooledBuffer>,

    pub(super) width: u32,
    pub(super) height: u32,
    pub(super) mip_count: u32,
    // MSAA sample count of the main depth the phase-1 kernel reduces (1 when the
    // world is single-sampled).
    sample_count: u32,
}

// Create the R32F mip-chain image (STORAGE + SAMPLED), GPU-local.
fn create_hiz_image(
    alloc: &DeviceAllocator,
    width: u32,
    height: u32,
    mip_count: u32,
) -> RenderResult<PooledImage> {
    let img_info = vk::ImageCreateInfo::default()
        .image_type(vk::ImageType::TYPE_2D)
        .extent(vk::Extent3D {
            width: width.max(1),
            height: height.max(1),
            depth: 1,
        })
        .mip_levels(mip_count.max(1))
        .array_layers(1)
        .format(vk::Format::R32_SFLOAT)
        .tiling(vk::ImageTiling::OPTIMAL)
        .initial_layout(vk::ImageLayout::UNDEFINED)
        .usage(vk::ImageUsageFlags::STORAGE | vk::ImageUsageFlags::SAMPLED)
        .sharing_mode(vk::SharingMode::EXCLUSIVE)
        .samples(vk::SampleCountFlags::TYPE_1);
    alloc
        .create_image(&img_info, vk::MemoryPropertyFlags::DEVICE_LOCAL)
        .map_err(|e| e.context("hiz image"))
}

// A single-level (`mip`) or all-mips (`base_mip = 0`, `count = mip_count`) 2D
// R32F view of the Hi-Z image.
fn create_hiz_view(
    device: &VkDevice,
    image: vk::Image,
    base_mip: u32,
    level_count: u32,
) -> RenderResult<vk::ImageView> {
    let info = vk::ImageViewCreateInfo::default()
        .image(image)
        .view_type(vk::ImageViewType::TYPE_2D)
        .format(vk::Format::R32_SFLOAT)
        .subresource_range(vk::ImageSubresourceRange {
            aspect_mask: vk::ImageAspectFlags::COLOR,
            base_mip_level: base_mip,
            level_count,
            base_array_layer: 0,
            layer_count: 1,
        });
    // SAFETY: the create-info and every slice it borrows are live for the call, and each handle it
    // names belongs to this device.
    unsafe { device.create_image_view(&info, None) }
        .map_err(|e| super::error::map_vk_result(e, "hiz view"))
}

// Build the phase-1 SPD and tail pipelines for the given MSAA mode. Returns the
// two pipelines; the layouts are created by the caller and outlive a hot-reload.
fn build_hiz_pipelines(
    device: &VkDevice,
    spd_layout: vk::PipelineLayout,
    spd_tail_layout: vk::PipelineLayout,
    sample_count: u32,
    hot_reload: bool,
) -> RenderResult<(OwnedPipeline, OwnedPipeline)> {
    // Phase 1 is a per-variant compile of the single-source `hiz_build.hlsl`:
    // the depth resource is a `Texture2DMS` when multisampled, a `Texture2D`
    // otherwise (a sampled image either way; the kernel reads texels by
    // coordinate, so no sampler is bound).
    let phase1_spv = if sample_count > 1 {
        super::builtin_shaders::HIZ_SPD_MSAA.compile(hot_reload)?
    } else {
        super::builtin_shaders::HIZ_SPD_SINGLE.compile(hot_reload)?
    };
    let tail_spv = super::builtin_shaders::HIZ_SPD_TAIL.compile(hot_reload)?;
    let phase1 = compute_pipeline(device, spd_layout, &phase1_spv, "hiz")?;
    let tail = compute_pipeline(device, spd_tail_layout, &tail_spv, "hiz")?;
    Ok((phase1, tail))
}

// Vulkan device + one-shot submission context threaded into the Hi-Z resource
// (re)creation calls. Bundles the handles needed to allocate GPU memory and run
// the layout-transition submit.
#[derive(Clone, Copy)]
pub(super) struct HiZDeviceCtx<'a> {
    pub(super) alloc: &'a DeviceAllocator,
    pub(super) device: &'a VkDevice,
    pub(super) command_pool: vk::CommandPool,
    pub(super) queue: vk::Queue,
}

// The render (depth) target the Hi-Z pyramid mirrors: its dimensions plus the
// per-frame main-depth views the phase-1 kernel reduces.
#[derive(Clone, Copy)]
pub(super) struct HiZTarget<'a> {
    pub(super) width: u32,
    pub(super) height: u32,
    pub(super) depth_views: &'a [vk::ImageView],
}

impl HiZResources {
    // The pyramid's all-mips sampled view, the image a cull set-1 ("read set")
    // binds at binding 0. Exposed so the reflection-probe bake can build a one-off
    // read set (with `hiz_enabled = 0`) from its OWN descriptor pool -- the probe
    // cull binds a valid set 1 without reading the pyramid, and without taking a
    // slot in this struct's pool. `read_set_layout` is already `pub(super)`.
    pub(super) fn read_set_view(&self) -> vk::ImageView {
        self.sampled_view
    }

    // Build every Hi-Z resource sized to the render (depth) resolution. Called
    // from the init path when the GPU-cull pipeline is active. `target.depth_views`
    // are the per-frame main-depth views the phase-1 kernel reduces.
    pub(super) fn new(
        ctx: HiZDeviceCtx,
        target: HiZTarget,
        sample_count: u32,
        frames: usize,
        // When set, allocate the phase-2 cull-read sets + UBOs for two-pass
        // occlusion (`Cull2`). Gated on the world's `occlusion_two_pass`.
        two_pass: bool,
        hot_reload: bool,
    ) -> RenderResult<Self> {
        // `command_pool`, `queue`, and `depth_views` are only needed by the
        // `create_image_and_sets` call below, which takes `ctx` / `target` whole.
        let HiZDeviceCtx { alloc, device, .. } = ctx;
        let HiZTarget { width, height, .. } = target;
        // Set layouts.
        // Phase 1: binding 0 depth (sampled image, read by coordinate -- no
        // sampler), binding 1 the mips it writes.
        let spd_set_layout = create_set_layout_counted(
            device,
            &[
                (0, vk::DescriptorType::SAMPLED_IMAGE, 1),
                (1, vk::DescriptorType::STORAGE_IMAGE, hiz_spd::LEVELS),
            ],
        )?;
        // Tail: binding 0 the mips it writes, whose first element is also the
        // level it reduces.
        let spd_tail_set_layout = create_set_layout_counted(
            device,
            &[(0, vk::DescriptorType::STORAGE_IMAGE, hiz_spd::LEVELS)],
        )?;
        let read_set_layout = create_descriptor_set_layout(device, &read_set_bindings())?;

        // Pipeline layouts (shared 16-byte push range for both build kernels).
        let push_range = vk::PushConstantRange::default()
            .stage_flags(vk::ShaderStageFlags::COMPUTE)
            .offset(0)
            .size(std::mem::size_of::<HizSpdParams>() as u32);
        let spd_pipeline_layout =
            create_pipeline_layout(device, spd_set_layout.handle(), push_range)?;
        let spd_tail_pipeline_layout =
            create_pipeline_layout(device, spd_tail_set_layout.handle(), push_range)?;

        let (spd_pipeline, spd_tail_pipeline) = build_hiz_pipelines(
            device,
            spd_pipeline_layout.handle(),
            spd_tail_pipeline_layout.handle(),
            sample_count,
            hot_reload,
        )?;

        // Dedicated descriptor pool, sized for the worst-case mip count.
        let descriptor_pool = create_pool(device, frames, two_pass)?;

        // Per-frame cull-read uniform buffers (host-mapped). The phase-2 set
        // gets its own ring (`cull_ubos2`) when two-pass occlusion is active.
        let ubo_size = std::mem::size_of::<CullHizParams>() as u64;
        let alloc_ubo_ring = |count: usize| -> RenderResult<Vec<PooledBuffer>> {
            (0..count)
                .map(|_| {
                    alloc.create_buffer(
                        ubo_size,
                        vk::BufferUsageFlags::UNIFORM_BUFFER,
                        vk::MemoryPropertyFlags::HOST_VISIBLE
                            | vk::MemoryPropertyFlags::HOST_COHERENT,
                    )
                })
                .collect()
        };
        let cull_ubos = alloc_ubo_ring(frames)?;
        let cull_ubos2 = alloc_ubo_ring(if two_pass { frames } else { 0 })?;

        let mut res = Self {
            spd_pipeline,
            spd_tail_pipeline,
            spd_pipeline_layout,
            spd_tail_pipeline_layout,
            spd_set_layout,
            spd_tail_set_layout,
            read_set_layout,
            descriptor_pool,
            pyramid: PooledImage::null(),
            sampled_view: vk::ImageView::null(),
            mip_views: Vec::new(),
            spd_sets: Vec::new(),
            spd_tail_sets: Vec::new(),
            read_sets: Vec::new(),
            cull_ubos,
            read_sets2: Vec::new(),
            cull_ubos2,
            width,
            height,
            mip_count: 0,
            sample_count,
        };
        res.create_image_and_sets(ctx, target)?;
        Ok(res)
    }

    // (Re)create the mip-chain image + views and (re)allocate every set bound
    // to it. Resets the descriptor pool, so all Hi-Z sets are freshly
    // allocated; the cull-read UBO buffers themselves survive (only their
    // descriptors are rewritten). The caller must have idled the GPU.
    fn create_image_and_sets(&mut self, ctx: HiZDeviceCtx, target: HiZTarget) -> RenderResult<()> {
        let HiZDeviceCtx {
            alloc,
            device,
            command_pool,
            queue,
        } = ctx;
        let HiZTarget {
            width,
            height,
            depth_views,
        } = target;
        let requested = hiz_mip_count(width, height).min(MAX_HIZ_MIPS as u32).max(1);
        // The depth the two SPD dispatches actually reach, which is what the
        // image carries and what the cull is told: a mip the plan skipped is
        // never written, and sampling one would feed the cull uninitialized
        // memory.
        let mip_count = Plan::new(width, height, requested, 1).mip_count();
        let pyramid = create_hiz_image(alloc, width, height, mip_count)?;
        // Rest in SHADER_READ_ONLY so the cull-read descriptor's layout is
        // satisfied on the first frame (the cull kernel won't sample it -
        // `hiz_enabled` is 0 - but the descriptor layout must still match).
        one_shot_submit(device, command_pool, queue, |cmd| {
            transition_image_layout_range(
                device,
                cmd,
                pyramid.image(),
                LayoutTransition {
                    old_layout: vk::ImageLayout::UNDEFINED,
                    new_layout: vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL,
                    aspect: vk::ImageAspectFlags::COLOR,
                },
                SubresourceRange {
                    base_layer: 0,
                    layer_count: 1,
                    base_mip: 0,
                    mip_count,
                },
            );
        })?;

        let sampled_view = create_hiz_view(device, pyramid.image(), 0, mip_count)?;
        pyramid.attach_view(sampled_view);
        let mut mip_views = Vec::with_capacity(mip_count as usize);
        for mip in 0..mip_count {
            let view = create_hiz_view(device, pyramid.image(), mip, 1)?;
            pyramid.attach_view(view);
            mip_views.push(view);
        }

        // Reset the pool and reallocate every set (phase 1 / tail / read).
        // SAFETY: `descriptor_pool` was created from this device and every set allocated from it is
        // dropped here; the caller has already idled the device, so none is still in use.
        unsafe {
            device
                .reset_descriptor_pool(
                    self.descriptor_pool.handle(),
                    vk::DescriptorPoolResetFlags::empty(),
                )
                .map_err(|e| super::error::map_vk_result(e, "reset hiz pool"))?;
        }
        let frames = self.cull_ubos.len();
        let spd_layouts: Vec<_> = (0..frames).map(|_| self.spd_set_layout.handle()).collect();
        let spd_sets = alloc_descriptor_sets(device, self.descriptor_pool.handle(), &spd_layouts)?;
        let spd_tail_layouts = [self.spd_tail_set_layout.handle()];
        let spd_tail_sets =
            alloc_descriptor_sets(device, self.descriptor_pool.handle(), &spd_tail_layouts)?;
        let read_layouts: Vec<_> = (0..frames).map(|_| self.read_set_layout.handle()).collect();
        let read_sets =
            alloc_descriptor_sets(device, self.descriptor_pool.handle(), &read_layouts)?;
        // Phase-2 cull-read sets (two-pass occlusion), one per frame, only when
        // the phase-2 UBO ring was allocated.
        let two_pass = !self.cull_ubos2.is_empty();
        let read_layouts2: Vec<_> = (0..if two_pass { frames } else { 0 })
            .map(|_| self.read_set_layout.handle())
            .collect();
        let read_sets2 =
            alloc_descriptor_sets(device, self.descriptor_pool.handle(), &read_layouts2)?;

        // The mip array each dispatch binds. Elements past the last live mip
        // repeat it so the array is fully populated even where the kernel's
        // `level_count` stops it writing through those elements.
        let bound = |base_mip: u32| -> Vec<vk::DescriptorImageInfo> {
            Plan::bound_mips(base_mip, mip_count)
                .map(|m| {
                    vk::DescriptorImageInfo::default()
                        .image_layout(vk::ImageLayout::GENERAL)
                        .image_view(mip_views[m as usize])
                })
                .collect()
        };
        // Phase-1 sets: binding 0 = that frame's main depth, binding 1 = mips 0..6.
        let phase1_mips = bound(0);
        for (i, &set) in spd_sets.iter().enumerate() {
            let depth = depth_views[i.min(depth_views.len().saturating_sub(1))];
            SetWrites::new(set)
                .sampled_image(0, depth)
                .images(1, vk::DescriptorType::STORAGE_IMAGE, &phase1_mips)
                .apply(device);
        }
        // Tail set: mips 6..12, the first of which is the level it reduces.
        let tail_mips = bound(hiz_spd::LEVELS - 1);
        for &set in &spd_tail_sets {
            SetWrites::new(set)
                .images(0, vk::DescriptorType::STORAGE_IMAGE, &tail_mips)
                .apply(device);
        }
        // Read sets: the all-mips Hi-Z view and the per-frame cull UBO; the
        // phase-2 sets bind the phase-2 ring.
        let ubo_size = std::mem::size_of::<CullHizParams>() as u64;
        for (&set, ubo) in read_sets.iter().zip(&self.cull_ubos) {
            write_read_set(device, set, sampled_view, ubo.buffer(), ubo_size);
        }
        for (&set, ubo) in read_sets2.iter().zip(&self.cull_ubos2) {
            write_read_set(device, set, sampled_view, ubo.buffer(), ubo_size);
        }

        // Replacing the pooled image drops the previous lease: the old pyramid
        // and every view attached to it retire through the allocator.
        self.pyramid = pyramid;
        self.sampled_view = sampled_view;
        self.mip_views = mip_views;
        self.spd_sets = spd_sets;
        self.spd_tail_sets = spd_tail_sets;
        self.read_sets = read_sets;
        self.read_sets2 = read_sets2;
        self.width = width;
        self.height = height;
        self.mip_count = mip_count;
        Ok(())
    }

    // Recreate the image + views + sets at new render-target dimensions. The
    // pipelines, layouts and cull-read UBO buffers survive; the old
    // pyramid retires through the allocator when the new one replaces it. The
    // caller flips `hiz_valid` to false so the next cull dispatch ignores the
    // now-stale pyramid.
    pub(super) fn resize_to(&mut self, ctx: HiZDeviceCtx, target: HiZTarget) -> RenderResult<()> {
        self.create_image_and_sets(ctx, target)
    }

    // Swap freshly-rebuilt pipelines into the live resource. Used by the shader
    // hot-reload pass; the image, views, sets, and layouts are kept.
    pub(super) fn swap_pipelines(&mut self, spd: OwnedPipeline, tail: OwnedPipeline) {
        self.spd_pipeline = spd;
        self.spd_tail_pipeline = tail;
    }

    // Recompile the build pipelines from disk-resident source (hot-reload).
    // The MSAA mode is fixed at init, so it is reused here.
    pub(super) fn recompile_pipelines(
        &self,
        device: &VkDevice,
        hot_reload: bool,
    ) -> RenderResult<(OwnedPipeline, OwnedPipeline)> {
        build_hiz_pipelines(
            device,
            self.spd_pipeline_layout.handle(),
            self.spd_tail_pipeline_layout.handle(),
            self.sample_count,
            hot_reload,
        )
    }

    // Destroy every non-pooled GPU resource and drop the pooled ones (the
    // pyramid + its views and the cull UBO rings retire through the allocator).
    pub(super) fn destroy(&mut self, _device: &VkDevice) {
        self.pyramid = PooledImage::null();
        self.sampled_view = vk::ImageView::null();
        self.mip_views.clear();
        self.cull_ubos.clear();
        self.cull_ubos2.clear();
    }
}

impl crate::vulkan::context::VkContext {
    // Encode the Hi-Z build into `cmd`. Reads this frame's main depth
    // (`depth_images[frame_idx]`) and writes the mip chain that *next* frame's
    // cull dispatch consults. A no-op when no Hi-Z resource was built (GPU-cull
    // pipeline not active). Runs as the graph's `HizBuild` (mid-frame, phase-1
    // depth) or `HizFinal` (terminal, the frame's last depth version) node, so
    // the executor has already put main depth in SHADER_READ_ONLY and the
    // pyramid in GENERAL; only the barrier between the two dispatches below is
    // this encoder's.
    pub(in crate::vulkan) fn encode_hiz_build(&self, rec: &Recorder<'_>, frame_idx: usize) {
        let Some(hiz) = self.cull.hiz.as_ref() else {
            return;
        };
        if hiz.mip_count == 0 || hiz.mip_views.is_empty() {
            return;
        }
        let plan = Plan::new(
            hiz.width,
            hiz.height,
            hiz.mip_count,
            hiz.sample_count.max(1),
        );

        // Phase 1: main depth into mips 0..6.
        let spd_set = hiz.spd_sets[frame_idx.min(hiz.spd_sets.len().saturating_sub(1))];
        rec.bind_pipeline(vk::PipelineBindPoint::COMPUTE, &hiz.spd_pipeline);
        rec.bind_descriptor_sets(
            vk::PipelineBindPoint::COMPUTE,
            &hiz.spd_pipeline_layout,
            0,
            std::slice::from_ref(&spd_set),
            &[],
        );
        rec.push_constants(
            &hiz.spd_pipeline_layout,
            vk::ShaderStageFlags::COMPUTE,
            0,
            &plan.phase1.params,
        );
        rec.dispatch(plan.phase1.groups.0, plan.phase1.groups.1, 1);

        let (Some(tail), Some(&tail_set)) = (plan.tail, hiz.spd_tail_sets.first()) else {
            return;
        };
        // The tail reduces the mip 6 phase 1 just wrote, so the pyramid needs
        // one write -> read barrier here. It is the only one the build takes;
        // the graph owns the dependencies on either side of the node, and every
        // mip stays in GENERAL throughout.
        rec.pipeline_barrier(
            vk::PipelineStageFlags::COMPUTE_SHADER,
            vk::PipelineStageFlags::COMPUTE_SHADER,
            &[vk::MemoryBarrier::default()
                .src_access_mask(vk::AccessFlags::SHADER_WRITE)
                .dst_access_mask(vk::AccessFlags::SHADER_READ)],
            &[],
            &[],
        );
        rec.bind_pipeline(vk::PipelineBindPoint::COMPUTE, &hiz.spd_tail_pipeline);
        rec.bind_descriptor_sets(
            vk::PipelineBindPoint::COMPUTE,
            &hiz.spd_tail_pipeline_layout,
            0,
            std::slice::from_ref(&tail_set),
            &[],
        );
        rec.push_constants(
            &hiz.spd_tail_pipeline_layout,
            vk::ShaderStageFlags::COMPUTE,
            0,
            &tail.params,
        );
        rec.dispatch(tail.groups.0, tail.groups.1, 1);
    }
}

// Cull-read set (set 1 of the cull pipeline): the Hi-Z image, read by texel
// coordinate so no sampler is bound, and the cull UBO.
fn read_set_bindings() -> [Binding; 2] {
    let compute = vk::ShaderStageFlags::COMPUTE;
    [
        (0, vk::DescriptorType::SAMPLED_IMAGE, compute),
        (1, vk::DescriptorType::UNIFORM_BUFFER, compute),
    ]
}

// A compute set layout with an explicit descriptor count per binding: the SPD
// sets bind their destination mips as one array.
fn create_set_layout_counted(
    device: &VkDevice,
    bindings: &[(u32, vk::DescriptorType, u32)],
) -> RenderResult<OwnedSetLayout> {
    let binds: Vec<_> = bindings
        .iter()
        .map(|&(b, ty, count)| {
            vk::DescriptorSetLayoutBinding::default()
                .binding(b)
                .descriptor_type(ty)
                .descriptor_count(count)
                .stage_flags(vk::ShaderStageFlags::COMPUTE)
        })
        .collect();
    device
        .create_descriptor_set_layout(
            &vk::DescriptorSetLayoutCreateInfo::default().bindings(&binds),
        )
        .map_err(|e| super::error::map_vk_result(e, "hiz set layout"))
}

fn create_pipeline_layout(
    device: &VkDevice,
    set_layout: vk::DescriptorSetLayout,
    push_range: vk::PushConstantRange,
) -> RenderResult<OwnedPipelineLayout> {
    let layouts = [set_layout];
    device
        .create_pipeline_layout(
            &vk::PipelineLayoutCreateInfo::default()
                .set_layouts(&layouts)
                .push_constant_ranges(std::slice::from_ref(&push_range)),
        )
        .map_err(|e| super::error::map_vk_result(e, "hiz pipeline layout"))
}

fn create_pool(
    device: &VkDevice,
    frames: usize,
    two_pass: bool,
) -> RenderResult<OwnedDescriptorPool> {
    let f = frames as u32;
    // Two-pass occlusion adds one extra cull-read set per frame (phase 2),
    // each with an image + a UBO descriptor.
    let read_rings = if two_pass { 2 } else { 1 };
    // Phase 1 adds its depth (frames) and a mip array per set, as does the tail.
    let sizes = PoolSizes::default()
        .sets(&read_set_bindings(), read_rings * f)
        .add(vk::DescriptorType::SAMPLED_IMAGE, f)
        .add(vk::DescriptorType::STORAGE_IMAGE, (f + 1) * hiz_spd::LEVELS)
        .build();
    // phase 1 (frames) + cull-read (frames per read ring) + the one tail set.
    let max_sets = (1 + read_rings) * f + 1;
    device
        .create_descriptor_pool(
            &vk::DescriptorPoolCreateInfo::default()
                .pool_sizes(&sizes)
                .max_sets(max_sets),
        )
        .map_err(|e| super::error::map_vk_result(e, "hiz descriptor pool"))
}

// A read set (cull set 1) for an off-camera cull (a probe face, a planar
// mirror), allocated from the caller's `pool`: the pyramid is bound because the
// cull layout references set 1, but `hiz_enabled = 0` keeps the frustum-only
// cull from sampling a pyramid built for another view. The returned UBO backs
// the set.
pub(super) fn off_camera_read_set(
    alloc: &DeviceAllocator,
    device: &VkDevice,
    pool: vk::DescriptorPool,
    layout: vk::DescriptorSetLayout,
    view: vk::ImageView,
) -> RenderResult<(vk::DescriptorSet, PooledBuffer)> {
    let params = CullHizParams {
        prev_view_proj: [[0.0; 4]; 4],
        hiz_size: [1.0, 1.0],
        hiz_mip_count: 1,
        hiz_enabled: 0,
    };
    let size = std::mem::size_of::<CullHizParams>() as u64;
    let ubo = alloc.create_buffer(
        size,
        vk::BufferUsageFlags::UNIFORM_BUFFER,
        vk::MemoryPropertyFlags::HOST_VISIBLE | vk::MemoryPropertyFlags::HOST_COHERENT,
    )?;
    ubo.write_val(0, &params);
    let set = alloc_descriptor_sets(device, pool, std::slice::from_ref(&layout))?[0];
    write_read_set(device, set, view, ubo.buffer(), size);
    Ok((set, ubo))
}

// Re-point an off-camera read set at a rebuilt pyramid view: resize retires
// the view it captured, and the set must not dangle even though it is never
// sampled.
pub(super) fn rewrite_read_set_view(
    device: &VkDevice,
    set: vk::DescriptorSet,
    view: vk::ImageView,
) {
    SetWrites::new(set).sampled_image(0, view).apply(device);
}

fn write_read_set(
    device: &VkDevice,
    set: vk::DescriptorSet,
    view: vk::ImageView,
    ubo: vk::Buffer,
    ubo_size: u64,
) {
    SetWrites::new(set)
        .sampled_image(0, view)
        .uniform_buffer(1, ubo, ubo_size)
        .apply(device);
}

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

    // The `HizParams` / `CullHizParams` layout tests live with the structs in
    // `concinnity_core::render::uniforms::vulkan`.

    #[test]
    fn mip_count_power_of_two() {
        assert_eq!(hiz_mip_count(1, 1), 1);
        assert_eq!(hiz_mip_count(2, 2), 2);
        assert_eq!(hiz_mip_count(256, 256), 9);
        assert_eq!(hiz_mip_count(1024, 1024), 11);
    }

    #[test]
    fn mip_count_uses_larger_dimension() {
        assert_eq!(hiz_mip_count(1920, 1080), hiz_mip_count(1920, 1920));
        assert_eq!(hiz_mip_count(1920, 1080), 11);
        assert_eq!(hiz_mip_count(1280, 720), 11);
    }

    #[test]
    fn mip_count_clamps_zero() {
        assert_eq!(hiz_mip_count(0, 0), 1);
        assert_eq!(hiz_mip_count(0, 8), 4);
    }
}