Your browser sucks, JavaScript disabled! Gunnar's console needs JavaScript and WebAssembly. Git over SSH on port 2222 does not.

nordisk/facett

facett-map3d-flight/src/cdlod.rs

//! **CDLOD terrain quadtree** (roadmap §6) — a *Continuous Distance-dependent Level
//! Of Detail* selector draped on the already-landed real terrain
//! ([`crate::terrain`]). The pure-CPU **selection math** — which quadtree nodes to
//! draw at which LOD given the camera, plus the per-node **morph factor** that makes
//! the LOD transitions crack-free — lives here, headless + deterministically
//! inject-assertable (FC-7). The selected nodes feed the existing terrain draw
//! ([`crate::terrain::build_terrain_mesh`], reused not forked), so a near node is a
//! *small* geographic patch sampled at the same grid resolution → finer topography,
//! while a far node is a *large* patch at the same grid → coarser, cheaper mesh.
//!
//! ## Why CDLOD (Filip Strugar's scheme)
//! A flat terrain grid over a whole DEM tile is either too coarse near the camera
//! (blocky ridgelines) or ruinously heavy far away (millions of sub-pixel triangles).
//! CDLOD keeps the triangle count **roughly constant** by subdividing a quadtree only
//! where the camera is close: each LOD level owns a **distance range**, and a node is
//! drawn at the coarsest level whose range still reaches it. The classic two failure
//! modes it fixes:
//! - **T-junction cracks** between a fine node and its coarser neighbour (the finer
//!   edge has vertices the coarse edge lacks → a one-pixel gap you can see sky
//!   through). CDLOD guarantees **adjacent selected nodes differ by ≤ 1 level**, then
//!   **morphs** the finer node's extra edge vertices onto the coarse edge as distance
//!   grows, so the seam closes continuously (no popping).
//! - **LOD popping** as you fly: the [`morph_factor`] ramps `0 → 1` across the outer
//!   third of each node's range, blending a node's geometry toward its parent grid
//!   *before* the hard switch, so the change is imperceptible.
//!
//! ## What is CPU vs GPU here
//! - **CPU (this module, always available, RED-when-broken tested):** the quadtree
//!   [`Cdlod::select`] (node subdivision by distance), the per-node LOD + [`morph_factor`],
//!   the triangle budget, and the drape via [`Cdlod::build_meshes`]. All pure `f64`/`f32`
//!   math over the [`crate::origin::WorldPos`] metric frame — no GPU, no adapter.
//! - **GPU (`#[cfg(feature = "wgpu")]`, adapter-guarded):** the vertex **height-morph**
//!   the terrain vertex shader applies is proven to match the CPU [`morph_height`]
//!   oracle texel-for-texel by a compute pass ([`cdlod_morph.wgsl`]), self-skipping with
//!   a recorded sentinel when no adapter is present (the [`crate::scatter_gpu`] pattern).
//!   The full render-pass wiring into `facett-map3d`'s pipeline is the next roadmap step
//!   (like the scattering LUT's `sky.wgsl` sampling), so the landed 21-pass pipeline is
//!   untouched (L2 additive).

use crate::origin::WorldPos;
use crate::terrain::{Elevation, GeoRegion, TerrainMesh, build_terrain_mesh};

// ─────────────────────────────────────────────────────────────────────────────────
// Config
// ─────────────────────────────────────────────────────────────────────────────────

/// The CDLOD selection knobs. Defaults are a sensible flight-sim ladder; every field
/// is injectable so a test can pin an exact subdivision.
#[derive(Clone, Copy, Debug, PartialEq)]
pub struct CdlodConfig {
    /// Quadtree depth: the root is depth 0 (one node over the whole region), the finest
    /// leaves are depth `max_depth` (2^`max_depth` nodes per side). Also the number of
    /// LOD levels − 1 (the finest LOD is 0, the root's LOD is `max_depth`).
    pub max_depth: u32,
    /// The **finest** LOD's visibility range in metres: leaf (LOD-0) detail is required
    /// within this ground distance of the camera. Each coarser level's range is this
    /// times [`CdlodConfig::range_ratio`] per level.
    pub leaf_range_m: f64,
    /// Ratio between successive LOD ranges (`range(lod+1) = range(lod) · ratio`).
    /// The canonical CDLOD value is 2.0 (each level reaches twice as far).
    pub range_ratio: f64,
    /// Where in each node's range the morph starts, as a fraction of the outer edge:
    /// `morph_start = range(lod) · morph_start_ratio`. Below it a node is at full
    /// detail (morph 0); from it to the outer edge morph ramps `0 → 1`. Strugar uses
    /// ~0.66; the transition band is the outer `1 − ratio` of the range.
    pub morph_start_ratio: f64,
    /// Vertices per side of every node's terrain patch (the shared CDLOD grid mesh).
    /// The triangle budget is `selected_nodes · (grid−1)² · 2`. `≥ 2`.
    pub grid: usize,
}

impl Default for CdlodConfig {
    fn default() -> Self {
        Self { max_depth: 5, leaf_range_m: 1_500.0, range_ratio: 2.0, morph_start_ratio: 0.66, grid: 33 }
    }
}

