concinnity-device 0.18.65

// src/vulkan/hiz.rs
//
// 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` and the Metal one in `metal/hiz.rs`; every
// backend's kernels ship from the single-source
// `src/shaders/hiz_build.slang` (one variant compile per kernel):
//
//   * `hiz_init_msaa` / `hiz_init_single`: reduce the main depth into mip 0,
//     taking the MAX over every sample so the result is conservative.
//   * `hiz_downsample`: MAX-reduce 2x2 source texels into the next mip.
//
// 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 samples it via a `sampler2D`
// (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 crate::vulkan::owned::{
    OwnedDescriptorPool, OwnedPipeline, OwnedPipelineLayout, OwnedSampler, OwnedSetLayout, VkDevice,
};

use super::allocator::{DeviceAllocator, PooledBuffer, PooledImage};
use super::pipeline::spv_module;
use super::resources::alloc_descriptor_sets;
use super::texture::{
    LayoutTransition, SubresourceRange, one_shot_submit, transition_image_layout_range,
};

// Upper bound on the Hi-Z mip count, used to size the dedicated descriptor pool
// for the per-downsample-step 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;

// Compute threadgroup tile size for the Hi-Z build kernels (8x8, matching the
// DirectX `[numthreads(8, 8, 1)]` and the Metal `HIZ_TILE`).
const HIZ_TILE: u32 = 8;

// `HizParams` (Hi-Z build push constant) and `CullHizParams` (cull-side Hi-Z
// std140 UBO) are GPU-free layout structs that live in concinnity-render;
// re-export them so `crate::vulkan::hiz::{HizParams,CullHizParams}` are
// unchanged for the passes that fill them.
pub(in crate::vulkan) use crate::vulkan::uniforms::CullHizParams;
use concinnity_render::uniforms::HizParams;

// 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_pipeline`).
pub(super) struct HiZResources {
    // Build pipelines + their layouts (the init and downsample kernels bind
    // different set layouts, so each needs its own pipeline layout).
    init_pipeline: OwnedPipeline,
    downsample_pipeline: OwnedPipeline,
    init_pipeline_layout: OwnedPipelineLayout,
    downsample_pipeline_layout: OwnedPipelineLayout,
    init_set_layout: OwnedSetLayout,
    downsample_set_layout: OwnedSetLayout,

    // Cull-read set layout (set 1 of the cull pipeline): sampler2D Hi-Z +
    // 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 init dst (mip 0) and
    // the downsample src/dst. Length = `mip_count`.
    mip_views: Vec<vk::ImageView>,
    // Nearest sampler the cull kernel reads the Hi-Z with (texelFetch ignores
    // filtering, but a sampler is still required for the combined-image-sampler
    // binding).
    sampler: OwnedSampler,

    // Build sets: one init set per frame (depth differs per frame slot), one
    // downsample set per mip step (frame-independent, Hi-Z mips only).
    init_sets: Vec<vk::DescriptorSet>,
    downsample_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 sampler 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 init 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,
) -> Result<PooledImage, String> {
    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| format!("hiz image: {e}"))
}

// 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,
) -> Result<vk::ImageView, String> {
    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| format!("hiz view: {e}"))
}

// Build the init/downsample 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,
    init_layout: vk::PipelineLayout,
    downsample_layout: vk::PipelineLayout,
    sample_count: u32,
    hot_reload: bool,
) -> Result<(OwnedPipeline, OwnedPipeline), String> {
    // The init kernel is a per-variant compile of the single-source
    // `hiz_build.slang`: 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 ctx = super::builtins::Ctx::plain(hot_reload);
    let init_spv = if sample_count > 1 {
        super::slang_builtins::HIZ_INIT_MSAA.compile(&ctx)?
    } else {
        super::slang_builtins::HIZ_INIT_SINGLE.compile(&ctx)?
    };
    let downsample_spv = super::slang_builtins::HIZ_DOWNSAMPLE.compile(&ctx)?;
    let init = create_compute_pipeline(device, init_layout, &init_spv)?;
    let downsample = create_compute_pipeline(device, downsample_layout, &downsample_spv)?;
    Ok((init, downsample))
}

