rlx_core/render/scenes/voxel/

mod.rs

//! The voxel system (ADR-0268): a 3D cellular automaton in a cube of cells,
//! drawn as a glowing, absorbing volume — or as lit solid cells (ADR-0269) —
//! seen through the shared orbit camera.
//!
//! # The state
//!
//! A ping-pong pair of 3D textures of side `grid`, one `u32` a cell packing its
//! state and its age (`shader::STEP_SHADER`). One generation is one compute
//! dispatch that counts each cell's live neighbours and applies the rule
//! ([`rules`]). The pair is built when a preset is configured, never on a frame.
//!
//! `[voxel] grid` is content, not a resolution: a structure is a fixed number of
//! cells, so the grid decides how large every structure looks. A grid past the
//! scene's cap is clamped and announced from `configure` as a
//! [`CapOverflow`](super::CapOverflow), never silently reduced.
//!
//! # Generations, not frames
//!
//! `step_rate` is in generations per second, integrated over the injected `dt`
//! by the shared `common::GenerationClock`, so the automaton
//! advances by the same count in the same wall time at any refresh rate.
//!
//! # Every cell is drawn from the seed
//!
//! The field is seeded from an integer hash of each cell's coordinates and a
//! seed derived from the preset's pinned salt (ADR-0051) — never from a clock —
//! and each step is integer counting, so a volume is a pure function of its
//! config and the sequence of bound values and `dt`s it was driven with. Each
//! `reseed` rise refills a ball whose centre and cells are hashed from the rise
//! count and the salt.
//!
//! # The rule list and the music
//!
//! `[voxel] rules` is a list of one to eight rules, and the structural `rule`
//! parameter indexes it, so a preset moves the rule on the bar through `[hold]`
//! without a mask ever passing through an `f32` (ADR-0268). The rule a frame
//! picks runs every generation that frame encodes. The spectrum lights the
//! cube's radial shells, bass at the centre, by `1 + shell_gain * level`: an
//! exact identity at zero gain, and the simulation never reads it.
//!
//! # The present
//!
//! A fullscreen march: each pixel's ray is walked cell by cell through the cube
//! `[-1, 1]^3` under emission-absorption (`shader::MARCH_SHADER`), and stops
//! once its transmittance is under one 8-bit step. The ray is recovered from the
//! camera's world-to-clip rows, whose aspect is the render target's (ADR-0037).
//! `focus` and `aperture` are inert here: the circle of confusion is per
//! primitive (ADR-0257), and a march has no primitives.
//!
//! The march draws into a target of its own — the tier's
//! [`voxel_march_scale`](crate::render::TierConfig::voxel_march_scale) of the
//! render target, held to its
//! [`voxel_march_cap`](crate::render::TierConfig::voxel_march_cap) — and a
//! plain stretch presents it. The target is a resolution, not a shape: the
//! march's aspect stays the render target's whatever the target's own is.
//!
//! # Solid
//!
//! Under `[voxel] present = "solid"` (ADR-0269) the same march stops at the
//! first live cell. The decay stages in front of it glow and absorb as they do
//! under `glow`; the live cell is shaded as a face — its normal the face the
//! ray entered through, exact from the traversal's last step — under one
//! directional key light and an ambient term:
//!
//! ```text
//! shade = palette(age) * brightness * shell * (ambient * occ + key * max(0, n . l))
//! ```
//!
//! composited under the light gathered so far, at full coverage. The light is
//! given in camera space (`light_dir`), so the lit side turns with an orbit.
//! `occ` is the face's voxel ambient occlusion: each of its four corners is
//! closed by the live cells beside it in the layer in front of the face — two
//! sides close it fully, otherwise each of the two sides and the diagonal a
//! third — and the four are interpolated across the face at the hit point and
//! mixed in by `ao`.
//!
//! # Bounce
//!
//! `bounce` (ADR-0272) lets a solid face take light from the live cells in
//! front of it. A light field of side `grid / LIGHT_CELL` holds, per texel, the
//! live fraction of the `LIGHT_CELL`³ block of cells it covers and that block's
//! summed age weighted the same way; a second texture holds it blurred once over
//! 3³ texels by a 1-2-1 tent on each axis. The march samples the blurred field
//! trilinearly half a cell in front of the face it hit, and adds
//!
//! ```text
//! bounce * palette(coord at L.age) * L.intensity * brightness * fog * shell
//! ```
//!
//! where `L.age` is the sampled age over the sampled intensity. The colour is
//! looked up at shade time, so only the intensity and the age wait for the
//! next refill. The field has no occlusion: light passes through walls.
//!
//! The field is refilled beside the brick pass, after every frame that changed
//! the state, and on the first frame `bounce` turns on after frames that did
//! not fill it. It exists only once a frame under `solid` has `bounce > 0`:
//! before that the march binds a one-texel stand-in and the uniform skips the
//! fetch, so a preset that never binds `bounce` builds and reads nothing.
//!
//! # The section plane
//!
//! `cut`, `cut_yaw` and `cut_pitch` (ADR-0271) hide every cell whose centre `c`,
//! in the cube's `[-1, 1]^3` frame, lies beyond one plane:
//!
//! ```text
//! dot(c, n) > sqrt(3) * (1 - 2 * cut)
//! n = (cos(cut_pitch) * sin(cut_yaw), sin(cut_pitch), cos(cut_pitch) * cos(cut_yaw))
//! ```
//!
//! so at `cut = 0.5` the plane passes through the centre and at 1 it removes
//! every cell. The normal is in the cube's frame, not the camera's, so the
//! section turns with the structure. A cut cell is dead to the march, under
//! either present, and to the solid present's occlusion, which reads it as
//! open. The automaton still runs it: only the view is cut. At `cut = 0` the
//! test is switched off by a flag, so every pixel is as it is without the
//! params.
//!
//! # Bricks
//!
//! After every frame that changed the state, a pass marks each 8³ brick of
//! cells that holds any non-dead cell. The march jumps a whole empty brick at
//! once, landing on exactly the cell, crossing parameters and position stepping
//! through it would have (the crossings are computed from the boundary index,
//! never accumulated), so skipping bricks moves no pixel. A brick is skipped
//! on its own empty flag alone: one holding cut cells is marched cell by cell,
//! and the cut is applied there.

// Hot-path panic-denial pragma (Plan 0002 Phase 2, extended to scenes by Plan
// 0003 Phase 0). Encodes its passes every displayed frame.
#![deny(
    clippy::unwrap_used,
    clippy::expect_used,
    clippy::indexing_slicing,
    clippy::panic,
    clippy::unreachable
)]

pub mod rules;
mod shader;

use super::common::{self, PaletteParams, PanParams};
use super::lines::GeneratorConfig;
use super::{FamilyParam, FamilyRange, ParamGroup, ParamKind, ParamSpec, Scene, default_of};
use crate::dsp::AnalysisFrame;
use crate::render::camera::{self, CameraParams};
use crate::render::gpu;
use crate::render::palette::{self, Palette};
use crate::render::tier::GridScale;
pub use rules::{DEFAULT_RULE, Neighbourhood, RosterRule, Rule, RuleList};
use std::f32::consts::{FRAC_PI_2, PI, TAU};

/// The grid an absent `[voxel] grid` means.
pub const DEFAULT_GRID: u32 = 64;
/// The smallest grid the loader accepts. Below it a seeded ball is a handful of
/// cells and every rule's structure is cut by the faces.
pub const MIN_GRID: u32 = 16;
/// The largest grid the loader accepts and any tier runs: 2.1 M cells, two
/// 4-byte texels each across the ping-pong pair, 16.8 MB.
pub const MAX_GRID: u32 = 128;

/// The fastest `step_rate` the clock integrates, in generations per second.
/// Well past the declared range's top, so it bounds a runaway binding rather
/// than an author.
pub const MAX_STEP_RATE: f32 = 120.0;

/// Where the age field stops counting: generations since a cell last changed
/// state, saturating here.
const AGE_CAP: u32 = 1023;
/// The generations over which a live cell's colour slides the whole of
/// `age_tint` along the palette.
const AGE_SPAN: f32 = 48.0;

/// The texel every cell is stored in: the state in the low byte and the age
/// above it, an integer format a compute pass can write.
const STATE_FORMAT: wgpu::TextureFormat = wgpu::TextureFormat::R32Uint;

/// Each side of the step pass's workgroup, in cells — `@workgroup_size(4, 4,
/// 4)` in `shader::STEP_SHADER`.
const WORKGROUP: u32 = 4;

/// Each side of a brick, in cells — `BRICK` in `shader::MARCH_SHADER` and the
/// brick pass's workgroup, which scans one brick.
const BRICK: u32 = 8;

/// Each side of the block of cells one light-field texel covers —
/// `LIGHT_CELL` in `shader::LIGHT_SPLAT_SHADER` and `shader::MARCH_SHADER`.
const LIGHT_CELL: u32 = 4;

/// The light field's texel: intensity in `r`, intensity-weighted age in `g`.
/// Filterable, so the march samples it trilinearly, and a storage format.
const LIGHT_FORMAT: wgpu::TextureFormat = wgpu::TextureFormat::Rgba16Float;

/// Each side of the light passes' workgroup, in texels — `@workgroup_size(4,
/// 4, 4)` in both light shaders.
const LIGHT_WORKGROUP: u32 = 4;

/// The largest `bounce` the march reads: well past the declared range, so it
/// bounds a runaway binding rather than an author.
const MAX_BOUNCE: f32 = 16.0;

/// `shader::MARCH_SHADER` with the constants it is written against prepended;
/// the vertex prelude goes in front of this.
fn march_body() -> String {
    format!(
        "const MAX_MARCH: i32 = {};\nconst AGE_SPAN: f32 = {:?};\nconst BRICK: i32 = {};\n\
         const LIGHT_CELL: f32 = {:?};\n{}",
        3 * MAX_GRID,
        AGE_SPAN,
        BRICK,
        LIGHT_CELL as f32,
        shader::MARCH_SHADER
    )
}

/// The step the march target's axes are rounded to (`grid::grid_size`).
const MARCH_STEP: u32 = 8;

/// The radius of the sphere bounding the cube `[-1, 1]^3`: the volume the
/// camera's `fog` is measured across.
const CUBE_RADIUS: f32 = 1.732_050_8;

/// Mixed into the preset's salt before it seeds the field, so a salt of `0`
/// still hashes to a field rather than to the hash's fixed point. The ASCII
/// bytes of "VOXL"; opaque, since changing it moves every seeded field.
const FIELD_SEED_MIX: u32 = 0x564F_584C;
/// Mixed into the salt for the stream of reseed balls, so a ball is not drawn
/// from the field's own sequence. The ASCII bytes of "RSVX".
const STAMP_SEED_MIX: u32 = 0x5253_5658;