/// A quadtree node selected for drawing this frame: the geographic patch it covers, its
/// tree `depth` + LOD `level`, the camera ground distance to it, and its **morph
/// factor** (`0` = full detail, `1` = fully morphed onto the parent grid). Everything
/// the terrain draw needs to build + morph the patch.
#[derive(Clone, Copy, Debug, PartialEq)]
pub struct SelectedNode {
    /// The geographic rectangle this node covers.
    pub region: GeoRegion,
    /// Tree depth (0 = root, `max_depth` = finest leaf). Node **size** halves per depth.
    pub depth: u32,
    /// LOD level (0 = finest grid, `max_depth` = coarsest). The grid **density** level.
    pub level: u32,
    /// Camera ground distance (metres) to the nearest point of the node's world AABB.
    pub dist_m: f64,
    /// Crack-free morph factor `[0, 1]`: how far this node's geometry is blended toward
    /// its parent (coarser) grid. `0` near the camera, `1` at the node's outer range.
    pub morph: f32,
}

// ─────────────────────────────────────────────────────────────────────────────────
// The selector
// ─────────────────────────────────────────────────────────────────────────────────

/// A CDLOD quadtree over a fixed [`GeoRegion`], evaluated against a camera each frame.
/// Holds no per-node state — [`Cdlod::select`] walks the tree geometrically (children
/// are the four lon/lat quadrants), so selection is a pure function of the camera
/// (FC-7 deterministic).
#[derive(Clone, Copy, Debug, PartialEq)]
pub struct Cdlod {
    /// The whole terrain extent the quadtree tiles.
    pub region: GeoRegion,
    /// The selection knobs.
    pub config: CdlodConfig,
}

impl Cdlod {
    /// A CDLOD selector over `region` with `config`.
    #[must_use]
    pub fn new(region: GeoRegion, config: CdlodConfig) -> Self {
        Self { region, config }
    }

    /// The root LOD level (= `max_depth`): the coarsest level, owned by the root node.
    #[must_use]
    pub fn max_level(&self) -> u32 {
        self.config.max_depth
    }

    /// The visibility range (metres) of LOD `level`: `leaf_range_m · ratio^level`. The
    /// finest level 0 has the smallest range; each coarser level reaches `ratio`× farther.
    #[must_use]
    pub fn range(&self, level: u32) -> f64 {
        self.config.leaf_range_m * self.config.range_ratio.powi(level as i32)
    }

    /// The **morph factor** for a node at `level` whose nearest point is `dist_m` away:
    /// `0` while inside `range(level)·morph_start_ratio` (full detail), ramping linearly
    /// to `1` at the outer edge `range(level)` and clamped beyond. The root level has no
    /// coarser parent to morph toward, so its morph is `0`.
    #[must_use]
    pub fn morph_factor(&self, level: u32, dist_m: f64) -> f32 {
        if level >= self.max_level() {
            return 0.0; // root: nothing coarser to blend to
        }
        let end = self.range(level);
        let start = end * self.config.morph_start_ratio;
        if end <= start {
            return 0.0;
        }
        let t = (dist_m - start) / (end - start);
        t.clamp(0.0, 1.0) as f32
    }

    /// **Select the quadtree nodes to draw** for a camera at `camera` (its `f64` world
    /// position; only the horizontal position matters for the ground-distance ranges).
    /// Returns the selected nodes; their regions **exactly tile** [`Cdlod::region`] with
    /// no gaps or overlaps (crack-free coverage), and spatially-adjacent nodes differ by
    /// at most one depth level. Deterministic: identical camera ⇒ identical `Vec`
    /// (nodes emitted in a fixed pre-order, so a scripted flythrough is reproducible).
    #[must_use]
    pub fn select(&self, camera: WorldPos) -> Vec<SelectedNode> {
        let mut out = Vec::new();
        self.select_node(self.region, 0, self.max_level(), camera, &mut out);
        out
    }

    /// Recursive CDLOD node selection (Strugar): draw the node whole at `level` unless
    /// the *finer* level's range reaches it, in which case subdivide — children the
    /// finer range reaches recurse, children it does not are drawn at this `level`
    /// (guaranteeing the ≤ 1-level neighbour invariant).
    fn select_node(&self, region: GeoRegion, depth: u32, level: u32, camera: WorldPos, out: &mut Vec<SelectedNode>) {
        let dist = ground_dist_to_region(region, camera);
        // Finest level, or the tree bottom: this node is a leaf — draw it.
        if level == 0 || depth >= self.config.max_depth {
            out.push(self.make_node(region, depth, 0, dist));
            return;
        }
        let finer_range = self.range(level - 1);
        if dist > finer_range {
            // The finer detail is not needed anywhere in this node → draw it whole here.
            out.push(self.make_node(region, depth, level, dist));
            return;
        }
        // Subdivide. Children the finer range reaches recurse to level−1; the rest are
        // drawn at this (parent) level so neighbours never differ by more than one level.
        for child in quadrants(region) {
            let cdist = ground_dist_to_region(child, camera);
            if cdist > finer_range {
                out.push(self.make_node(child, depth + 1, level, cdist));
            } else {
                self.select_node(child, depth + 1, level - 1, camera, out);
            }
        }
    }