fn create_compute_pipeline(
    device: &VkDevice,
    layout: vk::PipelineLayout,
    spv: &[u8],
) -> Result<OwnedPipeline, String> {
    let module = spv_module(device, spv)?;
    let entry = std::ffi::CString::new("main").unwrap();
    let stage = vk::PipelineShaderStageCreateInfo::default()
        .stage(vk::ShaderStageFlags::COMPUTE)
        .module(module.handle())
        .name(&entry);
    let info = vk::ComputePipelineCreateInfo::default()
        .stage(stage)
        .layout(layout);
    let pipeline = crate::vulkan::pipeline_cache::create_compute_pipeline(device, &info)
        .map_err(|e| format!("create hiz pipeline: {e}"))?;
    Ok(pipeline)
}

// 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 init kernel reduces into mip 0.
#[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 + its sampler, the two resources 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 sampling the
    // pyramid, and without taking a slot in this struct's pool. `read_set_layout`
    // is already `pub(super)`.
    pub(super) fn read_set_sources(&self) -> (vk::ImageView, vk::Sampler) {
        (self.sampled_view, self.sampler.handle())
    }

    // 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 init 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,
    ) -> Result<Self, String> {
        // `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.
        // Init: binding 0 depth (sampled image, read by coordinate -- no
        // sampler), binding 1 dst-mip storage image.
        let init_set_layout = create_set_layout(
            device,
            &[
                (0, vk::DescriptorType::SAMPLED_IMAGE),
                (1, vk::DescriptorType::STORAGE_IMAGE),
            ],
        )?;
        // Downsample: binding 0 src-mip storage image, binding 1 dst-mip.
        let downsample_set_layout = create_set_layout(
            device,
            &[
                (0, vk::DescriptorType::STORAGE_IMAGE),
                (1, vk::DescriptorType::STORAGE_IMAGE),
            ],
        )?;
        // Cull-read (set 1 of the cull pipeline): sampler2D Hi-Z + UBO.
        let read_set_layout = create_set_layout(
            device,
            &[
                (0, vk::DescriptorType::COMBINED_IMAGE_SAMPLER),
                (1, vk::DescriptorType::UNIFORM_BUFFER),
            ],
        )?;

        // 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::<HizParams>() as u32);
        let init_pipeline_layout =
            create_pipeline_layout(device, init_set_layout.handle(), push_range)?;
        let downsample_pipeline_layout =
            create_pipeline_layout(device, downsample_set_layout.handle(), push_range)?;

        let (init_pipeline, downsample_pipeline) = build_hiz_pipelines(
            device,
            init_pipeline_layout.handle(),
            downsample_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| -> Result<Vec<PooledBuffer>, String> {
            (0..count)
                .map(|_| {
                    alloc
                        .create_buffer(
                            ubo_size,
                            vk::BufferUsageFlags::UNIFORM_BUFFER,
                            vk::MemoryPropertyFlags::HOST_VISIBLE
                                | vk::MemoryPropertyFlags::HOST_COHERENT,
                        )
                        .map_err(String::from)
                })
                .collect()
        };
        let cull_ubos = alloc_ubo_ring(frames)?;
        let cull_ubos2 = alloc_ubo_ring(if two_pass { frames } else { 0 })?;

        let mut res = Self {
            init_pipeline,
            downsample_pipeline,
            init_pipeline_layout,
            downsample_pipeline_layout,
            init_set_layout,
            downsample_set_layout,
            read_set_layout,
            descriptor_pool,
            pyramid: PooledImage::null(),
            sampled_view: vk::ImageView::null(),
            mip_views: Vec::new(),
            sampler: create_sampler(device)?,
            init_sets: Vec::new(),
            downsample_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,
    ) -> Result<(), String> {
        let HiZDeviceCtx {
            alloc,
            device,
            command_pool,
            queue,
        } = ctx;
        let HiZTarget {
            width,
            height,
            depth_views,
        } = target;
        let mip_count = hiz_mip_count(width, height).min(MAX_HIZ_MIPS as u32).max(1);
        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 (init / downsample / 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| format!("reset hiz pool: {e}"))?;
        }
        let frames = self.cull_ubos.len();
        let init_layouts: Vec<_> = (0..frames).map(|_| self.init_set_layout.handle()).collect();
        let init_sets =
            alloc_descriptor_sets(device, self.descriptor_pool.handle(), &init_layouts)?;
        let downsample_layouts: Vec<_> = (1..mip_count)
            .map(|_| self.downsample_set_layout.handle())
            .collect();
        let downsample_sets =
            alloc_descriptor_sets(device, self.descriptor_pool.handle(), &downsample_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)?;

        // Init sets: binding 0 = that frame's main depth, binding 1 = mip 0.
        for (i, &set) in init_sets.iter().enumerate() {
            let depth = depth_views[i.min(depth_views.len().saturating_sub(1))];
            write_sampled_image(device, set, 0, depth);
            write_storage_image(device, set, 1, mip_views[0]);
        }
        // Downsample sets: step m reads mip m-1, writes mip m.
        for (step, &set) in downsample_sets.iter().enumerate() {
            let m = step + 1;
            write_storage_image(device, set, 0, mip_views[m - 1]);
            write_storage_image(device, set, 1, mip_views[m]);
        }
        // Read sets: binding 0 = all-mips Hi-Z sampler, binding 1 = cull UBO.
        for (i, &set) in read_sets.iter().enumerate() {
            write_sampler(device, set, 0, sampled_view, self.sampler.handle());
            write_uniform_buffer(
                device,
                set,
                1,
                self.cull_ubos[i].buffer(),
                std::mem::size_of::<CullHizParams>() as u64,
            );
        }
        // Phase-2 read sets: same pyramid sampler, the phase-2 per-frame UBO.
        for (i, &set) in read_sets2.iter().enumerate() {
            write_sampler(device, set, 0, sampled_view, self.sampler.handle());
            write_uniform_buffer(
                device,
                set,
                1,
                self.cull_ubos2[i].buffer(),
                std::mem::size_of::<CullHizParams>() as u64,
            );
        }

        // 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.init_sets = init_sets;
        self.downsample_sets = downsample_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, sampler, 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) -> Result<(), String> {
        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, init: OwnedPipeline, downsample: OwnedPipeline) {
        self.init_pipeline = init;
        self.downsample_pipeline = downsample;
    }

    // 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,
    ) -> Result<(OwnedPipeline, OwnedPipeline), String> {
        build_hiz_pipelines(
            device,
            self.init_pipeline_layout.handle(),
            self.downsample_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 per-mip chain below is this encoder's.
    pub(in crate::vulkan) fn encode_hiz_build(&self, cmd: vk::CommandBuffer, 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 device = &self.device;

        // 1. Init: mip 0 from main depth (MAX over MSAA samples when on).
        let init_params = HizParams {
            dst_width: hiz.width,
            dst_height: hiz.height,
            src_mip: 0,
            sample_count: hiz.sample_count.max(1),
        };
        // SAFETY: `cmd` is a command buffer in the recording state, and every handle and slice
        // these commands name is live for the call.
        unsafe {
            device.cmd_bind_pipeline(
                cmd,
                vk::PipelineBindPoint::COMPUTE,
                hiz.init_pipeline.handle(),
            );
            device.cmd_bind_descriptor_sets(
                cmd,
                vk::PipelineBindPoint::COMPUTE,
                hiz.init_pipeline_layout.handle(),
                0,
                std::slice::from_ref(&hiz.init_sets[frame_idx]),
                &[],
            );
            device.cmd_push_constants(
                cmd,
                hiz.init_pipeline_layout.handle(),
                vk::ShaderStageFlags::COMPUTE,
                0,
                as_bytes(&init_params),
            );
            device.cmd_dispatch(
                cmd,
                hiz.width.div_ceil(HIZ_TILE),
                hiz.height.div_ceil(HIZ_TILE),
                1,
            );
        }

        // 2. Downsample chain. Each step reads the prior mip and writes the
        //    next, with a compute write -> read barrier between dispatches.
        //    Finer than the graph's one-state-per-resource granularity, so this
        //    one stays inline.
        let mut cur_w = hiz.width;
        let mut cur_h = hiz.height;
        for mip in 1..hiz.mip_count {
            // SAFETY: `cmd` is a command buffer in the recording state, and every handle and slice
            // these commands name is live for the call.
            unsafe {
                device.cmd_pipeline_barrier(
                    cmd,
                    vk::PipelineStageFlags::COMPUTE_SHADER,
                    vk::PipelineStageFlags::COMPUTE_SHADER,
                    vk::DependencyFlags::empty(),
                    &[],
                    &[],
                    &[hiz_image_barrier(
                        hiz.pyramid.image(),
                        hiz.mip_count,
                        vk::ImageLayout::GENERAL,
                        vk::ImageLayout::GENERAL,
                        vk::AccessFlags::SHADER_WRITE,
                        vk::AccessFlags::SHADER_READ,
                    )],
                );
            }
            let next_w = (cur_w / 2).max(1);
            let next_h = (cur_h / 2).max(1);
            let params = HizParams {
                dst_width: next_w,
                dst_height: next_h,
                src_mip: mip - 1,
                sample_count: 0,
            };
            // SAFETY: `cmd` is a command buffer in the recording state, and every handle and slice
            // these commands name is live for the call.
            unsafe {
                device.cmd_bind_pipeline(
                    cmd,
                    vk::PipelineBindPoint::COMPUTE,
                    hiz.downsample_pipeline.handle(),
                );
                device.cmd_bind_descriptor_sets(
                    cmd,
                    vk::PipelineBindPoint::COMPUTE,
                    hiz.downsample_pipeline_layout.handle(),
                    0,
                    std::slice::from_ref(&hiz.downsample_sets[(mip - 1) as usize]),
                    &[],
                );
                device.cmd_push_constants(
                    cmd,
                    hiz.downsample_pipeline_layout.handle(),
                    vk::ShaderStageFlags::COMPUTE,
                    0,
                    as_bytes(&params),
                );
                device.cmd_dispatch(cmd, next_w.div_ceil(HIZ_TILE), next_h.div_ceil(HIZ_TILE), 1);
            }
            cur_w = next_w;
            cur_h = next_h;
        }
    }
}

fn as_bytes<T: bytemuck::NoUninit>(v: &T) -> &[u8] {
    bytemuck::bytes_of(v)
}

fn hiz_image_barrier(
    image: vk::Image,
    mip_count: u32,
    old: vk::ImageLayout,
    new: vk::ImageLayout,
    src: vk::AccessFlags,
    dst: vk::AccessFlags,
) -> vk::ImageMemoryBarrier<'static> {
    vk::ImageMemoryBarrier::default()
        .src_access_mask(src)
        .dst_access_mask(dst)
        .old_layout(old)
        .new_layout(new)
        .src_queue_family_index(vk::QUEUE_FAMILY_IGNORED)
        .dst_queue_family_index(vk::QUEUE_FAMILY_IGNORED)
        .image(image)
        .subresource_range(vk::ImageSubresourceRange {
            aspect_mask: vk::ImageAspectFlags::COLOR,
            base_mip_level: 0,
            level_count: mip_count,
            base_array_layer: 0,
            layer_count: 1,
        })
}