/// `reseed` rises past this to refill one ball — edge-triggered, so a value
/// held high refills once, not once per frame.
const RESEED_THRESHOLD: f32 = 0.5;

/// The most radial shells the spectrum can be laid across: four `vec4`s of the
/// march's uniform.
pub const MAX_SHELLS: u32 = 16;

/// The largest `shell_gain` the march reads: well past the declared range, so
/// it bounds a runaway binding rather than an author.
const MAX_SHELL_GAIN: f32 = 16.0;

/// `[voxel] present` — how the march draws a live cell (ADR-0269).
#[derive(Debug, Clone, Copy, PartialEq, Eq, Default)]
pub enum Present {
    /// Every cell emits and absorbs; the ray runs on through the volume.
    #[default]
    Glow,
    /// The ray stops at the first live cell and shades its face under one
    /// camera-space key light and an ambient term. The decay stages in front
    /// of it still glow.
    Solid,
}

impl Present {
    /// Every present, in the order the reference lists them.
    pub const ALL: [Present; 2] = [Present::Glow, Present::Solid];

    /// The name a preset writes.
    pub fn as_str(self) -> &'static str {
        match self {
            Present::Glow => "glow",
            Present::Solid => "solid",
        }
    }

    /// The present `name` spells, or `None` for a name outside the roster.
    pub fn from_name(name: &str) -> Option<Self> {
        Self::ALL.into_iter().find(|p| p.as_str() == name)
    }
}

/// `[voxel]` — the structural configuration, fixed while the preset is loaded
/// and delivered through `Scene::configure`.
#[derive(Debug, Clone, Copy, PartialEq)]
pub struct VoxelConfig {
    /// Cells per side of the cube, validated at load into
    /// [`MIN_GRID`]`..=`[`MAX_GRID`].
    pub grid: u32,
    /// The rules a preset lists, compiled.
    pub rules: RuleList,
    /// The seeded ball's radius, as a fraction of the cube's half-side.
    pub seed_radius: f32,
    /// The probability a cell inside the seeded ball is live.
    pub seed_fill: f32,
    /// `true`: the cube is a 3-torus. `false`: every cell past a face is dead.
    pub wrap: bool,
    /// How many radial shells the spectrum lights, `0..=`[`MAX_SHELLS`]; `0`
    /// lights none.
    pub shells: u32,
    /// How a live cell is drawn: as glowing light, or as a lit solid face.
    pub present: Present,
    /// The preset's pinned salt (ADR-0051), which every seeded cell is drawn
    /// from.
    pub salt: u32,
}

impl Default for VoxelConfig {
    fn default() -> Self {
        let (seed_radius, seed_fill) = rules::DEFAULT_RULE.seed();
        Self {
            grid: DEFAULT_GRID,
            rules: RuleList::default(),
            seed_radius,
            seed_fill,
            // Dead outside: a structure crossing a face of a torus reappears on
            // the opposite one, which an orbiting camera reads as a cut.
            wrap: false,
            shells: 0,
            present: Present::Glow,
            salt: 0,
        }
    }
}

const DEFAULT_STEP_RATE: f32 = default_of(PARAMS, "step_rate");
const DEFAULT_DENSITY: f32 = default_of(PARAMS, "density");
const DEFAULT_TRAIL: f32 = default_of(PARAMS, "trail");
const DEFAULT_AGE_TINT: f32 = default_of(PARAMS, "age_tint");
const DEFAULT_HUE: f32 = default_of(PARAMS, "hue");
const DEFAULT_HUE_SPREAD: f32 = default_of(PARAMS, "hue_spread");
const DEFAULT_BRIGHTNESS: f32 = default_of(PARAMS, "brightness");
const DEFAULT_ZOOM: f32 = default_of(PARAMS, "zoom");
const DEFAULT_RESEED_RADIUS: f32 = default_of(PARAMS, "reseed_radius");
const DEFAULT_SHELL_GAIN: f32 = default_of(PARAMS, "shell_gain");
const DEFAULT_KEY: f32 = default_of(PARAMS, "key");
const DEFAULT_AMBIENT: f32 = default_of(PARAMS, "ambient");
const DEFAULT_AO: f32 = default_of(PARAMS, "ao");
const DEFAULT_BOUNCE: f32 = default_of(PARAMS, "bounce");
const DEFAULT_LIGHT_YAW: f32 = default_of(PARAMS, "light_yaw");
const DEFAULT_LIGHT_PITCH: f32 = default_of(PARAMS, "light_pitch");
const DEFAULT_CUT: f32 = default_of(PARAMS, "cut");
const DEFAULT_CUT_YAW: f32 = default_of(PARAMS, "cut_yaw");
const DEFAULT_CUT_PITCH: f32 = default_of(PARAMS, "cut_pitch");

// The shared camera block (ADR-0258), re-declared with this system's own doc
// lines where they differ: the editor schema keys a row by the whole
// declaration. Default, range and kind stay the shared block's.
const FOCUS: ParamSpec = ParamSpec {
    doc: "Inert on voxel: depth of field blurs each drawn primitive, and the march draws none.",
    ..camera::FOCUS
};
const APERTURE: ParamSpec = ParamSpec {
    doc: "Inert on voxel: depth of field blurs each drawn primitive, and the march draws none.",
    ..camera::APERTURE
};
const FOG: ParamSpec = ParamSpec {
    doc: "Fades the volume toward black with depth: at 1 the cube's farthest point is black and its nearest keeps its light. 0 is off.",
    ..camera::FOG
};

/// The parameter names this scene consumes — the vocabulary a preset binding is
/// checked against at load (ADR-0020). **Keep in sync with `set_param` below**;
/// `declared_params_match_set_param` in `core/tests/suite/preset.rs` fails if the
/// two drift.
pub const PARAMS: &[ParamSpec] = &[
    ParamSpec {
        name: "rule",
        default: 0.0,
        range: Some([0.0, (rules::MAX_RULES - 1) as f32]),
        doc: "Which of the [voxel] rules runs, counting from 0; past the list's end it runs \
              the last. A change takes effect at the next generation.",
        kind: ParamKind::Structural,
        group: ParamGroup::Shape,
        main: true,
    },
    ParamSpec {
        name: "step_rate",
        default: 8.0,
        range: Some([0.0, 30.0]),
        doc: "How many generations the automaton runs per second, whatever the frame rate; 0 \
              freezes it.",
        kind: ParamKind::Modal,
        group: ParamGroup::Motion,
        main: true,
    },
    ParamSpec {
        name: "reseed",
        default: 0.0,
        range: Some([0.0, 1.0]),
        doc: "A rise past 0.5 fills one ball of the cube with fresh seeded cells, once per \
              rise; bind an onset or a beat to it.",
        kind: ParamKind::Modal,
        group: ParamGroup::Motion,
        main: false,
    },
    ParamSpec {
        name: "reseed_radius",
        default: 0.25,
        range: Some([0.05, 1.0]),
        doc: "The radius of the ball a reseed fills, as a fraction of the cube's half-side.",
        kind: ParamKind::Modal,
        group: ParamGroup::Shape,
        main: false,
    },
    ParamSpec {
        name: "cut",
        default: 0.0,
        range: Some([0.0, 1.0]),
        doc: "How much of the cube a section plane hides, so the inside shows: 0.5 cuts it \
              through the centre, 1 hides it all, 0 is off. Only the view is cut; the \
              automaton runs on.",
        kind: ParamKind::Modal,
        group: ParamGroup::Shape,
        main: true,
    },
    ParamSpec {
        name: "cut_yaw",
        default: 0.0,
        range: Some([0.0, TAU]),
        doc: "Which way the section plane faces around the cube's vertical axis, in radians: \
              the side it hides. It turns with the cube, not the camera.",
        kind: ParamKind::Modal,
        group: ParamGroup::Shape,
        main: false,
    },
    ParamSpec {
        name: "cut_pitch",
        default: 0.0,
        range: Some([-FRAC_PI_2, FRAC_PI_2]),
        doc: "How far the section plane's hidden side tilts up toward the cube's top, in \
              radians; negative tilts it down.",
        kind: ParamKind::Modal,
        group: ParamGroup::Shape,
        main: false,
    },
    ParamSpec {
        name: "shell_gain",
        default: 0.0,
        range: Some([0.0, 4.0]),
        doc: "How strongly the spectrum lights the cube's radial shells, bass at the centre \
              and treble at the faces; 0 leaves the light as the cells give it.",
        kind: ParamKind::Modal,
        group: ParamGroup::Light,
        main: false,
    },
    ParamSpec {
        name: "density",
        default: 1.0,
        range: Some([0.0, 8.0]),
        doc: "How strongly the volume absorbs the light behind it, per unit of the cube's \
              width; 0 is pure additive glow with no front or back.",
        kind: ParamKind::Modal,
        group: ParamGroup::Light,
        main: true,
    },
    ParamSpec {
        name: "trail",
        default: 0.6,
        range: Some([0.0, 1.0]),
        doc: "How much of its light a decaying cell keeps at each stage of its decay; 0 draws \
              only the live cells.",
        kind: ParamKind::Modal,
        group: ParamGroup::Light,
        main: false,
    },
    ParamSpec {
        name: "age_tint",
        default: 0.3,
        range: Some([0.0, 1.0]),
        doc: "How far along the palette a cell's colour travels as it ages; a decaying cell \
              takes the far end.",
        kind: ParamKind::Modal,
        group: ParamGroup::Colour,
        main: false,
    },
    common::hue(0.0),
    ParamSpec {
        name: "hue_spread",
        default: 0.3,
        range: Some([0.0, 1.0]),
        doc: "How far along the palette the colour travels from the cube's centre to its \
              corners.",
        kind: ParamKind::Modal,
        group: ParamGroup::Colour,
        main: false,
    },
    common::brightness(1.0),
    ParamSpec {
        name: "key",
        default: 1.0,
        range: Some([0.0, 4.0]),
        doc: "How strongly the key light lights a solid face turned toward it; 0 leaves only \
              the ambient light.",
        kind: ParamKind::Modal,
        group: ParamGroup::Light,
        main: true,
    },
    ParamSpec {
        name: "ambient",
        default: 0.3,
        range: Some([0.0, 1.0]),
        doc: "The light every solid face gets whichever way it turns; 1 with no key light \
              draws each face in its own colour.",
        kind: ParamKind::Modal,
        group: ParamGroup::Light,
        main: false,
    },
    ParamSpec {
        name: "ao",
        default: 1.0,
        range: Some([0.0, 1.0]),
        doc: "How far the cells around a solid face darken its ambient light toward the \
              creases and inner corners; 0 is none.",
        kind: ParamKind::Modal,
        group: ParamGroup::Light,
        main: false,
    },
    ParamSpec {
        name: "bounce",
        default: 0.0,
        range: Some([0.0, 2.0]),
        doc: "How much light a solid face takes from the live cells in front of it, so a \
              structure glows onto itself; soft, reaching a few cells and through thin \
              walls. 0 is off.",
        kind: ParamKind::Modal,
        group: ParamGroup::Light,
        main: false,
    },
    ParamSpec {
        name: "light_yaw",
        default: -0.6,
        range: Some([-PI, PI]),
        doc: "Where the key light stands around the viewer, in radians: 0 behind the camera, \
              positive to its right. It turns with the camera.",
        kind: ParamKind::Modal,
        group: ParamGroup::Light,
        main: false,
    },
    ParamSpec {
        name: "light_pitch",
        default: 0.6,
        range: Some([-1.5, 1.5]),
        doc: "How far above the viewer the key light stands, in radians; negative lights from \
              below. It turns with the camera.",
        kind: ParamKind::Modal,
        group: ParamGroup::Light,
        main: false,
    },
    camera::YAW,
    camera::PITCH,
    camera::DISTANCE,
    camera::FOV,
    FOCUS,
    APERTURE,
    FOG,
    common::SATURATION,
    common::PALETTE_MIX,
    common::PALETTE_STEPS,
    common::zoom(1.0),
    common::PAN_X,
    common::PAN_Y,
];