    #[inline]
    fn make_node(&self, region: GeoRegion, depth: u32, level: u32, dist_m: f64) -> SelectedNode {
        SelectedNode { region, depth, level, dist_m, morph: self.morph_factor(level, dist_m) }
    }

    /// The total triangle count the selection draws: `nodes · (grid−1)² · 2`. CDLOD
    /// keeps this roughly constant regardless of terrain size — the headline win.
    #[must_use]
    pub fn triangle_count(&self, nodes: usize) -> usize {
        let g = self.config.grid.max(2);
        nodes * (g - 1) * (g - 1) * 2
    }

    /// **Drape the selected nodes on the real terrain** (roadmap §6): build a
    /// [`TerrainMesh`] per selected node by sampling `elev` over the node's region at the
    /// configured grid, rebased through `origin` (reusing [`build_terrain_mesh`], not
    /// forking it). A near node is a small patch → finer topography; a far node is a big
    /// patch at the same grid → coarser. Returns `(SelectedNode, TerrainMesh)` pairs so
    /// the caller keeps each node's morph factor for the vertex shader.
    #[must_use]
    pub fn build_meshes(
        &self,
        camera: WorldPos,
        elev: &impl Elevation,
        origin: crate::origin::FloatingOrigin,
        vertical_exaggeration: f64,
    ) -> Vec<(SelectedNode, TerrainMesh)> {
        let g = self.config.grid.max(2);
        self.select(camera)
            .into_iter()
            .map(|node| {
                let mesh = build_terrain_mesh(elev, node.region, g, g, origin, vertical_exaggeration);
                (node, mesh)
            })
            .collect()
    }
}

/// The four lon/lat quadrants of a region, split at its midpoints — the quadtree
/// children. Order is fixed (SW, SE, NW, NE) for deterministic selection.
#[must_use]
fn quadrants(r: GeoRegion) -> [GeoRegion; 4] {
    let mid_lon = 0.5 * (r.min_lon + r.max_lon);
    let mid_lat = 0.5 * (r.min_lat + r.max_lat);
    [
        GeoRegion::new(r.min_lon, r.min_lat, mid_lon, mid_lat), // SW
        GeoRegion::new(mid_lon, r.min_lat, r.max_lon, mid_lat), // SE
        GeoRegion::new(r.min_lon, mid_lat, mid_lon, r.max_lat), // NW
        GeoRegion::new(mid_lon, mid_lat, r.max_lon, r.max_lat), // NE
    ]
}

/// Camera **ground distance** (metres) to the nearest point of a region's world AABB.
/// The region's lon/lat corners project (monotonically) to a Web-Mercator `x`/`y`
/// rectangle; the distance is `0` if the camera is horizontally over the region, else
/// the distance to the nearest edge/corner. Reuses the [`WorldPos`] metric frame the
/// floating origin + streamer already key on.
#[must_use]
fn ground_dist_to_region(r: GeoRegion, camera: WorldPos) -> f64 {
    // Mercator is monotone in lon and lat, so min/max lon,lat map to min/max x,y.
    let sw = WorldPos::from_lonlat(r.min_lon, r.min_lat, 0.0);
    let ne = WorldPos::from_lonlat(r.max_lon, r.max_lat, 0.0);
    let (min_x, max_x) = (sw.x.min(ne.x), sw.x.max(ne.x));
    let (min_y, max_y) = (sw.y.min(ne.y), sw.y.max(ne.y));
    let dx = (camera.x - camera.x.clamp(min_x, max_x)).abs();
    let dy = (camera.y - camera.y.clamp(min_y, max_y)).abs();
    (dx * dx + dy * dy).sqrt()
}

// ─────────────────────────────────────────────────────────────────────────────────
// The morph height oracle (the CPU reference the GPU vertex shader mirrors)
// ─────────────────────────────────────────────────────────────────────────────────

/// **The CDLOD height morph** (the crack-closing blend): a vertex's height morphs from
/// its own grid value `h` toward the **parent grid** value `h_parent` (the height the
/// coarser neighbour has at that spot) by the node's `morph` factor. At `morph = 0` the
/// full-detail height is kept; at `morph = 1` the vertex sits exactly on the parent
/// grid, so the shared edge with a coarser neighbour matches vertex-for-vertex → no
/// crack. This is the single scalar op the terrain vertex shader runs per vertex; the
/// GPU pass ([`cdlod_morph.wgsl`]) is proven equal to this oracle.
#[must_use]
pub fn morph_height(h: f32, h_parent: f32, morph: f32) -> f32 {
    h + morph.clamp(0.0, 1.0) * (h_parent - h)
}

#[cfg(test)]
mod tests {
    use super::*;
    use crate::origin::FloatingOrigin;
    use crate::terrain::ProceduralTerrain;

    /// A ~11 km square test region near Stockholm (0.1° ≈ 11 km lat).
    fn region() -> GeoRegion {
        GeoRegion::new(18.0, 59.25, 18.1, 59.35)
    }

    /// The camera directly over the region centre at a given altitude (only the
    /// horizontal position matters for the ranges; altitude is carried for realism).
    fn cam_over_center(r: GeoRegion, alt_m: f64) -> WorldPos {
        let (clon, clat) = r.center();
        WorldPos::from_lonlat(clon, clat, alt_m)
    }