fn create_set_layout(
    device: &VkDevice,
    bindings: &[(u32, vk::DescriptorType)],
) -> Result<OwnedSetLayout, String> {
    let binds: Vec<_> = bindings
        .iter()
        .map(|&(b, ty)| {
            vk::DescriptorSetLayoutBinding::default()
                .binding(b)
                .descriptor_type(ty)
                .descriptor_count(1)
                .stage_flags(vk::ShaderStageFlags::COMPUTE)
        })
        .collect();
    device
        .create_descriptor_set_layout(
            &vk::DescriptorSetLayoutCreateInfo::default().bindings(&binds),
        )
        .map_err(|e| format!("hiz set layout: {e}"))
}

fn create_pipeline_layout(
    device: &VkDevice,
    set_layout: vk::DescriptorSetLayout,
    push_range: vk::PushConstantRange,
) -> Result<OwnedPipelineLayout, String> {
    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| format!("hiz pipeline layout: {e}"))
}

fn create_pool(
    device: &VkDevice,
    frames: usize,
    two_pass: bool,
) -> Result<OwnedDescriptorPool, String> {
    let f = frames as u32;
    // Two-pass occlusion adds one extra cull-read set per frame (phase 2),
    // each with a sampler + a UBO descriptor.
    let read_rings = if two_pass { 2 } else { 1 };
    let sizes = [
        // cull-read Hi-Z (frames per read ring).
        vk::DescriptorPoolSize::default()
            .ty(vk::DescriptorType::COMBINED_IMAGE_SAMPLER)
            .descriptor_count(read_rings * f),
        // init depth (frames): read by coordinate, no sampler.
        vk::DescriptorPoolSize::default()
            .ty(vk::DescriptorType::SAMPLED_IMAGE)
            .descriptor_count(f),
        // init dst (frames) + downsample src+dst (2 per step).
        vk::DescriptorPoolSize::default()
            .ty(vk::DescriptorType::STORAGE_IMAGE)
            .descriptor_count(f + 2 * MAX_HIZ_MIPS as u32),
        // cull-read UBO (frames per read ring).
        vk::DescriptorPoolSize::default()
            .ty(vk::DescriptorType::UNIFORM_BUFFER)
            .descriptor_count(read_rings * f),
    ];
    // init (frames) + cull-read (frames per read ring) + downsample (per mip).
    let max_sets = (1 + read_rings) * f + MAX_HIZ_MIPS as u32;
    device
        .create_descriptor_pool(
            &vk::DescriptorPoolCreateInfo::default()
                .pool_sizes(&sizes)
                .max_sets(max_sets),
        )
        .map_err(|e| format!("hiz descriptor pool: {e}"))
}