/// One row of [`FAMILY_PARAMS`]: a parameter only `solid` reads, with its
/// range there, in [`Present::ALL`]'s order.
macro_rules! solid_only {
    ($name:literal: $range:expr) => {
        FamilyParam {
            name: $name,
            ranges: &[
                FamilyRange {
                    family: "glow",
                    range: None,
                },
                FamilyRange {
                    family: "solid",
                    range: $range,
                },
            ],
        }
    };
}

/// Every parameter only one present reads (ADR-0180 rule 4): the light, which
/// a glowing volume has no surface to take. A parameter missing from here
/// reads the same on both.
pub const FAMILY_PARAMS: &[FamilyParam] = &[
    solid_only!("key": Some([0.0, 4.0])),
    solid_only!("ambient": Some([0.0, 1.0])),
    solid_only!("ao": Some([0.0, 1.0])),
    solid_only!("bounce": Some([0.0, 2.0])),
    solid_only!("light_yaw": Some([-PI, PI])),
    solid_only!("light_pitch": Some([-1.5, 1.5])),
];

/// The seed the whole field is hashed from, for a preset whose salt is `salt`.
pub(crate) fn field_seed(salt: u32) -> u32 {
    salt ^ FIELD_SEED_MIX
}

/// A fill probability as the step shader compares it: the fraction of the
/// 24-bit hash range under which a cell is seeded live.
pub(crate) fn live_threshold(fill: f32) -> u32 {
    (fill.clamp(0.0, 1.0) * 16_777_216.0) as u32
}

/// The seeded ball's radius squared in half-cells², as the step shader compares
/// a cell's offset from the centre: `seed_radius` of the half-side is
/// `seed_radius * grid` half-cells.
pub(crate) fn seed_radius_sq(seed_radius: f32, grid: u32) -> u32 {
    let r = seed_radius.clamp(0.0, 2.0) * grid as f32;
    (r * r) as u32
}

/// One round of the lowbias32 mixer `gpu::HASH_WGSL` defines, on the CPU, for
/// the few values a reseed draws per rise.
fn mix32(v: u32) -> u32 {
    let mut h = v;
    h ^= h >> 16;
    h = h.wrapping_mul(0x7FEB_352D);
    h ^= h >> 15;
    h = h.wrapping_mul(0x846C_A68B);
    h ^= h >> 16;
    h
}

/// One scheduled `reseed`: the ball's centre in cells, its radius squared in
/// cells², and the seed its cells are hashed from.
#[derive(Debug, Clone, Copy, PartialEq, Eq)]
pub(crate) struct Stamp {
    pub(crate) centre: [u32; 3],
    pub(crate) radius_sq: u32,
    pub(crate) seed: u32,
}

/// The ball the `rise`-th reseed of a preset salted `salt` fills on a grid of
/// `grid`, at `radius` of the half-side. Hashed from the rise count and the
/// salt, never from a clock, so a preset replays its balls identically.
pub(crate) fn stamp(salt: u32, rise: u32, grid: u32, radius: f32) -> Stamp {
    let base = mix32(rise ^ mix32(salt ^ STAMP_SEED_MIX));
    let side = grid.max(1);
    let r = radius.clamp(0.0, 2.0) * side as f32 * 0.5;
    Stamp {
        centre: [
            mix32(base ^ 1) % side,
            mix32(base ^ 2) % side,
            mix32(base ^ 3) % side,
        ],
        radius_sq: (r * r) as u32,
        seed: mix32(base ^ 4),
    }
}

/// `spectrum` laid across `shells` radial bands, bass at the centre: band `i`
/// is the mean of its share of the bins, held finite and non-negative. Bands
/// past `shells` stay zero.
pub(crate) fn shell_levels(spectrum: &[f32], shells: u32) -> [f32; MAX_SHELLS as usize] {
    let mut out = [0.0; MAX_SHELLS as usize];
    let s = (shells.min(MAX_SHELLS) as usize).min(spectrum.len());
    if s == 0 {
        return out;
    }
    for (i, level) in out.iter_mut().take(s).enumerate() {
        let lo = i * spectrum.len() / s;
        let hi = ((i + 1) * spectrum.len() / s).max(lo + 1);
        let band = spectrum.get(lo..hi).unwrap_or(&[]);
        let mean = band.iter().sum::<f32>() / band.len().max(1) as f32;
        *level = if mean.is_finite() { mean.max(0.0) } else { 0.0 };
    }
    out
}

/// A bound `step_rate` as the clock integrates it: clamped into
/// `0..=`[`MAX_STEP_RATE`], the declared default where it is not finite.
pub(crate) fn applied_step_rate(value: f32) -> f32 {
    if value.is_finite() {
        value.clamp(0.0, MAX_STEP_RATE)
    } else {
        DEFAULT_STEP_RATE
    }
}

/// `value` where it is finite, `fallback` where it is not.
fn finite_or(value: f32, fallback: f32) -> f32 {
    if value.is_finite() { value } else { fallback }
}

/// `config` with its grid held to `cap`, and the clamp to report if that moved
/// it. A grid is structural, so this runs once, at `configure`.
pub(crate) fn grid_clamp(
    config: VoxelConfig,
    cap: u32,
) -> (VoxelConfig, Option<super::CapOverflow>) {
    if config.grid <= cap {
        return (config, None);
    }
    let overflow = super::CapOverflow {
        dropped: (config.grid - cap) as usize,
        context: super::OverflowContext::Voxels(config.grid),
        cap: cap as usize,
    };
    (
        VoxelConfig {
            grid: cap,
            ..config
        },
        Some(overflow),
    )
}

/// The three non-zero rows of `view`'s world-to-clip matrix — screen x, screen
/// y and view depth — each as `[x, y, z, translation]`. The matrix is stored
/// column-major, and its third row is zero (no depth attachment, ADR-0257).
fn rows(view: &camera::CameraView) -> [[f32; 4]; 3] {
    let m = view.view_proj;
    let row = |r: usize| {
        let at = |c: usize| m.get(c).and_then(|col| col.get(r)).copied().unwrap_or(0.0);
        [at(0), at(1), at(2), at(3)]
    };
    [row(0), row(1), row(3)]
}

/// The eye: the one point every non-zero row of `view` maps to zero, solved by
/// Cramer's rule. The three rows are the camera's right, up and view axes,
/// scaled and sheared by the lens and the pan, so they are never coplanar.
fn eye(view: &camera::CameraView) -> [f32; 3] {
    let [r0, r1, r3] = rows(view);
    let (a, b, c) = (xyz(r0), xyz(r1), xyz(r3));
    let rhs = [-r0[3], -r1[3], -r3[3]];
    let det = dot(a, cross(b, c));
    if det.abs() < 1e-12 {
        return [0.0, 0.0, view.distance];
    }
    // Column `i` of the system replaced by the right-hand side.
    let solve = |i: usize| {
        let swap = |row: [f32; 3], k: usize| {
            let mut r = row;
            if let Some(slot) = r.get_mut(i) {
                *slot = rhs.get(k).copied().unwrap_or(0.0);
            }
            r
        };
        dot(swap(a, 0), cross(swap(b, 1), swap(c, 2))) / det
    };
    [solve(0), solve(1), solve(2)]
}

/// The unit world direction toward a key light standing at `light_yaw` around
/// the viewer and `light_pitch` above them, in `view`'s camera space: at `(0,
/// 0)` it shines from the eye along the view axis. The camera's basis is taken
/// from the rows, so the light turns with every orbit, and the lens's scale and
/// the pan's shear are removed from them first.
pub(crate) fn light_dir(view: &camera::CameraView, light_yaw: f32, light_pitch: f32) -> [f32; 3] {
    let [r0, _, r3] = rows(view);
    let normalize = |v: [f32; 3]| {
        let len = dot(v, v).sqrt();
        if len > 1e-12 {
            [v[0] / len, v[1] / len, v[2] / len]
        } else {
            v
        }
    };
    let forward = normalize(xyz(r3));
    // Screen x is `right * sx + forward * pan`: drop the forward share.
    let r = xyz(r0);
    let along = dot(r, forward);
    let right = normalize([
        r[0] - along * forward[0],
        r[1] - along * forward[1],
        r[2] - along * forward[2],
    ]);
    let up = cross(right, forward);
    let (sy, cy) = finite_or(light_yaw, DEFAULT_LIGHT_YAW).sin_cos();
    let (sp, cp) = finite_or(light_pitch, DEFAULT_LIGHT_PITCH)
        .clamp(-FRAC_PI_2, FRAC_PI_2)
        .sin_cos();
    let (x, y, z) = (sy * cp, sp, cy * cp);
    // `z` points back at the eye, against the view axis.
    normalize([
        right[0] * x + up[0] * y - forward[0] * z,
        right[1] * x + up[1] * y - forward[1] * z,
        right[2] * x + up[2] * y - forward[2] * z,
    ])
}