    /// Two regions are **edge-adjacent** if they share a border segment (touch along a
    /// vertical or horizontal edge with overlapping extent) — used to prove the ≤1-level
    /// neighbour invariant.
    fn edge_adjacent(a: GeoRegion, b: GeoRegion) -> bool {
        let eps = 1e-9;
        let x_overlap = a.min_lon < b.max_lon - eps && b.min_lon < a.max_lon - eps;
        let y_overlap = a.min_lat < b.max_lat - eps && b.min_lat < a.max_lat - eps;
        // Share a vertical edge (a.right == b.left or vice-versa) with y overlap.
        let vshare = ((a.max_lon - b.min_lon).abs() < eps || (b.max_lon - a.min_lon).abs() < eps) && y_overlap;
        // Share a horizontal edge with x overlap.
        let hshare = ((a.max_lat - b.min_lat).abs() < eps || (b.max_lat - a.min_lat).abs() < eps) && x_overlap;
        vshare || hshare
    }

    /// **INJECT-ASSERT (subdivision follows the camera — RED-when-broken):** flying low
    /// over the centre must subdivide the quadtree to its finest leaf (a LOD-0 node
    /// exists) and select **more** nodes than a camera parked far out at altitude (which
    /// collapses toward the coarse root). If `select` stopped keying on distance both
    /// counts would be equal and this fails.
    #[test]
    fn subdivision_follows_camera_distance() {
        let cd = Cdlod::new(region(), CdlodConfig::default());

        let near = cd.select(cam_over_center(region(), 200.0));
        let finest_near = near.iter().map(|n| n.level).min().unwrap();
        let deepest_near = near.iter().map(|n| n.depth).max().unwrap();
        assert_eq!(finest_near, 0, "low over centre reaches the finest LOD (level 0)");
        assert_eq!(deepest_near, cd.config.max_depth, "…and the deepest leaf depth");

        // Park the camera ~500 km away: even the coarsest level's range can't reach, so
        // the whole region collapses to the single root node.
        let far_cam = WorldPos::from_lonlat(24.0, 59.3, 30_000.0);
        let far = cd.select(far_cam);
        assert_eq!(far.len(), 1, "a far camera draws just the root node, got {}", far.len());
        assert_eq!(far[0].level, cd.max_level(), "the lone far node is the coarsest (root) LOD");

        assert!(near.len() > far.len(), "near view selects more nodes ({} > {})", near.len(), far.len());

        #[cfg(feature = "testmatrix")]
        facett_core::testmatrix::emit(
            "facett-map3d-flight",
            "cdlod_subdivision",
            finest_near == 0 && deepest_near == cd.config.max_depth && near.len() > far.len(),
            &format!(
                "near_nodes={} near_finest_lod={finest_near} near_deepest_depth={deepest_near} far_nodes={} tris_near={} tris_far={}",
                near.len(),
                far.len(),
                cd.triangle_count(near.len()),
                cd.triangle_count(far.len()),
            ),
        );
    }

    /// **INJECT-ASSERT (crack-free coverage — the watertight tiling, RED-when-broken):**
    /// the selected node regions **exactly tile** the whole terrain extent — their
    /// lon/lat areas sum to the region area (no gaps, no overlaps), the union bounds
    /// equal the region bounds, and no two selected regions overlap. A selection bug
    /// (dropped child / double-emit) breaks the area sum immediately.
    #[test]
    fn selection_tiles_the_region_watertight() {
        let cd = Cdlod::new(region(), CdlodConfig::default());
        let sel = cd.select(cam_over_center(region(), 350.0));
        assert!(sel.len() > 4, "an interesting selection has several nodes, got {}", sel.len());

        let region_area = (cd.region.max_lon - cd.region.min_lon) * (cd.region.max_lat - cd.region.min_lat);
        let sum: f64 = sel.iter().map(|n| (n.region.max_lon - n.region.min_lon) * (n.region.max_lat - n.region.min_lat)).sum();
        assert!(
            (sum - region_area).abs() < region_area * 1e-9,
            "selected areas tile the region exactly (Σ={sum}, region={region_area})"
        );

        // Union bounds == region bounds.
        let (mut lo_lon, mut lo_lat, mut hi_lon, mut hi_lat) = (f64::MAX, f64::MAX, f64::MIN, f64::MIN);
        for n in &sel {
            lo_lon = lo_lon.min(n.region.min_lon);
            lo_lat = lo_lat.min(n.region.min_lat);
            hi_lon = hi_lon.max(n.region.max_lon);
            hi_lat = hi_lat.max(n.region.max_lat);
        }
        assert!((lo_lon - cd.region.min_lon).abs() < 1e-12 && (hi_lon - cd.region.max_lon).abs() < 1e-12);
        assert!((lo_lat - cd.region.min_lat).abs() < 1e-12 && (hi_lat - cd.region.max_lat).abs() < 1e-12);

        // No two selected regions overlap (interiors disjoint).
        let eps = 1e-9;
        for (i, a) in sel.iter().enumerate() {
            for b in &sel[i + 1..] {
                let x_ov = a.region.min_lon < b.region.max_lon - eps && b.region.min_lon < a.region.max_lon - eps;
                let y_ov = a.region.min_lat < b.region.max_lat - eps && b.region.min_lat < a.region.max_lat - eps;
                assert!(!(x_ov && y_ov), "selected regions must not overlap: {:?} vs {:?}", a.region, b.region);
            }
        }
    }