fn create_sampler(device: &VkDevice) -> Result<OwnedSampler, String> {
    let info = vk::SamplerCreateInfo::default()
        .mag_filter(vk::Filter::NEAREST)
        .min_filter(vk::Filter::NEAREST)
        .mipmap_mode(vk::SamplerMipmapMode::NEAREST)
        .address_mode_u(vk::SamplerAddressMode::CLAMP_TO_EDGE)
        .address_mode_v(vk::SamplerAddressMode::CLAMP_TO_EDGE)
        .address_mode_w(vk::SamplerAddressMode::CLAMP_TO_EDGE)
        .min_lod(0.0)
        .max_lod(MAX_HIZ_MIPS as f32);
    device
        .create_sampler(&info)
        .map_err(|e| format!("hiz sampler: {e}"))
}

fn write_sampler(
    device: &VkDevice,
    set: vk::DescriptorSet,
    binding: u32,
    view: vk::ImageView,
    sampler: vk::Sampler,
) {
    let info = vk::DescriptorImageInfo::default()
        .image_layout(vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL)
        .image_view(view)
        .sampler(sampler);
    let write = vk::WriteDescriptorSet::default()
        .dst_set(set)
        .dst_binding(binding)
        .descriptor_type(vk::DescriptorType::COMBINED_IMAGE_SAMPLER)
        .image_info(std::slice::from_ref(&info));
    // SAFETY: `writes` and the buffer/image infos it borrows are live for the call, and every set
    // and resource it names belongs to this device.
    unsafe { device.update_descriptor_sets(std::slice::from_ref(&write), &[]) };
}