/// The section plane as the march reads it: `n.xyz` the unit normal toward
/// the hidden side in the cube's frame, and `w` the plane's offset `d`, so a
/// cell whose centre `c` has `dot(c, n) > d` is cut. `None` at `cut <= 0`,
/// which the march takes as no plane at all. A `cut` past 1 is held there.
pub(crate) fn cut_plane(cut: f32, cut_yaw: f32, cut_pitch: f32) -> Option<[f32; 4]> {
    let cut = finite_or(cut, DEFAULT_CUT).min(1.0);
    if cut <= 0.0 {
        return None;
    }
    let (sy, cy) = finite_or(cut_yaw, DEFAULT_CUT_YAW).sin_cos();
    let (sp, cp) = finite_or(cut_pitch, DEFAULT_CUT_PITCH)
        .clamp(-FRAC_PI_2, FRAC_PI_2)
        .sin_cos();
    Some([cp * sy, sp, cp * cy, CUBE_RADIUS * (1.0 - 2.0 * cut)])
}

fn xyz(v: [f32; 4]) -> [f32; 3] {
    [v[0], v[1], v[2]]
}

fn dot(a: [f32; 3], b: [f32; 3]) -> f32 {
    a[0] * b[0] + a[1] * b[1] + a[2] * b[2]
}

fn cross(a: [f32; 3], b: [f32; 3]) -> [f32; 3] {
    [
        a[1] * b[2] - a[2] * b[1],
        a[2] * b[0] - a[0] * b[2],
        a[0] * b[1] - a[1] * b[0],
    ]
}

/// The one uniform the seed and step passes read, written once a frame.
///
/// **One buffer for both, and that is load-bearing**, for `cellular`'s reason:
/// which pass is running is a constant compiled into its pipeline, never a
/// field here, so every pass is bound to identical resources and an adapter
/// cannot hand one pass another's uniform (ADR-0058).
#[repr(C)]
#[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
struct StepParams {
    /// x: grid, y: wrap (1 torus), z: the field's seed, w: live threshold.
    a: [u32; 4],
    /// x: birth mask, y: survive mask, z: states, w: neighbourhood.
    b: [u32; 4],
    /// x: the seed ball's radius squared in half-cells², y: the reseed ball's
    /// seed; zw unused.
    c: [u32; 4],
    /// xyz: the reseed ball's centre (cells), w: its radius squared (cells²).
    d: [u32; 4],
}

/// The march's uniform, written once a frame.
#[repr(C)]
#[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
struct MarchParams {
    r0: [f32; 4],
    r1: [f32; 4],
    r3: [f32; 4],
    /// xyz: the eye; w: grid.
    eye: [f32; 4],
    /// x: hue, y: brightness, z: saturation, w: palette_mix.
    a: [f32; 4],
    /// x: palette_steps, y: occlude, z: density, w: trail.
    b: [f32; 4],
    /// x: age_tint, y: hue_spread, z: fog, w: shell_gain.
    c: [f32; 4],
    /// x: the cube's nearest view depth, y: its depth span, z: shells, w: 1 to
    /// jump empty bricks.
    d: [f32; 4],
    /// xyz: the unit world direction toward the key light; w: 1 under
    /// `solid`.
    light: [f32; 4],
    /// x: key, y: ambient, z: ao (0..1); w: 1 when the section plane is on.
    lit: [f32; 4],
    /// The section plane ([`cut_plane`]): xyz its unit normal, w its offset.
    /// Zero, and never read, while `lit.w` is 0.
    cut: [f32; 4],
    /// x: the bounce the march adds, 0 to skip the light field's fetch — always
    /// 0 under `glow`; yzw unused.
    bounce: [f32; 4],
    /// The shells' levels, four to a `vec4`, bass first.
    shells: [[f32; 4]; (MAX_SHELLS / 4) as usize],
}

/// The march target for a render target of `target` pixels: `scale` of it on
/// each axis, held to `cap` by one factor and rounded to [`MARCH_STEP`]
/// (`grid::grid_size`, ADR-0037). `scale` is the tier's march scale times the
/// renderer's own grid scale, floored at the smallest scale a grid takes.
pub(crate) fn march_size(target: (u32, u32), scale: f32, cap: (u32, u32)) -> (u32, u32) {
    let scale =
        GridScale::new(scale.clamp(GridScale::MIN, GridScale::MAX)).unwrap_or(GridScale::FULL);
    crate::render::grid::grid_size(target, scale, cap, MARCH_STEP)
}

/// What one pass of the step shader does: its `MODE` constant, compiled in.
#[derive(Clone, Copy)]
enum StepPass {
    /// Advance every cell one generation by the rule.
    Step,
    /// Hash every cell from the field's seed.
    Seed,
    /// Hash the cells inside the scheduled reseed ball from its own seed.
    Stamp,
}

impl StepPass {
    fn mode(self) -> u32 {
        match self {
            StepPass::Step => 0,
            StepPass::Seed => 1,
            StepPass::Stamp => 2,
        }
    }

    fn label(self) -> &'static str {
        match self {
            StepPass::Step => "voxel-step",
            StepPass::Seed => "voxel-seed",
            StepPass::Stamp => "voxel-stamp",
        }
    }

    /// This pass's shader: its constants, the hash and the body. Built at
    /// construction only.
    fn source(self) -> String {
        format!(
            "const MODE: u32 = {}u;\nconst AGE_CAP: u32 = {}u;\n{}{}{}",
            self.mode(),
            AGE_CAP,
            gpu::HASH_WGSL,
            common::CELL_HASH_WGSL,
            shader::STEP_SHADER
        )
    }
}

/// One of the pair of state textures, with the two views the passes take of it.
struct StateTexture {
    /// Held beside its views; read back by the probes.
    #[cfg_attr(not(test), allow(dead_code))]
    texture: wgpu::Texture,
    /// Sampled by `textureLoad`, as the step's source and the march's field.
    read: wgpu::TextureView,
    /// Written by the step pass.
    write: wgpu::TextureView,
}

impl StateTexture {
    fn new(device: &wgpu::Device, grid: u32, label: &str) -> Self {
        let texture = device.create_texture(&wgpu::TextureDescriptor {
            label: Some(label),
            size: wgpu::Extent3d {
                width: grid,
                height: grid,
                depth_or_array_layers: grid,
            },
            mip_level_count: 1,
            sample_count: 1,
            dimension: wgpu::TextureDimension::D3,
            format: STATE_FORMAT,
            // The copies only so a probe can read the cells back and write a
            // structure in.
            usage: wgpu::TextureUsages::TEXTURE_BINDING
                | wgpu::TextureUsages::STORAGE_BINDING
                | wgpu::TextureUsages::COPY_SRC
                | wgpu::TextureUsages::COPY_DST,
            view_formats: &[],
        });
        let read = texture.create_view(&wgpu::TextureViewDescriptor::default());
        let write = texture.create_view(&wgpu::TextureViewDescriptor::default());
        Self {
            texture,
            read,
            write,
        }
    }
}

/// The light field (ADR-0272): the splat texture each refill reduces the cells
/// into, the blurred texture the march samples, and the two passes between
/// them. Built on the first frame that reads it, never at `configure`.
struct LightField {
    /// Held beside the views the passes take of them.
    _splat: wgpu::Texture,
    _blurred: wgpu::Texture,
    /// The blurred field, sampled by the march.
    blurred_view: wgpu::TextureView,
    splat_pipeline: wgpu::ComputePipeline,
    blur_pipeline: wgpu::ComputePipeline,
    /// Splat the cells of A, and of B.
    splat_from_a: wgpu::BindGroup,
    splat_from_b: wgpu::BindGroup,
    blur_bg: wgpu::BindGroup,
    /// Texels a side.
    side: u32,
}