    /// **INJECT-ASSERT (the ≤1-level neighbour invariant — the anti-crack guarantee):**
    /// any two spatially edge-adjacent selected nodes differ in depth by at most one, so
    /// a T-junction spans a single subdivision the morph can close. Break the "draw the
    /// non-recursed children at the parent level" rule and adjacent depths jump by 2+,
    /// re-opening cracks — this catches it.
    #[test]
    fn adjacent_nodes_differ_by_at_most_one_level() {
        let cd = Cdlod::new(region(), CdlodConfig::default());
        // An off-centre low camera makes an asymmetric tree (near + far quadrants meet).
        let cam = WorldPos::from_lonlat(18.02, 59.27, 250.0);
        let sel = cd.select(cam);
        let mut adjacencies = 0usize;
        for (i, a) in sel.iter().enumerate() {
            for b in &sel[i + 1..] {
                if edge_adjacent(a.region, b.region) {
                    adjacencies += 1;
                    let dd = a.depth.abs_diff(b.depth);
                    assert!(dd <= 1, "edge-adjacent nodes differ by ≤1 depth, got {dd} ({:?} d{} vs {:?} d{})", a.region, a.depth, b.region, b.depth);
                }
            }
        }
        assert!(adjacencies > 0, "the asymmetric selection has adjacent nodes to check");
    }

    /// **INJECT-ASSERT (the morph factor ramp — RED-when-broken):** the crack-free morph
    /// is `0` at/inside `morph_start`, `1` at/beyond the outer range edge, and strictly
    /// monotonic increasing across the transition band. A wrong morph formula (popping)
    /// fails the endpoints or the monotonicity.
    #[test]
    fn morph_factor_ramps_zero_to_one_across_the_band() {
        let cd = Cdlod::new(region(), CdlodConfig::default());
        let level = 1u32; // a mid level (root is exempt)
        let end = cd.range(level);
        let start = end * cd.config.morph_start_ratio;

        assert_eq!(cd.morph_factor(level, 0.0), 0.0, "full detail at the camera");
        assert_eq!(cd.morph_factor(level, start), 0.0, "morph starts at morph_start");
        assert!((cd.morph_factor(level, (start + end) * 0.5) - 0.5).abs() < 1e-5, "half morph at the band midpoint");
        assert_eq!(cd.morph_factor(level, end), 1.0, "fully morphed at the outer edge");
        assert_eq!(cd.morph_factor(level, end * 2.0), 1.0, "clamped beyond the edge");

        // Strictly monotonic across the band.
        let mut prev = -1.0f32;
        for i in 0..=20 {
            let d = start + (end - start) * i as f64 / 20.0;
            let m = cd.morph_factor(level, d);
            assert!(m >= prev, "morph is monotonic non-decreasing across the band");
            prev = m;
        }
        // The root level never morphs (no coarser parent).
        assert_eq!(cd.morph_factor(cd.max_level(), 0.0), 0.0);
        assert_eq!(cd.morph_factor(cd.max_level(), 1e6), 0.0, "root has no parent to morph toward");

        #[cfg(feature = "testmatrix")]
        facett_core::testmatrix::emit(
            "facett-map3d-flight",
            "cdlod_morph_factor",
            cd.morph_factor(level, start) == 0.0 && cd.morph_factor(level, end) == 1.0,
            &format!("level={level} start_m={start:.1} end_m={end:.1} mid_morph={:.3}", cd.morph_factor(level, (start + end) * 0.5)),
        );
    }

    /// **INJECT-ASSERT (the morph HEIGHT blend closes the seam — the oracle):** the
    /// per-vertex height morph (the op the GPU vertex shader mirrors) is the exact linear
    /// blend `h → h_parent`: identity at morph 0, the parent height at morph 1, halfway
    /// at 0.5. This is the crack-closing math the GPU parity proof validates against.
    #[test]
    fn morph_height_blends_to_the_parent_grid() {
        assert!((morph_height(120.0, 80.0, 0.0) - 120.0).abs() < 1e-6, "morph 0 keeps the fine height");
        assert!((morph_height(120.0, 80.0, 1.0) - 80.0).abs() < 1e-6, "morph 1 sits on the parent grid");
        assert!((morph_height(120.0, 80.0, 0.5) - 100.0).abs() < 1e-6, "morph 0.5 is the midpoint");
        assert!((morph_height(50.0, 50.0, 0.7) - 50.0).abs() < 1e-6, "equal heights never move (flat edge)");
        // Clamped: an out-of-range morph doesn't overshoot the parent.
        assert!((morph_height(120.0, 80.0, 2.0) - 80.0).abs() < 1e-6, "morph clamps at 1");
    }