// Sampled-image write with no sampler: the init kernel reads the depth by
// texel coordinate, so only the image view is bound.
fn write_sampled_image(
    device: &VkDevice,
    set: vk::DescriptorSet,
    binding: u32,
    view: vk::ImageView,
) {
    let info = vk::DescriptorImageInfo::default()
        .image_layout(vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL)
        .image_view(view);
    let write = vk::WriteDescriptorSet::default()
        .dst_set(set)
        .dst_binding(binding)
        .descriptor_type(vk::DescriptorType::SAMPLED_IMAGE)
        .image_info(std::slice::from_ref(&info));
    // SAFETY: `writes` and the buffer/image infos it borrows are live for the call, and every set
    // and resource it names belongs to this device.
    unsafe { device.update_descriptor_sets(std::slice::from_ref(&write), &[]) };
}

fn write_storage_image(
    device: &VkDevice,
    set: vk::DescriptorSet,
    binding: u32,
    view: vk::ImageView,
) {
    let info = vk::DescriptorImageInfo::default()
        .image_layout(vk::ImageLayout::GENERAL)
        .image_view(view);
    let write = vk::WriteDescriptorSet::default()
        .dst_set(set)
        .dst_binding(binding)
        .descriptor_type(vk::DescriptorType::STORAGE_IMAGE)
        .image_info(std::slice::from_ref(&info));
    // SAFETY: `writes` and the buffer/image infos it borrows are live for the call, and every set
    // and resource it names belongs to this device.
    unsafe { device.update_descriptor_sets(std::slice::from_ref(&write), &[]) };
}

fn write_uniform_buffer(
    device: &VkDevice,
    set: vk::DescriptorSet,
    binding: u32,
    buffer: vk::Buffer,
    range: u64,
) {
    let info = vk::DescriptorBufferInfo::default()
        .buffer(buffer)
        .offset(0)
        .range(range);
    let write = vk::WriteDescriptorSet::default()
        .dst_set(set)
        .dst_binding(binding)
        .descriptor_type(vk::DescriptorType::UNIFORM_BUFFER)
        .buffer_info(std::slice::from_ref(&info));
    // SAFETY: `writes` and the buffer/image infos it borrows are live for the call, and every set
    // and resource it names belongs to this device.
    unsafe { device.update_descriptor_sets(std::slice::from_ref(&write), &[]) };
}

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

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

    #[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);
    }
}