impl LightField {
    fn build(device: &wgpu::Device, grid: u32, a: &StateTexture, b: &StateTexture) -> Self {
        let side = grid.div_ceil(LIGHT_CELL);
        let field = |label: &str| {
            device.create_texture(&wgpu::TextureDescriptor {
                label: Some(label),
                size: wgpu::Extent3d {
                    width: side,
                    height: side,
                    depth_or_array_layers: side,
                },
                mip_level_count: 1,
                sample_count: 1,
                dimension: wgpu::TextureDimension::D3,
                format: LIGHT_FORMAT,
                usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::STORAGE_BINDING,
                view_formats: &[],
            })
        };
        let splat = field("voxel-light-splat");
        let blurred = field("voxel-light-blurred");
        let splat_read = splat.create_view(&wgpu::TextureViewDescriptor::default());
        let splat_write = splat.create_view(&wgpu::TextureViewDescriptor::default());
        let blurred_write = blurred.create_view(&wgpu::TextureViewDescriptor::default());
        let blurred_view = blurred.create_view(&wgpu::TextureViewDescriptor::default());

        // These two layouts and `voxel-brick-layout` all bind one 3D texture and
        // one 3D storage texture to compute, and must still be three distinct
        // shapes (ADR-0058; `tonemap/tests.rs` asserts it). Hence the order and
        // the source bound twice: a tidier arrangement re-collides a pair. Every
        // entry is a full literal, because that scan reads only the spellings
        // its `MARKERS` lists and a local helper is not one of them.
        // `[StorageTexture<rgba16float>, Texture<u32>]`: the splat, the cells.
        let splat_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
            label: Some("voxel-light-splat-layout"),
            entries: &[
                wgpu::BindGroupLayoutEntry {
                    binding: 0,
                    visibility: wgpu::ShaderStages::COMPUTE,
                    ty: wgpu::BindingType::StorageTexture {
                        access: wgpu::StorageTextureAccess::WriteOnly,
                        format: LIGHT_FORMAT,
                        view_dimension: wgpu::TextureViewDimension::D3,
                    },
                    count: None,
                },
                wgpu::BindGroupLayoutEntry {
                    binding: 1,
                    visibility: wgpu::ShaderStages::COMPUTE,
                    ty: wgpu::BindingType::Texture {
                        sample_type: wgpu::TextureSampleType::Uint,
                        view_dimension: wgpu::TextureViewDimension::D3,
                        multisampled: false,
                    },
                    count: None,
                },
            ],
        });
        // `[Texture<f32>, StorageTexture<rgba16float>, Texture<f32>]`: the
        // splat, the blurred, and the splat again, which the shader never reads.
        let blur_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
            label: Some("voxel-light-blur-layout"),
            entries: &[
                wgpu::BindGroupLayoutEntry {
                    binding: 0,
                    visibility: wgpu::ShaderStages::COMPUTE,
                    ty: wgpu::BindingType::Texture {
                        sample_type: wgpu::TextureSampleType::Float { filterable: false },
                        view_dimension: wgpu::TextureViewDimension::D3,
                        multisampled: false,
                    },
                    count: None,
                },
                wgpu::BindGroupLayoutEntry {
                    binding: 1,
                    visibility: wgpu::ShaderStages::COMPUTE,
                    ty: wgpu::BindingType::StorageTexture {
                        access: wgpu::StorageTextureAccess::WriteOnly,
                        format: LIGHT_FORMAT,
                        view_dimension: wgpu::TextureViewDimension::D3,
                    },
                    count: None,
                },
                wgpu::BindGroupLayoutEntry {
                    binding: 2,
                    visibility: wgpu::ShaderStages::COMPUTE,
                    ty: wgpu::BindingType::Texture {
                        sample_type: wgpu::TextureSampleType::Float { filterable: false },
                        view_dimension: wgpu::TextureViewDimension::D3,
                        multisampled: false,
                    },
                    count: None,
                },
            ],
        });
        // `views[i]` is bound at binding `i`.
        let bind = |layout: &wgpu::BindGroupLayout, views: &[&wgpu::TextureView]| {
            let entries: Vec<wgpu::BindGroupEntry<'_>> = (0u32..)
                .zip(views)
                .map(|(binding, view)| wgpu::BindGroupEntry {
                    binding,
                    resource: wgpu::BindingResource::TextureView(view),
                })
                .collect();
            device.create_bind_group(&wgpu::BindGroupDescriptor {
                label: Some("voxel-light-bg"),
                layout,
                entries: &entries,
            })
        };
        let splat_from_a = bind(&splat_layout, &[&splat_write, &a.read]);
        let splat_from_b = bind(&splat_layout, &[&splat_write, &b.read]);
        let blur_bg = bind(&blur_layout, &[&splat_read, &blurred_write, &splat_read]);
        let pipeline = |label: &str, layout: &wgpu::BindGroupLayout, source: String| {
            let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
                label: Some(label),
                source: wgpu::ShaderSource::Wgsl(source.into()),
            });
            let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
                label: Some(label),
                bind_group_layouts: &[Some(layout)],
                immediate_size: 0,
            });
            device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
                label: Some(label),
                layout: Some(&pipeline_layout),
                module: &module,
                entry_point: Some("main"),
                compilation_options: Default::default(),
                cache: None,
            })
        };
        let splat_pipeline = pipeline(
            "voxel-light-splat",
            &splat_layout,
            format!(
                "const LIGHT_CELL: u32 = {LIGHT_CELL}u;\nconst AGE_SPAN: f32 = {AGE_SPAN:?};\n{}",
                shader::LIGHT_SPLAT_SHADER
            ),
        );
        let blur_pipeline = pipeline(
            "voxel-light-blur",
            &blur_layout,
            shader::LIGHT_BLUR_SHADER.to_string(),
        );
        Self {
            _splat: splat,
            _blurred: blurred,
            blurred_view,
            splat_pipeline,
            blur_pipeline,
            splat_from_a,
            splat_from_b,
            blur_bg,
            side,
        }
    }

    /// Refill the field from the generation `reading_a` names: the splat, then
    /// the blur, each its own pass so the blur reads a complete splat.
    fn encode(&self, encoder: &mut wgpu::CommandEncoder, reading_a: bool) {
        let groups = self.side.div_ceil(LIGHT_WORKGROUP);
        {
            let mut cpass = gpu::compute_pass(encoder, "voxel-light-splat");
            cpass.set_pipeline(&self.splat_pipeline);
            cpass.set_bind_group(
                0,
                if reading_a {
                    &self.splat_from_a
                } else {
                    &self.splat_from_b
                },
                &[],
            );
            cpass.dispatch_workgroups(groups, groups, groups);
        }
        let mut cpass = gpu::compute_pass(encoder, "voxel-light-blur");
        cpass.set_pipeline(&self.blur_pipeline);
        cpass.set_bind_group(0, &self.blur_bg, &[]);
        cpass.dispatch_workgroups(groups, groups, groups);
    }
}

/// The march's own target and the bind group that presents it, rebuilt when
/// the size the frame asks for changes.
struct MarchTarget {
    size: (u32, u32),
    /// Held beside its view.
    _texture: wgpu::Texture,
    view: wgpu::TextureView,
    present_bg: wgpu::BindGroup,
}

/// The GPU-side state, built when a preset is configured and rebuilt when one
/// asks for a different grid.
struct Resources {
    /// The grid these textures were built for.
    grid: u32,
    /// The pair, kept so the bind groups' views stay valid.
    #[cfg_attr(not(test), allow(dead_code))]
    a: StateTexture,
    #[cfg_attr(not(test), allow(dead_code))]
    b: StateTexture,
    /// Whether `a` holds the current generation.
    reading_a: bool,
    seed_pipeline: wgpu::ComputePipeline,
    stamp_pipeline: wgpu::ComputePipeline,
    step_pipeline: wgpu::ComputePipeline,
    step_uniform: wgpu::Buffer,
    /// Read A write B, and read B write A.
    step_from_a: wgpu::BindGroup,
    step_from_b: wgpu::BindGroup,
    /// One texel a brick: non-zero where the brick holds a non-dead cell.
    /// Held beside the views the passes take of it.
    _bricks: wgpu::Texture,
    /// The march's view of the bricks, kept so its bind groups can be rebuilt
    /// when the light field arrives.
    bricks_read: wgpu::TextureView,
    brick_pipeline: wgpu::ComputePipeline,
    /// Mark the bricks of A, and of B.
    brick_from_a: wgpu::BindGroup,
    brick_from_b: wgpu::BindGroup,
    march_layout: wgpu::BindGroupLayout,
    march_pipeline: wgpu::RenderPipeline,
    march_uniform: wgpu::Buffer,
    /// Each bound to the light field once it exists, to `light_stand_in`
    /// before.
    march_from_a: wgpu::BindGroup,
    march_from_b: wgpu::BindGroup,
    /// A one-texel field the march binds until the light field is built; the
    /// uniform never lets it be read. Held beside the bind groups' view of it.
    _light_stand_in: wgpu::Texture,
    /// Clamps at the cube's faces and filters trilinearly.
    light_sampler: wgpu::Sampler,
    /// Built on the first frame under `solid` with `bounce > 0`.
    light: Option<LightField>,
    /// Whether `light` holds the current generation. Cleared by every frame
    /// that changes the state without refilling it.
    light_fresh: bool,
    present_layout: wgpu::BindGroupLayout,
    present_pipeline: wgpu::RenderPipeline,
    present_sampler: wgpu::Sampler,
    /// Built by the first `render`, at the size that frame asks for.
    target: Option<MarchTarget>,
    /// The shared gradient LUT pair (ADR-0021). A fresh pair is dirty, so a
    /// (re)build uploads on its first frame.
    luts: palette::LutPair,
}