    /// **INJECT-ASSERT (triangle budget stays bounded — the CDLOD headline):** the near
    /// view draws more triangles than the far view, but the near count is a small
    /// multiple of a single node's mesh (constant-ish, NOT the millions a flat full-res
    /// DEM grid would need). Proves the LOD actually caps the load.
    #[test]
    fn triangle_budget_is_bounded_and_scales_with_view() {
        let cd = Cdlod::new(region(), CdlodConfig::default());
        let near = cd.select(cam_over_center(region(), 200.0));
        let far = cd.select(WorldPos::from_lonlat(24.0, 59.3, 30_000.0));
        let tri_near = cd.triangle_count(near.len());
        let tri_far = cd.triangle_count(far.len());
        assert!(tri_near > tri_far, "near view is heavier ({tri_near} > {tri_far})");
        // The headline: CDLOD draws far fewer nodes than a UNIFORM finest-resolution
        // grid would (4^max_depth leaf nodes, one node-mesh each). Prove the saving is
        // real — the near selection is a fraction of the uniform-finest node count, and
        // its triangle budget is a fraction of the uniform-finest triangle count.
        let uniform_leaves = 4usize.pow(cd.config.max_depth); // full-res leaf-node count
        assert!(
            near.len() < uniform_leaves / 4,
            "CDLOD selects far fewer nodes than a uniform finest grid ({} < {}/4)",
            near.len(),
            uniform_leaves
        );
        assert!(
            tri_near < cd.triangle_count(uniform_leaves) / 4,
            "the near triangle budget is a fraction of a uniform finest grid ({tri_near} < {})",
            cd.triangle_count(uniform_leaves) / 4
        );
    }

    /// **INJECT-ASSERT (deterministic selection — FC-7):** the same camera yields a
    /// byte-identical node list twice, so a scripted flythrough is reproducible.
    #[test]
    fn selection_is_deterministic() {
        let cd = Cdlod::new(region(), CdlodConfig::default());
        let cam = WorldPos::from_lonlat(18.03, 59.28, 400.0);
        assert_eq!(cd.select(cam), cd.select(cam), "identical camera ⇒ identical selection");
    }

    /// **INJECT-ASSERT (the drape reuses the terrain builder — roadmap §6):** each
    /// selected node builds a real terrain patch over its own region at the shared grid,
    /// so a near (small) node samples finer topography than a far (large) node, the total
    /// triangle count equals the budget, and every patch has relief (not the flat quad).
    #[test]
    fn build_meshes_drapes_each_node_on_the_terrain() {
        let cd = Cdlod::new(region(), CdlodConfig::default());
        let elev = ProceduralTerrain::default();
        let cam = cam_over_center(region(), 220.0);
        let origin = FloatingOrigin::snapped(cam, 256.0);
        let meshes = cd.build_meshes(cam, &elev, origin, 1.0);
        assert!(!meshes.is_empty(), "the drape produced patches");

        let total_tris: usize = meshes.iter().map(|(_, m)| m.triangle_count()).sum();
        assert_eq!(total_tris, cd.triangle_count(meshes.len()), "total triangles == the CDLOD budget");

        // Each patch is a real (grid×grid) mesh with vertical relief.
        let g = cd.config.grid;
        for (node, m) in &meshes {
            assert_eq!(m.vertex_count(), g * g, "each node is a grid×grid patch");
            assert!(m.min.iter().chain(m.max.iter()).all(|v| v.is_finite()));
            assert!((0.0..=1.0).contains(&node.morph), "morph in [0,1]");
        }
        // The whole drape has relief (procedural terrain isn't flat).
        let any_relief = meshes.iter().any(|(_, m)| m.max[1] - m.min[1] > 1.0);
        assert!(any_relief, "the draped terrain has real topography, not a flat plane");

        #[cfg(feature = "testmatrix")]
        facett_core::testmatrix::emit(
            "facett-map3d-flight",
            "cdlod_drape",
            total_tris == cd.triangle_count(meshes.len()) && any_relief,
            &format!("patches={} total_tris={total_tris} grid={g}", meshes.len()),
        );
    }
}

// ─────────────────────────────────────────────────────────────────────────────────
// GPU seam — the vertex height-morph, proven equal to the CPU oracle (adapter-guarded)
// ─────────────────────────────────────────────────────────────────────────────────

/// The GPU CDLOD vertex height-morph compute pass (behind `wgpu`): a `@workgroup_size`
/// compute that applies [`morph_height`] to an array of `(h, h_parent)` pairs under one
/// `morph` factor, so the height-morph the terrain **vertex shader** will run is proven
/// (by readback) to be the exact same blend as the CPU selection math's oracle. Mirrors
/// the [`crate::scatter_gpu`] compute+readback+parity+self-skip pattern; the full render
/// pass into `facett-map3d`'s pipeline is the next roadmap step (L2 additive).
#[cfg(feature = "wgpu")]
pub mod gpu {
    use super::morph_height;

    /// The morph compute entry point + blend. Included so a build validates the WGSL.
    pub const CDLOD_MORPH_WGSL: &str = include_str!("cdlod_morph.wgsl");

    /// Uniform params for the morph pass: `morph` factor + element `count` (`f32`/`u32`).
    #[repr(C)]
    #[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
    struct MorphParams {
        morph: f32,
        count: u32,
        _pad: [u32; 2],
    }