impl Resources {
    fn build(device: &wgpu::Device, surface_format: wgpu::TextureFormat, grid: u32) -> Self {
        let a = StateTexture::new(device, grid, "voxel-state-a");
        let b = StateTexture::new(device, grid, "voxel-state-b");
        let side = grid.div_ceil(BRICK);
        let bricks = device.create_texture(&wgpu::TextureDescriptor {
            label: Some("voxel-bricks"),
            size: wgpu::Extent3d {
                width: side,
                height: side,
                depth_or_array_layers: side,
            },
            mip_level_count: 1,
            sample_count: 1,
            dimension: wgpu::TextureDimension::D3,
            format: STATE_FORMAT,
            usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::STORAGE_BINDING,
            view_formats: &[],
        });
        let bricks_read = bricks.create_view(&wgpu::TextureViewDescriptor::default());
        let bricks_write = bricks.create_view(&wgpu::TextureViewDescriptor::default());
        let step_uniform = gpu::uniform_buffer(
            device,
            "voxel-step-params",
            std::mem::size_of::<StepParams>(),
        );
        let march_uniform = gpu::uniform_buffer(
            device,
            "voxel-march-params",
            std::mem::size_of::<MarchParams>(),
        );

        // `[Texture, StorageTexture, Uniform]`, all compute-visible: the source
        // generation, the one being written, the uniform.
        let step_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
            label: Some("voxel-step-layout"),
            entries: &[
                wgpu::BindGroupLayoutEntry {
                    binding: 0,
                    visibility: wgpu::ShaderStages::COMPUTE,
                    ty: wgpu::BindingType::Texture {
                        sample_type: wgpu::TextureSampleType::Uint,
                        view_dimension: wgpu::TextureViewDimension::D3,
                        multisampled: false,
                    },
                    count: None,
                },
                wgpu::BindGroupLayoutEntry {
                    binding: 1,
                    visibility: wgpu::ShaderStages::COMPUTE,
                    ty: wgpu::BindingType::StorageTexture {
                        access: wgpu::StorageTextureAccess::WriteOnly,
                        format: STATE_FORMAT,
                        view_dimension: wgpu::TextureViewDimension::D3,
                    },
                    count: None,
                },
                gpu::uniform(2, wgpu::ShaderStages::COMPUTE),
            ],
        });
        let step_bind = |src: &StateTexture, dst: &StateTexture| {
            device.create_bind_group(&wgpu::BindGroupDescriptor {
                label: Some("voxel-step-bg"),
                layout: &step_layout,
                entries: &[
                    wgpu::BindGroupEntry {
                        binding: 0,
                        resource: wgpu::BindingResource::TextureView(&src.read),
                    },
                    wgpu::BindGroupEntry {
                        binding: 1,
                        resource: wgpu::BindingResource::TextureView(&dst.write),
                    },
                    wgpu::BindGroupEntry {
                        binding: 2,
                        resource: step_uniform.as_entire_binding(),
                    },
                ],
            })
        };
        let step_from_a = step_bind(&a, &b);
        let step_from_b = step_bind(&b, &a);
        let step_pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
            label: Some("voxel-step-pipeline-layout"),
            bind_group_layouts: &[Some(&step_layout)],
            immediate_size: 0,
        });
        let compute = |pass: StepPass| {
            let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
                label: Some(pass.label()),
                source: wgpu::ShaderSource::Wgsl(pass.source().into()),
            });
            device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
                label: Some(pass.label()),
                layout: Some(&step_pipeline_layout),
                module: &module,
                entry_point: Some("main"),
                compilation_options: Default::default(),
                cache: None,
            })
        };
        let seed_pipeline = compute(StepPass::Seed);
        let stamp_pipeline = compute(StepPass::Stamp);
        let step_pipeline = compute(StepPass::Step);

        // `[Texture, StorageTexture]`, compute-visible: the generation to mark,
        // and the bricks. Its size comes from the texture, so it takes no
        // uniform.
        let brick_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
            label: Some("voxel-brick-layout"),
            entries: &[
                wgpu::BindGroupLayoutEntry {
                    binding: 0,
                    visibility: wgpu::ShaderStages::COMPUTE,
                    ty: wgpu::BindingType::Texture {
                        sample_type: wgpu::TextureSampleType::Uint,
                        view_dimension: wgpu::TextureViewDimension::D3,
                        multisampled: false,
                    },
                    count: None,
                },
                wgpu::BindGroupLayoutEntry {
                    binding: 1,
                    visibility: wgpu::ShaderStages::COMPUTE,
                    ty: wgpu::BindingType::StorageTexture {
                        access: wgpu::StorageTextureAccess::WriteOnly,
                        format: STATE_FORMAT,
                        view_dimension: wgpu::TextureViewDimension::D3,
                    },
                    count: None,
                },
            ],
        });
        let brick_bind = |src: &StateTexture| {
            device.create_bind_group(&wgpu::BindGroupDescriptor {
                label: Some("voxel-brick-bg"),
                layout: &brick_layout,
                entries: &[
                    wgpu::BindGroupEntry {
                        binding: 0,
                        resource: wgpu::BindingResource::TextureView(&src.read),
                    },
                    wgpu::BindGroupEntry {
                        binding: 1,
                        resource: wgpu::BindingResource::TextureView(&bricks_write),
                    },
                ],
            })
        };
        let brick_from_a = brick_bind(&a);
        let brick_from_b = brick_bind(&b);
        let brick_module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
            label: Some("voxel-bricks"),
            source: wgpu::ShaderSource::Wgsl(
                format!("const BRICK: u32 = {BRICK}u;\n{}", shader::BRICK_SHADER).into(),
            ),
        });
        let brick_pipeline_layout =
            device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
                label: Some("voxel-brick-pipeline-layout"),
                bind_group_layouts: &[Some(&brick_layout)],
                immediate_size: 0,
            });
        let brick_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
            label: Some("voxel-bricks"),
            layout: Some(&brick_pipeline_layout),
            module: &brick_module,
            entry_point: Some("main"),
            compilation_options: Default::default(),
            cache: None,
        });

        let luts = palette::LutPair::new(device, "voxel");
        // `[Uniform+size, Texture, Texture, Texture, Texture, Sampler, Texture3D,
        // Sampler]`: the uniform first and sized, so the shape is not
        // `cellular-present-layout`'s field-LUTs-sampler-uniform (ADR-0058).
        let march_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
            label: Some("voxel-march-layout"),
            entries: &[
                wgpu::BindGroupLayoutEntry {
                    binding: 0,
                    visibility: wgpu::ShaderStages::FRAGMENT,
                    ty: wgpu::BindingType::Buffer {
                        ty: wgpu::BufferBindingType::Uniform,
                        has_dynamic_offset: false,
                        min_binding_size: wgpu::BufferSize::new(
                            std::mem::size_of::<MarchParams>() as u64
                        ),
                    },
                    count: None,
                },
                wgpu::BindGroupLayoutEntry {
                    binding: 1,
                    visibility: wgpu::ShaderStages::FRAGMENT,
                    ty: wgpu::BindingType::Texture {
                        sample_type: wgpu::TextureSampleType::Uint,
                        view_dimension: wgpu::TextureViewDimension::D3,
                        multisampled: false,
                    },
                    count: None,
                },
                wgpu::BindGroupLayoutEntry {
                    binding: 2,
                    visibility: wgpu::ShaderStages::FRAGMENT,
                    ty: wgpu::BindingType::Texture {
                        sample_type: wgpu::TextureSampleType::Uint,
                        view_dimension: wgpu::TextureViewDimension::D3,
                        multisampled: false,
                    },
                    count: None,
                },
                gpu::texture(3, true),
                gpu::texture(4, true),
                gpu::sampler(5),
                wgpu::BindGroupLayoutEntry {
                    binding: 6,
                    visibility: wgpu::ShaderStages::FRAGMENT,
                    ty: wgpu::BindingType::Texture {
                        sample_type: wgpu::TextureSampleType::Float { filterable: true },
                        view_dimension: wgpu::TextureViewDimension::D3,
                        multisampled: false,
                    },
                    count: None,
                },
                gpu::sampler(7),
            ],
        });
        let light_stand_in = device.create_texture(&wgpu::TextureDescriptor {
            label: Some("voxel-light-stand-in"),
            size: wgpu::Extent3d {
                width: 1,
                height: 1,
                depth_or_array_layers: 1,
            },
            mip_level_count: 1,
            sample_count: 1,
            dimension: wgpu::TextureDimension::D3,
            format: LIGHT_FORMAT,
            usage: wgpu::TextureUsages::TEXTURE_BINDING,
            view_formats: &[],
        });
        let light_stand_in_view =
            light_stand_in.create_view(&wgpu::TextureViewDescriptor::default());
        let light_sampler = device.create_sampler(&wgpu::SamplerDescriptor {
            label: Some("voxel-light-sampler"),
            address_mode_u: wgpu::AddressMode::ClampToEdge,
            address_mode_v: wgpu::AddressMode::ClampToEdge,
            address_mode_w: wgpu::AddressMode::ClampToEdge,
            mag_filter: wgpu::FilterMode::Linear,
            min_filter: wgpu::FilterMode::Linear,
            ..Default::default()
        });
        let march_bind = |field: &StateTexture| {
            march_bind_group(
                device,
                &march_layout,
                &march_uniform,
                &field.read,
                &bricks_read,
                &luts,
                &light_stand_in_view,
                &light_sampler,
            )
        };
        let march_from_a = march_bind(&a);
        let march_from_b = march_bind(&b);
        let march_shader =
            gpu::fullscreen_shader(device, "voxel-march", gpu::FULLSCREEN_VS_NDC, &march_body());
        // Into the march's own target, which is cleared every frame, so the
        // march writes its premultiplied colour and coverage as they are.
        let march_pipeline = gpu::fullscreen_pipeline(
            device,
            &march_shader,
            &[&march_layout],
            surface_format,
            wgpu::BlendState::REPLACE,
            "voxel-march",
        );

        // `[Sampler:VERTEX_FRAGMENT, Texture]`: the sampler is declared to both
        // stages only so the shape is not `bloom-blur-layout`'s
        // `[Sampler, Texture]` (ADR-0058).
        let present_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
            label: Some("voxel-present-layout"),
            entries: &[
                wgpu::BindGroupLayoutEntry {
                    binding: 0,
                    visibility: wgpu::ShaderStages::VERTEX_FRAGMENT,
                    ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering),
                    count: None,
                },
                gpu::texture(1, true),
            ],
        });
        let present_shader = gpu::fullscreen_shader(
            device,
            "voxel-present",
            gpu::FULLSCREEN_VS_UV_FLIPPED,
            shader::PRESENT_SHADER,
        );
        let present_pipeline = gpu::fullscreen_pipeline(
            device,
            &present_shader,
            &[&present_layout],
            surface_format,
            // Premultiplied OVER the backdrop (ADR-0201): the march's light and
            // the coverage its absorption computed.
            wgpu::BlendState::PREMULTIPLIED_ALPHA_BLENDING,
            "voxel-present",
        );
        let present_sampler = device.create_sampler(&wgpu::SamplerDescriptor {
            label: Some("voxel-present-sampler"),
            mag_filter: wgpu::FilterMode::Linear,
            min_filter: wgpu::FilterMode::Linear,
            ..Default::default()
        });

        Self {
            grid,
            a,
            b,
            reading_a: true,
            seed_pipeline,
            stamp_pipeline,
            step_pipeline,
            step_uniform,
            step_from_a,
            step_from_b,
            _bricks: bricks,
            bricks_read,
            brick_pipeline,
            brick_from_a,
            brick_from_b,
            march_layout,
            march_pipeline,
            march_uniform,
            march_from_a,
            march_from_b,
            _light_stand_in: light_stand_in,
            light_sampler,
            light: None,
            light_fresh: false,
            present_layout,
            present_pipeline,
            present_sampler,
            target: None,
            luts,
        }
    }

    /// Build the light field if it does not exist, and rebind the march to it.
    fn ensure_light(&mut self, device: &wgpu::Device) {
        if self.light.is_some() {
            return;
        }
        let light = LightField::build(device, self.grid, &self.a, &self.b);
        let bind = |field: &StateTexture| {
            march_bind_group(
                device,
                &self.march_layout,
                &self.march_uniform,
                &field.read,
                &self.bricks_read,
                &self.luts,
                &light.blurred_view,
                &self.light_sampler,
            )
        };
        self.march_from_a = bind(&self.a);
        self.march_from_b = bind(&self.b);
        self.light = Some(light);
        self.light_fresh = false;
    }

    /// Refill the light field from the current generation, if it exists.
    fn encode_light(&mut self, encoder: &mut wgpu::CommandEncoder) {
        if let Some(light) = self.light.as_ref() {
            light.encode(encoder, self.reading_a);
            self.light_fresh = true;
        }
    }

    /// Mark every brick of the current generation that holds a non-dead cell.
    fn encode_bricks(&self, encoder: &mut wgpu::CommandEncoder) {
        let side = self.grid.div_ceil(BRICK);
        let mut cpass = gpu::compute_pass(encoder, "voxel-bricks");
        cpass.set_pipeline(&self.brick_pipeline);
        cpass.set_bind_group(
            0,
            if self.reading_a {
                &self.brick_from_a
            } else {
                &self.brick_from_b
            },
            &[],
        );
        cpass.dispatch_workgroups(side, side, side);
    }

    /// Hold the march's target at `size`, rebuilding it only when the size
    /// moved.
    fn ensure_target(
        &mut self,
        device: &wgpu::Device,
        format: wgpu::TextureFormat,
        size: (u32, u32),
    ) {
        if self.target.as_ref().is_none_or(|t| t.size != size) {
            let texture = device.create_texture(&wgpu::TextureDescriptor {
                label: Some("voxel-march-target"),
                size: wgpu::Extent3d {
                    width: size.0,
                    height: size.1,
                    depth_or_array_layers: 1,
                },
                mip_level_count: 1,
                sample_count: 1,
                dimension: wgpu::TextureDimension::D2,
                format,
                usage: wgpu::TextureUsages::RENDER_ATTACHMENT
                    | wgpu::TextureUsages::TEXTURE_BINDING,
                view_formats: &[],
            });
            let view = texture.create_view(&wgpu::TextureViewDescriptor::default());
            let present_bg = device.create_bind_group(&wgpu::BindGroupDescriptor {
                label: Some("voxel-present-bg"),
                layout: &self.present_layout,
                entries: &[
                    wgpu::BindGroupEntry {
                        binding: 0,
                        resource: wgpu::BindingResource::Sampler(&self.present_sampler),
                    },
                    wgpu::BindGroupEntry {
                        binding: 1,
                        resource: wgpu::BindingResource::TextureView(&view),
                    },
                ],
            });
            self.target = Some(MarchTarget {
                size,
                _texture: texture,
                view,
                present_bg,
            });
        }
    }

    /// The texture holding the current generation.
    #[cfg(test)]
    fn current(&self) -> &wgpu::Texture {
        if self.reading_a {
            &self.a.texture
        } else {
            &self.b.texture
        }
    }

    /// Encode one `pass` of the step shader: read the current generation,
    /// write the other texture, swap. Its own compute pass, so the write is
    /// complete before the next pass samples it.
    fn encode_step(&mut self, encoder: &mut wgpu::CommandEncoder, pass: StepPass) {
        let pipeline = match pass {
            StepPass::Step => &self.step_pipeline,
            StepPass::Seed => &self.seed_pipeline,
            StepPass::Stamp => &self.stamp_pipeline,
        };
        let groups = self.grid.div_ceil(WORKGROUP);
        {
            let mut cpass = gpu::compute_pass(encoder, pass.label());
            cpass.set_pipeline(pipeline);
            cpass.set_bind_group(
                0,
                if self.reading_a {
                    &self.step_from_a
                } else {
                    &self.step_from_b
                },
                &[],
            );
            cpass.dispatch_workgroups(groups, groups, groups);
        }
        self.reading_a = !self.reading_a;
    }
}

/// The march's bind group over the generation `field`, its bricks, the palette
/// LUTs, and `light` — the light field, or the stand-in before it exists.
#[allow(clippy::too_many_arguments)]
fn march_bind_group(
    device: &wgpu::Device,
    layout: &wgpu::BindGroupLayout,
    uniform: &wgpu::Buffer,
    field: &wgpu::TextureView,
    bricks: &wgpu::TextureView,
    luts: &palette::LutPair,
    light: &wgpu::TextureView,
    light_sampler: &wgpu::Sampler,
) -> wgpu::BindGroup {
    let [lut_a, lut_b, lut_sampler] = luts.bind_entries(3, 4, 5);
    device.create_bind_group(&wgpu::BindGroupDescriptor {
        label: Some("voxel-march-bg"),
        layout,
        entries: &[
            wgpu::BindGroupEntry {
                binding: 0,
                resource: uniform.as_entire_binding(),
            },
            wgpu::BindGroupEntry {
                binding: 1,
                resource: wgpu::BindingResource::TextureView(field),
            },
            wgpu::BindGroupEntry {
                binding: 2,
                resource: wgpu::BindingResource::TextureView(bricks),
            },
            lut_a,
            lut_b,
            lut_sampler,
            wgpu::BindGroupEntry {
                binding: 6,
                resource: wgpu::BindingResource::TextureView(light),
            },
            wgpu::BindGroupEntry {
                binding: 7,
                resource: wgpu::BindingResource::Sampler(light_sampler),
            },
        ],
    })
}

/// A 3D automaton in a cube, marched as an emitting, absorbing volume, or as
/// lit solid cells inside the glow of their own decay.
pub struct VoxelScene {
    /// Cloned device handle that builds [`Resources`] at `configure`.
    device: wgpu::Device,
    surface_format: wgpu::TextureFormat,
    res: Option<Resources>,
    /// The `[voxel]` table of the preset last configured.
    config: VoxelConfig,
    /// The tier's [`voxel_grid`](crate::render::TierConfig::voxel_grid): the
    /// largest grid a preset runs on.
    grid_cap: u32,
    /// The tier's [`voxel_march_scale`](crate::render::TierConfig::voxel_march_scale).
    march_scale: GridScale,
    /// The tier's [`voxel_march_cap`](crate::render::TierConfig::voxel_march_cap).
    march_cap: (u32, u32),
    /// The render target's size, handed in every frame.
    target: (u32, u32),
    /// The share of the renderer's grid scale the target does not already
    /// carry, handed in every frame (ADR-0245).
    field_scale: GridScale,
    /// Whether the march jumps empty bricks. Always on; a probe turns it off to
    /// hold the jump to the step it replaces.
    skip_bricks: bool,
    /// Whether the next `render` seeds the field before stepping it — set by
    /// every `configure`.
    needs_seed: bool,
    clock: common::GenerationClock,
    /// This frame's `dt`, stored by `advance` and integrated in `update`, where
    /// this frame's `step_rate` has landed.
    dt: f32,
    /// Generations `update` scheduled for the next `render` to encode.
    pending_generations: u32,
    /// How many reseeds have risen since `configure`: the count the next ball
    /// is hashed from.
    rises: u32,
    /// A ball scheduled by a `reseed` rising edge for the next `render`.
    pending_stamp: Option<Stamp>,
    /// Last frame's `reseed`, for the rising edge.
    prev_reseed: f32,
    /// This frame's spectrum laid across the shells, written by `update`.
    levels: [f32; MAX_SHELLS as usize],
    rule: f32,
    step_rate: f32,
    reseed: f32,
    reseed_radius: f32,
    shell_gain: f32,
    density: f32,
    trail: f32,
    age_tint: f32,
    hue_spread: f32,
    key: f32,
    ambient: f32,
    ao: f32,
    bounce: f32,
    /// How many times the light field has been refilled since this scene was
    /// built. Diagnostic: a preset that never binds `bounce` keeps it at 0.
    light_builds: u32,
    light_yaw: f32,
    light_pitch: f32,
    cut: f32,
    cut_yaw: f32,
    cut_pitch: f32,
    zoom: f32,
    colour: PaletteParams,
    pan: PanParams,
    camera: CameraParams,
    /// How much of this scene's coverage the backdrop resolves against
    /// (ADR-0085), set by the renderer every frame.
    occlude: f32,
    /// The active baked palette, held because `set_palette` can arrive before
    /// the resources exist.
    palette: Palette,
}

impl VoxelScene {
    /// The CPU-side state, holding a preset's grid to `grid_cap` and marching at
    /// `march_scale` of the render target held to `march_cap` — the tier's
    /// three voxel values. GPU resources are built at the first `configure`.
    pub fn new(
        device: &wgpu::Device,
        surface_format: wgpu::TextureFormat,
        grid_cap: u32,
        march_scale: GridScale,
        march_cap: (u32, u32),
    ) -> Self {
        Self {
            device: device.clone(),
            surface_format,
            res: None,
            config: VoxelConfig::default(),
            grid_cap: grid_cap.clamp(MIN_GRID, MAX_GRID),
            march_scale,
            march_cap,
            target: (1, 1),
            field_scale: GridScale::FULL,
            skip_bricks: true,
            needs_seed: true,
            clock: common::GenerationClock::default(),
            dt: 0.0,
            pending_generations: 0,
            rises: 0,
            pending_stamp: None,
            prev_reseed: 0.0,
            levels: [0.0; MAX_SHELLS as usize],
            rule: 0.0,
            step_rate: DEFAULT_STEP_RATE,
            reseed: 0.0,
            reseed_radius: DEFAULT_RESEED_RADIUS,
            shell_gain: DEFAULT_SHELL_GAIN,
            density: DEFAULT_DENSITY,
            trail: DEFAULT_TRAIL,
            age_tint: DEFAULT_AGE_TINT,
            hue_spread: DEFAULT_HUE_SPREAD,
            key: DEFAULT_KEY,
            ambient: DEFAULT_AMBIENT,
            ao: DEFAULT_AO,
            bounce: DEFAULT_BOUNCE,
            light_builds: 0,
            light_yaw: DEFAULT_LIGHT_YAW,
            light_pitch: DEFAULT_LIGHT_PITCH,
            cut: DEFAULT_CUT,
            cut_yaw: DEFAULT_CUT_YAW,
            cut_pitch: DEFAULT_CUT_PITCH,
            zoom: DEFAULT_ZOOM,
            colour: PaletteParams::new(DEFAULT_HUE, DEFAULT_BRIGHTNESS),
            pan: PanParams::default(),
            camera: CameraParams::default(),
            occlude: crate::render::post::DEFAULT_OCCLUDE,
            palette: Palette::default_spectrum(),
        }
    }

    /// This frame's step uniform: the rule `rule` picks, the field's seed, and
    /// the ball `stamp` refills if one is scheduled.
    fn step_params(&self, stamp: Option<Stamp>) -> StepParams {
        let rule = self.config.rules.pick(self.rule);
        let stamp = stamp.unwrap_or(Stamp {
            centre: [0, 0, 0],
            radius_sq: 0,
            seed: 0,
        });
        StepParams {
            a: [
                self.config.grid,
                u32::from(self.config.wrap),
                field_seed(self.config.salt),
                live_threshold(self.config.seed_fill),
            ],
            b: [
                rule.birth,
                rule.survive,
                rule.states,
                rule.neighbourhood.index(),
            ],
            c: [
                seed_radius_sq(self.config.seed_radius, self.config.grid),
                stamp.seed,
                0,
                0,
            ],
            d: [
                stamp.centre[0],
                stamp.centre[1],
                stamp.centre[2],
                stamp.radius_sq,
            ],
        }
    }

    /// How many times the light field has been refilled. See
    /// [`Self::light_builds`].
    #[cfg(test)]
    pub(crate) fn light_build_count(&self) -> u32 {
        self.light_builds
    }

    /// The bounce the march adds this frame: the bound value, held finite and
    /// inside `0..=`[`MAX_BOUNCE`], and 0 under `glow`, which reads no light.
    fn applied_bounce(&self) -> f32 {
        match self.config.present {
            Present::Glow => 0.0,
            Present::Solid => finite_or(self.bounce, DEFAULT_BOUNCE).clamp(0.0, MAX_BOUNCE),
        }
    }