    /// The GPU CDLOD height-morph pipeline. Build once, dispatch per morph factor.
    pub struct CdlodMorphGpu {
        pipeline: wgpu::ComputePipeline,
        bgl: wgpu::BindGroupLayout,
    }

    impl CdlodMorphGpu {
        /// Build the compute pipeline from `cdlod_morph.wgsl` (validates the WGSL on build).
        #[must_use]
        pub fn new(device: &wgpu::Device) -> Self {
            let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
                label: Some("flight_cdlod_morph"),
                source: wgpu::ShaderSource::Wgsl(CDLOD_MORPH_WGSL.into()),
            });
            let bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
                label: Some("flight_cdlod_morph_bgl"),
                entries: &[
                    wgpu::BindGroupLayoutEntry {
                        binding: 0,
                        visibility: wgpu::ShaderStages::COMPUTE,
                        ty: wgpu::BindingType::Buffer {
                            ty: wgpu::BufferBindingType::Uniform,
                            has_dynamic_offset: false,
                            min_binding_size: None,
                        },
                        count: None,
                    },
                    wgpu::BindGroupLayoutEntry {
                        binding: 1,
                        visibility: wgpu::ShaderStages::COMPUTE,
                        ty: wgpu::BindingType::Buffer {
                            ty: wgpu::BufferBindingType::Storage { read_only: true },
                            has_dynamic_offset: false,
                            min_binding_size: None,
                        },
                        count: None,
                    },
                    wgpu::BindGroupLayoutEntry {
                        binding: 2,
                        visibility: wgpu::ShaderStages::COMPUTE,
                        ty: wgpu::BindingType::Buffer {
                            ty: wgpu::BufferBindingType::Storage { read_only: false },
                            has_dynamic_offset: false,
                            min_binding_size: None,
                        },
                        count: None,
                    },
                ],
            });
            let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
                label: Some("flight_cdlod_morph_pipeline"),
                layout: Some(&device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
                    label: Some("flight_cdlod_morph_pll"),
                    bind_group_layouts: &[Some(&bgl)],
                    immediate_size: 0,
                })),
                module: &shader,
                entry_point: Some("cs_morph"),
                compilation_options: Default::default(),
                cache: None,
            });
            Self { pipeline, bgl }
        }

        /// Dispatch the morph over `pairs` (each `[h, h_parent]`) with `morph`, reading the
        /// morphed heights back (headless blocking `device.poll`). One `f32` out per pair.
        #[must_use]
        pub fn morph(&self, device: &wgpu::Device, queue: &wgpu::Queue, pairs: &[[f32; 2]], morph: f32) -> Vec<f32> {
            let count = pairs.len().max(1);
            let in_bytes = (count * std::mem::size_of::<[f32; 2]>()) as u64;
            let out_bytes = (count * std::mem::size_of::<f32>()) as u64;

            let in_buf = device.create_buffer(&wgpu::BufferDescriptor {
                label: Some("flight_cdlod_morph_in"),
                size: in_bytes.max(8),
                usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
                mapped_at_creation: false,
            });
            queue.write_buffer(&in_buf, 0, bytemuck::cast_slice(pairs));
            let out_buf = device.create_buffer(&wgpu::BufferDescriptor {
                label: Some("flight_cdlod_morph_out"),
                size: out_bytes.max(4),
                usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC,
                mapped_at_creation: false,
            });
            let params = MorphParams { morph, count: pairs.len() as u32, _pad: [0, 0] };
            let params_buf = device.create_buffer(&wgpu::BufferDescriptor {
                label: Some("flight_cdlod_morph_params"),
                size: std::mem::size_of::<MorphParams>() as u64,
                usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
                mapped_at_creation: false,
            });
            queue.write_buffer(&params_buf, 0, bytemuck::bytes_of(&params));

            let bind = device.create_bind_group(&wgpu::BindGroupDescriptor {
                label: Some("flight_cdlod_morph_bind"),
                layout: &self.bgl,
                entries: &[
                    wgpu::BindGroupEntry { binding: 0, resource: params_buf.as_entire_binding() },
                    wgpu::BindGroupEntry { binding: 1, resource: in_buf.as_entire_binding() },
                    wgpu::BindGroupEntry { binding: 2, resource: out_buf.as_entire_binding() },
                ],
            });

            let mut enc = device.create_command_encoder(&Default::default());
            {
                let mut cp = enc.begin_compute_pass(&wgpu::ComputePassDescriptor {
                    label: Some("flight_cdlod_morph_pass"),
                    timestamp_writes: None,
                });
                cp.set_pipeline(&self.pipeline);
                cp.set_bind_group(0, &bind, &[]);
                cp.dispatch_workgroups((count as u32).div_ceil(64), 1, 1);
            }
            queue.submit(Some(enc.finish()));

            let readback = device.create_buffer(&wgpu::BufferDescriptor {
                label: Some("flight_cdlod_morph_readback"),
                size: out_bytes.max(4),
                usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ,
                mapped_at_creation: false,
            });
            let mut enc = device.create_command_encoder(&Default::default());
            enc.copy_buffer_to_buffer(&out_buf, 0, &readback, 0, out_bytes.max(4));
            queue.submit(Some(enc.finish()));
            let slice = readback.slice(..);
            let (tx, rx) = std::sync::mpsc::channel();
            slice.map_async(wgpu::MapMode::Read, move |r| {
                let _ = tx.send(r);
            });
            device.poll(wgpu::PollType::wait_indefinitely()).ok();
            rx.recv().unwrap().unwrap();
            let data = slice.get_mapped_range().expect("wgpu: get_mapped_range on a mapped buffer");
            let out: Vec<f32> = bytemuck::cast_slice(&data)[..pairs.len()].to_vec();
            drop(data);
            readback.unmap();
            out
        }
    }

    /// The maximum per-element `|GPU − CPU|` of the morph vs the [`morph_height`] oracle.
    #[must_use]
    pub fn max_abs_error_vs_oracle(gpu: &[f32], pairs: &[[f32; 2]], morph: f32) -> f32 {
        gpu.iter()
            .zip(pairs)
            .map(|(&g, &[h, hp])| (g - morph_height(h, hp, morph)).abs())
            .fold(0.0f32, f32::max)
    }

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

        /// Headless compute device, guarded on compute support (`None` ⇒ caller self-skips).
        fn compute_device() -> Option<(wgpu::Device, wgpu::Queue)> {
            let instance = wgpu::Instance::default();
            let adapter = pollster::block_on(instance.request_adapter(&wgpu::RequestAdapterOptions {
                power_preference: wgpu::PowerPreference::default(),
                force_fallback_adapter: false,
                compatible_surface: None,
                apply_limit_buckets: false,
            }))
            .ok()?;
            if !adapter.get_downlevel_capabilities().flags.contains(wgpu::DownlevelFlags::COMPUTE_SHADERS) {
                return None;
            }
            pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor {
                label: Some("flight-cdlod-morph-proof"),
                required_features: wgpu::Features::empty(),
                required_limits: adapter.limits(),
                memory_hints: wgpu::MemoryHints::default(),
                experimental_features: wgpu::ExperimentalFeatures::disabled(),
                trace: wgpu::Trace::Off,
            }))
            .ok()
        }

        /// INJECT-ASSERT (runs GPU-less): the WGSL exposes the named entry point + the
        /// uniform block is the size the Rust `MorphParams` declares.
        #[test]
        fn shader_entry_point_and_uniform_size() {
            assert!(CDLOD_MORPH_WGSL.contains("fn cs_morph"), "compute entry present");
            assert!(CDLOD_MORPH_WGSL.contains("@workgroup_size(64)"), "workgroup size present");
            assert_eq!(std::mem::size_of::<MorphParams>(), 16, "params is one vec4-worth");
        }

        /// **GPU-PARITY PROOF (self-skips with a recorded sentinel when no adapter):** the
        /// GPU height-morph matches the CPU [`morph_height`] oracle to `f32` precision over
        /// a sweep of `(h, h_parent)` pairs and morph factors — proving the terrain vertex
        /// shader's crack-closing blend is the same function as the CPU selection math.
        #[test]
        fn gpu_morph_matches_cpu_oracle() {
            // THE GPU TURN, held before the bring-up and across the whole body —
            // `facett_core::render::gputurn`, the same one facett-core's probes and
            // facett-map's kittest files take (LAW 5: never a second lock). This
            // hand-rolled `compute_device()` took no turn, which is the shape that
            // SIGSEGV'd `facett_graph3d`'s two kittest tests 17 times in 30 runs on
            // 2026-09-03 and deadlocked `facett-core --lib` for 13 min on 2026-08-31.
            let _gpu = facett_core::render::gputurn::hold();
            let pairs: Vec<[f32; 2]> = (0..256).map(|i| [100.0 + i as f32, 100.0 - 0.5 * i as f32]).collect();
            let Some((device, queue)) = compute_device() else {
                eprintln!("[cdlod] no compute adapter — skipping GPU morph parity (CPU oracle still asserted GPU-less)");
                #[cfg(feature = "testmatrix")]
                facett_core::testmatrix::emit(
                    "facett-map3d-flight",
                    "cdlod_morph_gpu_parity",
                    true,
                    "gpu_absent skipped=true cpu_oracle_gated_elsewhere",
                );
                return;
            };
            let gpu = CdlodMorphGpu::new(&device);
            let mut worst = 0.0f32;
            for &morph in &[0.0f32, 0.25, 0.5, 0.66, 1.0] {
                let out = gpu.morph(&device, &queue, &pairs, morph);
                assert_eq!(out.len(), pairs.len(), "one morphed height per pair");
                worst = worst.max(max_abs_error_vs_oracle(&out, &pairs, morph));
            }
            #[cfg(feature = "testmatrix")]
            facett_core::testmatrix::emit(
                "facett-map3d-flight",
                "cdlod_morph_gpu_parity",
                worst < 1e-3,
                &format!("gpu_present max_abs_err={worst:.7} tol=1e-3 pairs={}", pairs.len()),
            );
            assert!(worst < 1e-3, "GPU morph == CPU oracle (worst err {worst})");
        }
    }
}

#[cfg(feature = "wgpu")]
pub use gpu::{CdlodMorphGpu, max_abs_error_vs_oracle};