    /// This frame's march uniform, for a render target of `aspect`.
    fn march_params(&self, aspect: f32) -> MarchParams {
        let zoom = if self.zoom.is_finite() && self.zoom > 1e-3 {
            self.zoom
        } else {
            DEFAULT_ZOOM
        };
        let view = self.camera.camera().view(
            aspect,
            zoom,
            [finite_or(self.pan.x, 0.0), finite_or(self.pan.y, 0.0)],
        );
        let [r0, r1, r3] = rows(&view);
        let [ex, ey, ez] = eye(&view);
        let cut = cut_plane(self.cut, self.cut_yaw, self.cut_pitch);
        MarchParams {
            r0,
            r1,
            r3,
            eye: [ex, ey, ez, self.config.grid as f32],
            a: [
                self.colour.hue,
                self.colour.brightness,
                self.colour.saturation,
                self.colour.mix,
            ],
            b: [
                palette::band_steps(self.colour.steps),
                self.occlude,
                finite_or(self.density, DEFAULT_DENSITY).max(0.0),
                finite_or(self.trail, DEFAULT_TRAIL).clamp(0.0, 1.0),
            ],
            c: [
                finite_or(self.age_tint, DEFAULT_AGE_TINT),
                finite_or(self.hue_spread, DEFAULT_HUE_SPREAD),
                finite_or(self.camera.fog, 0.0).clamp(0.0, 1.0),
                finite_or(self.shell_gain, DEFAULT_SHELL_GAIN).clamp(0.0, MAX_SHELL_GAIN),
            ],
            d: [
                view.distance - CUBE_RADIUS,
                2.0 * CUBE_RADIUS,
                self.config.shells.min(MAX_SHELLS) as f32,
                if self.skip_bricks { 1.0 } else { 0.0 },
            ],
            light: {
                let [lx, ly, lz] = light_dir(&view, self.light_yaw, self.light_pitch);
                let solid = match self.config.present {
                    Present::Glow => 0.0,
                    Present::Solid => 1.0,
                };
                [lx, ly, lz, solid]
            },
            lit: [
                finite_or(self.key, DEFAULT_KEY).max(0.0),
                finite_or(self.ambient, DEFAULT_AMBIENT).max(0.0),
                finite_or(self.ao, DEFAULT_AO).clamp(0.0, 1.0),
                if cut.is_some() { 1.0 } else { 0.0 },
            ],
            cut: cut.unwrap_or([0.0; 4]),
            bounce: [self.applied_bounce(), 0.0, 0.0, 0.0],
            shells: {
                let mut out = [[0.0; 4]; (MAX_SHELLS / 4) as usize];
                for (slot, level) in out.iter_mut().flatten().zip(self.levels) {
                    *slot = level;
                }
                out
            },
        }
    }
}

impl Scene for VoxelScene {
    fn name(&self) -> &'static str {
        "voxel"
    }

    fn advance(&mut self, dt: f32) {
        self.dt = dt;
    }

    fn set_target_size(&mut self, width: u32, height: u32) {
        // Recorded, never built here (ADR-0030): `render` compares the size it
        // asks for with the target it holds.
        self.target = (width.max(1), height.max(1));
    }

    fn set_grid_scale(&mut self, scale: GridScale) {
        self.field_scale = scale;
    }

    fn set_occlude(&mut self, occlude: f32) {
        self.occlude = occlude;
    }

    fn set_palette(&mut self, palette: &Palette) {
        self.palette = palette.clone();
        if let Some(res) = self.res.as_mut() {
            res.luts.set(palette);
        }
    }

    fn configure(&mut self, cfg: &GeneratorConfig) -> Option<super::CapOverflow> {
        let GeneratorConfig::Voxel(config) = cfg else {
            return None;
        };
        // A switch starts the incoming preset from its own seed. The grid is
        // held to the cap here, once, and the clamp is returned for the
        // renderer to announce with the preset. The state pair is built here,
        // off the frame, and only when the grid changes.
        let (config, overflow) = grid_clamp(*config, self.grid_cap);
        self.config = config;
        if self.res.as_ref().is_none_or(|res| res.grid != config.grid) {
            let mut built = Resources::build(&self.device, self.surface_format, config.grid);
            built.luts.set(&self.palette);
            self.res = Some(built);
        }
        self.needs_seed = true;
        self.clock = common::GenerationClock::default();
        self.pending_generations = 0;
        // The incoming preset's balls are its own stream, from its first.
        self.rises = 0;
        self.pending_stamp = None;
        self.prev_reseed = 0.0;
        overflow
    }

    fn reset_params(&mut self) {
        self.rule = 0.0;
        self.step_rate = DEFAULT_STEP_RATE;
        self.reseed = 0.0;
        self.reseed_radius = DEFAULT_RESEED_RADIUS;
        self.shell_gain = DEFAULT_SHELL_GAIN;
        self.density = DEFAULT_DENSITY;
        self.trail = DEFAULT_TRAIL;
        self.age_tint = DEFAULT_AGE_TINT;
        self.hue_spread = DEFAULT_HUE_SPREAD;
        self.key = DEFAULT_KEY;
        self.ambient = DEFAULT_AMBIENT;
        self.ao = DEFAULT_AO;
        self.bounce = DEFAULT_BOUNCE;
        self.light_yaw = DEFAULT_LIGHT_YAW;
        self.light_pitch = DEFAULT_LIGHT_PITCH;
        self.cut = DEFAULT_CUT;
        self.cut_yaw = DEFAULT_CUT_YAW;
        self.cut_pitch = DEFAULT_CUT_PITCH;
        self.zoom = DEFAULT_ZOOM;
        self.colour.reset();
        self.pan.reset();
        self.camera.reset();
    }

    fn set_param(&mut self, name: &str, value: f32) {
        // The shared param blocks first, this scene's own names after
        // (`scenes::common`).
        if self.colour.set(name, value) || self.pan.set(name, value) {
            return;
        }
        match name {
            // The camera block's names, matched here rather than delegated to
            // `CameraParams::set`: the march declares all of it but `solid`,
            // which sorts strokes and a march has none, and a delegation would
            // answer that undeclared name too. `focus` and `aperture` are
            // stored and never read.
            "yaw" => self.camera.yaw = value,
            "pitch" => self.camera.pitch = value,
            "distance" => self.camera.distance = value,
            "fov" => self.camera.fov = value,
            "focus" => self.camera.focus = value,
            "aperture" => self.camera.aperture = value,
            "fog" => self.camera.fog = value,
            "rule" => self.rule = value,
            "step_rate" => self.step_rate = value,
            "reseed" => self.reseed = value,
            "reseed_radius" => self.reseed_radius = value,
            "shell_gain" => self.shell_gain = value,
            "density" => self.density = value,
            "trail" => self.trail = value,
            "age_tint" => self.age_tint = value,
            "hue_spread" => self.hue_spread = value,
            "key" => self.key = value,
            "ambient" => self.ambient = value,
            "ao" => self.ao = value,
            "bounce" => self.bounce = value,
            "light_yaw" => self.light_yaw = value,
            "light_pitch" => self.light_pitch = value,
            "cut" => self.cut = value,
            "cut_yaw" => self.cut_yaw = value,
            "cut_pitch" => self.cut_pitch = value,
            "zoom" => self.zoom = value,
            _ => {}
        }
    }

    fn update(&mut self, frame: &AnalysisFrame) {
        self.pending_generations = self
            .clock
            .advance(applied_step_rate(self.step_rate), self.dt);
        self.levels = shell_levels(&frame.spectrum, self.config.shells);
        // Rising edge only, so a beat flag held for several frames — or a
        // latch's hold — refills one ball rather than one per frame. A NaN
        // compares false both ways and is stored as zero, so it neither fires
        // nor arms a spurious edge on the next finite value.
        let reseed = finite_or(self.reseed, 0.0);
        if reseed >= RESEED_THRESHOLD && self.prev_reseed < RESEED_THRESHOLD {
            self.pending_stamp = Some(stamp(
                self.config.salt,
                self.rises,
                self.config.grid,
                finite_or(self.reseed_radius, DEFAULT_RESEED_RADIUS),
            ));
            self.rises = self.rises.wrapping_add(1);
        }
        self.prev_reseed = reseed;
    }

    fn render(
        &mut self,
        queue: &wgpu::Queue,
        encoder: &mut wgpu::CommandEncoder,
        view: &wgpu::TextureView,
        aspect: f32,
    ) {
        // The aspect is the render target's, handed in here (ADR-0037).
        let march = self.march_params(aspect);
        let seed = std::mem::replace(&mut self.needs_seed, false);
        let stamp = self.pending_stamp.take();
        let generations = std::mem::take(&mut self.pending_generations);
        // The rule `rule` picks this frame runs every generation below, so a
        // change lands at the next generation boundary.
        let step = self.step_params(stamp);
        let size = march_size(
            self.target,
            self.march_scale.get() * self.field_scale.get(),
            self.march_cap,
        );
        let Some(res) = self.res.as_mut() else {
            return;
        };
        res.luts.flush(queue);
        queue.write_buffer(&res.march_uniform, 0, bytemuck::bytes_of(&march));
        // One write serves every pass below: the seed, then the ball, then the
        // generations, in that order, so a ball scheduled on a preset's first
        // frame lands on its seeded field rather than under it.
        let changed = seed || stamp.is_some() || generations > 0;
        if changed {
            queue.write_buffer(&res.step_uniform, 0, bytemuck::bytes_of(&step));
        }
        if seed {
            res.encode_step(encoder, StepPass::Seed);
        }
        if stamp.is_some() {
            res.encode_step(encoder, StepPass::Stamp);
        }
        for _ in 0..generations {
            res.encode_step(encoder, StepPass::Step);
        }
        // The march reads only the frame's last generation, so the bricks are
        // marked once, after it.
        if changed {
            res.encode_bricks(encoder);
            res.light_fresh = false;
        }
        // The light field follows the bricks, and a field built or switched
        // back on this frame is filled this frame, state change or not.
        if march.bounce[0] > 0.0 {
            res.ensure_light(&self.device);
            if !res.light_fresh {
                res.encode_light(encoder);
                self.light_builds = self.light_builds.wrapping_add(1);
            }
        }

        res.ensure_target(&self.device, self.surface_format, size);
        let Some(target) = res.target.as_ref() else {
            return;
        };
        {
            // The march writes every pixel of its own target, which starts each
            // frame empty.
            let mut pass = gpu::color_pass(
                encoder,
                "voxel-march-pass",
                &target.view,
                wgpu::LoadOp::Clear(wgpu::Color::TRANSPARENT),
            );
            pass.set_pipeline(&res.march_pipeline);
            pass.set_bind_group(
                0,
                if res.reading_a {
                    &res.march_from_a
                } else {
                    &res.march_from_b
                },
                &[],
            );
            pass.draw(0..3, 0..1);
        }
        // A plain stretch over the engine backdrop (ADR-0018): a ray through
        // empty cells wrote no light and no coverage.
        let mut pass = gpu::color_pass(encoder, "voxel-present-pass", view, wgpu::LoadOp::Load);
        pass.set_pipeline(&res.present_pipeline);
        pass.set_bind_group(0, &target.present_bg, &[]);
        pass.draw(0..3, 0..1);
    }
}

#[cfg(test)]
mod tests;