facett-map3d-flight/src/extrude_gpu.rs
//! **GPU-compute building extrusion** (FLIGHT-SIM #3b / roadmap §3b) — the
//! `BuildingMeta` marshaller + `extrude.wgsl` compute pass that build extruded
//! building geometry on the GPU instead of the CPU mesh builder, with the landed CPU
//! [`facett_map3d::extrude_prism_mesh`] as the reference oracle.
//!
//! ## The shape of the pass
//! Refactor extrusion from a geometry **generator** into a data **marshaller**: the
//! CPU packs raw footprints into two tight storage buffers —
//! - a flat 2D **contour** buffer (`[f32; 2]` ring vertices, open + CCW), and
//! - a **`BuildingMeta`** buffer (`vertex_start`/`vertex_count`, cap-index slice,
//! output bases, `height`/`min_height`, `building_type`) —
//!
//! plus a flat **cap-index** buffer the §3a ear-clip precomputes (concave/H-shaped
//! footprints handled on the CPU; see [`facett_map3d::cap_triangulation`]). Then
//! `extrude.wgsl` runs **one invocation per building**, emitting walls + caps into
//! shared output vertex/normal/index buffers — the data an **indirect draw** consumes.
//!
//! ## The CPU path stays the oracle (L2 — additive, no regression)
//! Nothing in `facett-map3d`'s CPU `extrude_prism` changes; it is reused verbatim as
//! the single-prism oracle ([`facett_map3d::extrude_prism_mesh`]) and the fallback for
//! hosts with no adapter (wasm / the laptop tester / a deck pane). The compute shader
//! is a **twin** of it, proven equal by readback.
//!
//! ## Testability without a GPU (session constraint)
//! The **RED-when-broken** correctness gate runs GPU-**less**: the CPU
//! [`BuildingMarshal`] + the CPU reference extrusion [`BuildingMarshal::cpu_extrude`]
//! (which executes the exact op `extrude.wgsl` runs) is proven vertex-for-vertex equal
//! to the `extrude_prism_mesh` oracle for the same footprints — so a wrong offset,
//! centroid, height, or cap winding goes red on any box. The adapter-needing GPU
//! dispatch test ([`gpu`]) self-skips with a recorded sentinel when no adapter is
//! present (the [`crate::scatter_gpu`] / [`crate::cdlod`] pattern), and when it runs it
//! proves the GPU output equals that same CPU reference.
use facett_map3d::{Mesh3, cap_triangulation, open_cap_ring};
/// The compute entry point + extrusion. Included so a build validates the WGSL.
pub const EXTRUDE_WGSL: &str = include_str!("extrude.wgsl");
/// One building footprint to extrude: a closed/open ring in the design plane `(x, z)`,
/// a roof `height` and a base `min_height` (OSM `building:min_height`; `0.0` matches
/// the CPU `extrude_prism` bottom-at-ground), and a `building_type` styling tag.
#[derive(Clone, Debug, PartialEq)]
pub struct Footprint {
/// Ring vertices in the design plane `(x, z)` (OSM closing-duplicate tolerated).
pub ring: Vec<(f32, f32)>,
/// Roof height (design units) — the top cap sits at `y = height`.
pub height: f32,
/// Base height (design units) — the bottom cap sits at `y = min_height`.
pub min_height: f32,
/// Styling tag (passed through to the mesh consumer; geometry-independent).
pub building_type: u32,
}
impl Footprint {
/// A ground-based building (base at `y = 0`), matching the CPU `extrude_prism`.
#[must_use]
pub fn new(ring: Vec<(f32, f32)>, height: f32) -> Self {
Self { ring, height, min_height: 0.0, building_type: 0 }
}
}
/// The GPU-side per-building record — packed into the `BuildingMeta` storage buffer,
/// byte-for-byte the WGSL `struct BuildingMeta` in `extrude.wgsl`. All fields are
/// 4-byte (`u32`/`f32`); `#[repr(C)]`, size 40, no padding → `bytemuck::Pod` (the
/// derive is gated behind `wgpu`, so the marshaller compiles GPU-less).
#[repr(C)]
#[derive(Clone, Copy, Debug, PartialEq)]
#[cfg_attr(feature = "wgpu", derive(bytemuck::Pod, bytemuck::Zeroable))]
pub struct BuildingMeta {
/// Offset into the flat `contour` buffer of this building's first ring vertex.
pub vertex_start: u32,
/// `n` — this building's open, CCW ring vertex count.
pub vertex_count: u32,
/// Starting triangle index into the flat `caps` buffer (3 `u32` per triangle).
pub cap_start: u32,
/// Number of cap triangles (`n − 2` for a clean ring).
pub cap_count: u32,
/// Base index into the output vertex/normal buffers (this building emits `6n`).
pub out_vertex_base: u32,
/// Base index into the output index buffer (this building emits `(4n−4)·3`).
pub out_index_base: u32,
/// Roof height (design units).
pub height: f32,
/// Base height (design units).
pub min_height: f32,
/// Styling tag.
pub building_type: u32,
/// Padding to keep the struct 4-byte-clean for `bytemuck::Pod` (10 × 4 = 40 B).
pub _pad: u32,
}
/// The marshalled buffers for a batch of footprints: the flat contour vertices, the
/// flat cap-triangle indices (from the shared §3a ear-clip), the per-building
/// [`BuildingMeta`] records, and the output-buffer totals. This is the whole CPU-side
/// job of the GPU path — no geometry is generated here.
#[derive(Clone, Debug, Default)]
pub struct BuildingMarshal {
/// Flat 2D contour buffer: every building's open, CCW ring vertices concatenated.
pub contour: Vec<[f32; 2]>,
/// Flat cap-triangle indices (LOCAL to each building's ring, 3 `u32` per triangle).
pub caps: Vec<u32>,
/// Per-building records (skips degenerate `< 3`-vertex footprints).
pub meta: Vec<BuildingMeta>,
/// Total output vertices across all buildings (`Σ 6n`).
pub total_vertices: usize,
/// Total output indices across all buildings (`Σ (4n−4)·3`).
pub total_indices: usize,
}
impl BuildingMarshal {
/// Pack `footprints` into the storage buffers: open + CCW-normalise each ring (via
/// [`open_cap_ring`], the exact contour the caps index), precompute its cap
/// triangulation (via [`cap_triangulation`], the shared robust ear-clip), and
/// stamp its [`BuildingMeta`] with prefix-summed output bases. Degenerate
/// (`< 3`-vertex / collinear) footprints are dropped (no meta entry). Deterministic.
#[must_use]
pub fn from_footprints(footprints: &[Footprint]) -> Self {
let mut m = Self::default();
let (mut vbase, mut ibase, mut cap_tri) = (0u32, 0u32, 0u32);
for fp in footprints {
let ring = open_cap_ring(&fp.ring);
let n = ring.len();
if n < 3 {
continue; // degenerate footprint — no geometry
}
let cap = cap_triangulation(&fp.ring);
if cap.is_empty() {
continue; // collinear / zero-area ring
}
let vertex_start = m.contour.len() as u32;
for &(x, z) in &ring {
m.contour.push([x, z]);
}
let cap_start = cap_tri;
for t in &cap {
m.caps.extend_from_slice(&[t[0], t[1], t[2]]);
}
cap_tri += cap.len() as u32;
let out_verts = 6 * n as u32; // top n + bottom n + 4n walls
let out_idx = (12 * n as u32).saturating_sub(12); // (4n−4)·3
m.meta.push(BuildingMeta {
vertex_start,
vertex_count: n as u32,
cap_start,
cap_count: cap.len() as u32,
out_vertex_base: vbase,
out_index_base: ibase,
height: fp.height,
min_height: fp.min_height,
building_type: fp.building_type,
_pad: 0,
});
vbase += out_verts;
ibase += out_idx;
}
m.total_vertices = vbase as usize;
m.total_indices = ibase as usize;
m
}
/// The **CPU reference extrusion** — executes the exact op `extrude.wgsl` runs, on
/// the CPU, over the marshalled buffers, producing the same [`Mesh3`] the GPU pass
/// reads back. This is the GPU-less oracle-parity gate: proven equal to
/// [`extrude_prism_mesh`] per footprint (so it validates the marshalling) AND the
/// target the GPU readback is compared against.
#[must_use]
pub fn cpu_extrude(&self) -> Mesh3 {
let mut pos = vec![[0.0f32; 3]; self.total_vertices];
let mut nrm = vec![[0.0f32; 3]; self.total_vertices];
let mut idx = vec![0u32; self.total_indices];
for m in &self.meta {
extrude_one(&self.contour, &self.caps, m, &mut pos, &mut nrm, &mut idx);
}
// Every vertex this pass emits is a BUILDING wall or cap — that is the whole op
// (`extrude.wgsl`). See `facett_map3d::extrude::Mesh3::building`.
let building = vec![true; pos.len()];
Mesh3 { positions: pos, normals: nrm, colors: Vec::new(), indices: idx, building }
}
/// The design-plane `(x, z)` bounding box + `(min_height, height)` of one building's
/// output vertices in `mesh` (the marshalled AABB the render/cull path wants). Used
/// by the test to assert each extruded prism spans exactly its footprint × height.
#[must_use]
pub fn building_aabb(mesh: &Mesh3, m: &BuildingMeta) -> ([f32; 3], [f32; 3]) {
let base = m.out_vertex_base as usize;
let end = base + 6 * m.vertex_count as usize;
let mut lo = [f32::INFINITY; 3];
let mut hi = [f32::NEG_INFINITY; 3];
for p in &mesh.positions[base..end.min(mesh.positions.len())] {
for k in 0..3 {
lo[k] = lo[k].min(p[k]);
hi[k] = hi[k].max(p[k]);
}
}
(lo, hi)
}
}
/// The single-building extrusion op — the CPU statement of exactly what `cs_extrude`
/// in `extrude.wgsl` does (top ring, bottom ring, wall quads, cap + wall indices),
/// writing into the shared output buffers at `m`'s precomputed bases. Kept in one
/// place so the CPU reference and the WGSL twin cannot drift.
fn extrude_one(contour: &[[f32; 2]], caps: &[u32], m: &BuildingMeta, pos: &mut [[f32; 3]], nrm: &mut [[f32; 3]], idx: &mut [u32]) {
let n = m.vertex_count as usize;
if n < 3 {
return;
}
let vs = m.vertex_start as usize;
let vb = m.out_vertex_base as usize;
let (h, mh) = (m.height, m.min_height);
// Footprint centroid (mean of ring verts) — the outward wall-normal sense.
let (mut cx, mut cz) = (0.0f32, 0.0f32);
for i in 0..n {
cx += contour[vs + i][0];
cz += contour[vs + i][1];
}
let (cx, cz) = (cx / n as f32, cz / n as f32);
// Top ring (y = h, +Y) + bottom ring (y = mh, −Y).
for i in 0..n {
let [x, z] = contour[vs + i];
pos[vb + i] = [x, h, z];
nrm[vb + i] = [0.0, 1.0, 0.0];
pos[vb + n + i] = [x, mh, z];
nrm[vb + n + i] = [0.0, -1.0, 0.0];
}
// Side walls — one quad (4 own verts, one flat outward normal) per rim edge.
let wall_base = vb + 2 * n;
for i in 0..n {
let [x0, z0] = contour[vs + i];
let [x1, z1] = contour[vs + (i + 1) % n];
let (ex, ez) = (x1 - x0, z1 - z0);
let elen = (ex * ex + ez * ez).sqrt().max(1e-6);
let (mut nx, mut nz) = (ez / elen, -ex / elen);
let (mx, mz) = ((x0 + x1) * 0.5 - cx, (z0 + z1) * 0.5 - cz);
if nx * mx + nz * mz < 0.0 {
nx = -nx;
nz = -nz;
}
let wn = [nx, 0.0, nz];
let q = wall_base + 4 * i;
pos[q] = [x0, h, z0];
nrm[q] = wn; // i0t
pos[q + 1] = [x0, mh, z0];
nrm[q + 1] = wn; // i0b
pos[q + 2] = [x1, h, z1];
nrm[q + 2] = wn; // i1t
pos[q + 3] = [x1, mh, z1];
nrm[q + 3] = wn; // i1b
}
// Indices: top cap (as-is) · bottom cap (reversed) · wall quads.
let mut w = m.out_index_base as usize;
let vb32 = m.out_vertex_base;
let n32 = m.vertex_count;
for t in 0..m.cap_count as usize {
let ci = (m.cap_start as usize + t) * 3;
idx[w] = vb32 + caps[ci];
idx[w + 1] = vb32 + caps[ci + 1];
idx[w + 2] = vb32 + caps[ci + 2];
w += 3;
}
for t in 0..m.cap_count as usize {
let ci = (m.cap_start as usize + t) * 3;
idx[w] = vb32 + n32 + caps[ci];
idx[w + 1] = vb32 + n32 + caps[ci + 2];
idx[w + 2] = vb32 + n32 + caps[ci + 1];
w += 3;
}
for i in 0..n {
let q = (wall_base + 4 * i) as u32;
idx[w] = q;
idx[w + 1] = q + 2;
idx[w + 2] = q + 1;
idx[w + 3] = q + 2;
idx[w + 4] = q + 3;
idx[w + 5] = q + 1;
w += 6;
}
}
/// Geometry-only parity between two [`Mesh3`]es (the GPU output vs the CPU reference,
/// or the reference vs the oracle): position/normal max-abs error + exact index match.
#[derive(Clone, Copy, Debug)]
pub struct GeomParity {
/// Both meshes carry the same vertex count.
pub vertex_count_match: bool,
/// Both index buffers are element-for-element equal.
pub index_match: bool,
/// Max `|a − b|` over every position component (`INFINITY` if counts differ).
pub max_pos_err: f32,
/// Max `|a − b|` over every normal component (`INFINITY` if counts differ).
pub max_nrm_err: f32,
/// Triangle count (of `a`).
pub tri_count: usize,
}
/// Compare two meshes' geometry (positions, normals, indices). Colours are ignored
/// (geometry is colour-independent). Differing vertex counts → infinite error.
#[must_use]
pub fn geom_parity(a: &Mesh3, b: &Mesh3) -> GeomParity {
let vc = a.positions.len() == b.positions.len() && a.normals.len() == b.normals.len();
let index_match = a.indices == b.indices;
if !vc {
return GeomParity { vertex_count_match: false, index_match, max_pos_err: f32::INFINITY, max_nrm_err: f32::INFINITY, tri_count: a.triangle_count() };
}
let mut mp = 0.0f32;
let mut mn = 0.0f32;
for (pa, pb) in a.positions.iter().zip(&b.positions) {
for k in 0..3 {
mp = mp.max((pa[k] - pb[k]).abs());
}
}
for (na, nb) in a.normals.iter().zip(&b.normals) {
for k in 0..3 {
mn = mn.max((na[k] - nb[k]).abs());
}
}
GeomParity { vertex_count_match: true, index_match, max_pos_err: mp, max_nrm_err: mn, tri_count: a.triangle_count() }
}
/// The GPU-compute building-extrusion pass (behind `wgpu`): dispatches `extrude.wgsl`
/// over a [`BuildingMarshal`], reads the geometry back as a [`Mesh3`], and is proven
/// equal to [`BuildingMarshal::cpu_extrude`] (⇒ to the CPU oracle). Mirrors the
/// [`crate::scatter_gpu`] compute+readback+parity+self-skip pattern; the full indirect
/// draw into `facett-map3d`'s render pipeline is the next roadmap step (L2 additive).
#[cfg(feature = "wgpu")]
pub mod gpu {
use super::{BuildingMarshal, EXTRUDE_WGSL};
use facett_map3d::Mesh3;
/// Uniform params for the extrusion pass: `count` buildings + `vec4` padding,
/// matching `Params` in `extrude.wgsl`.
#[repr(C)]
#[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
struct ExtrudeParams {
count: u32,
_pad: [u32; 3],
}
/// The GPU building-extrusion pipeline. Build once, dispatch per footprint batch.
pub struct ExtrudeGpu {
pipeline: wgpu::ComputePipeline,
bgl: wgpu::BindGroupLayout,
}
impl ExtrudeGpu {
/// Build the compute pipeline from `extrude.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_extrude"),
source: wgpu::ShaderSource::Wgsl(EXTRUDE_WGSL.into()),
});
let storage = |binding: u32, read_only: bool| wgpu::BindGroupLayoutEntry {
binding,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Storage { read_only }, has_dynamic_offset: false, min_binding_size: None },
count: None,
};
let bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: Some("flight_extrude_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,
},
storage(1, true), // contour
storage(2, true), // caps
storage(3, true), // meta
storage(4, false), // out_pos
storage(5, false), // out_nrm
storage(6, false), // out_idx
],
});
let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: Some("flight_extrude_pipeline"),
layout: Some(&device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("flight_extrude_pll"),
bind_group_layouts: &[Some(&bgl)],
immediate_size: 0,
})),
module: &shader,
entry_point: Some("cs_extrude"),
compilation_options: Default::default(),
cache: None,
});
Self { pipeline, bgl }
}
/// Dispatch the extrusion over `marshal` and read the geometry back as a
/// [`Mesh3`] (headless blocking `device.poll` — deterministic, no clock).
#[must_use]
pub fn extrude(&self, device: &wgpu::Device, queue: &wgpu::Queue, marshal: &BuildingMarshal) -> Mesh3 {
let nv = marshal.total_vertices.max(1);
let ni = marshal.total_indices.max(1);
let pos_bytes = (nv * std::mem::size_of::<[f32; 4]>()) as u64;
let idx_bytes = (ni * std::mem::size_of::<u32>()) as u64;
let mk = |label: &str, data: &[u8], extra: wgpu::BufferUsages| {
use wgpu::util::DeviceExt;
device.create_buffer_init(&wgpu::util::BufferInitDescriptor { label: Some(label), contents: data, usage: wgpu::BufferUsages::STORAGE | extra })
};
// Pad the contour/caps/meta to at least one element so empty batches bind.
let contour: Vec<[f32; 2]> = if marshal.contour.is_empty() { vec![[0.0, 0.0]] } else { marshal.contour.clone() };
let caps: Vec<u32> = if marshal.caps.is_empty() { vec![0] } else { marshal.caps.clone() };
let meta = if marshal.meta.is_empty() { vec![super::BuildingMeta { vertex_start: 0, vertex_count: 0, cap_start: 0, cap_count: 0, out_vertex_base: 0, out_index_base: 0, height: 0.0, min_height: 0.0, building_type: 0, _pad: 0 }] } else { marshal.meta.clone() };
let contour_buf = mk("flight_extrude_contour", bytemuck::cast_slice(&contour), wgpu::BufferUsages::empty());
let caps_buf = mk("flight_extrude_caps", bytemuck::cast_slice(&caps), wgpu::BufferUsages::empty());
let meta_buf = mk("flight_extrude_meta", bytemuck::cast_slice(&meta), wgpu::BufferUsages::empty());
let out_pos = device.create_buffer(&wgpu::BufferDescriptor { label: Some("flight_extrude_pos"), size: pos_bytes, usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC, mapped_at_creation: false });
let out_nrm = device.create_buffer(&wgpu::BufferDescriptor { label: Some("flight_extrude_nrm"), size: pos_bytes, usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC, mapped_at_creation: false });
let out_idx = device.create_buffer(&wgpu::BufferDescriptor { label: Some("flight_extrude_idx"), size: idx_bytes, usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC, mapped_at_creation: false });
let params = ExtrudeParams { count: marshal.meta.len() as u32, _pad: [0; 3] };
let params_buf = device.create_buffer(&wgpu::BufferDescriptor { label: Some("flight_extrude_params"), size: std::mem::size_of::<ExtrudeParams>() as u64, usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, mapped_at_creation: false });
queue.write_buffer(¶ms_buf, 0, bytemuck::bytes_of(¶ms));
let bind = device.create_bind_group(&wgpu::BindGroupDescriptor {
label: Some("flight_extrude_bind"),
layout: &self.bgl,
entries: &[
wgpu::BindGroupEntry { binding: 0, resource: params_buf.as_entire_binding() },
wgpu::BindGroupEntry { binding: 1, resource: contour_buf.as_entire_binding() },
wgpu::BindGroupEntry { binding: 2, resource: caps_buf.as_entire_binding() },
wgpu::BindGroupEntry { binding: 3, resource: meta_buf.as_entire_binding() },
wgpu::BindGroupEntry { binding: 4, resource: out_pos.as_entire_binding() },
wgpu::BindGroupEntry { binding: 5, resource: out_nrm.as_entire_binding() },
wgpu::BindGroupEntry { binding: 6, resource: out_idx.as_entire_binding() },
],
});
let mut enc = device.create_command_encoder(&Default::default());
{
let mut cp = enc.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some("flight_extrude_pass"), timestamp_writes: None });
cp.set_pipeline(&self.pipeline);
cp.set_bind_group(0, &bind, &[]);
cp.dispatch_workgroups((marshal.meta.len().max(1) as u32).div_ceil(64), 1, 1);
}
queue.submit(Some(enc.finish()));
let positions = read_vec4_xyz(device, queue, &out_pos, pos_bytes, marshal.total_vertices);
let normals = read_vec4_xyz(device, queue, &out_nrm, pos_bytes, marshal.total_vertices);
let indices = read_u32(device, queue, &out_idx, idx_bytes, marshal.total_indices);
let building = vec![true; positions.len()];
Mesh3 { positions, normals, colors: Vec::new(), indices, building }
}
}
/// Read a `vec4<f32>` storage buffer back and drop the `w` component → `[f32; 3]`.
fn read_vec4_xyz(device: &wgpu::Device, queue: &wgpu::Queue, buf: &wgpu::Buffer, bytes: u64, count: usize) -> Vec<[f32; 3]> {
let raw: Vec<[f32; 4]> = read_mapped(device, queue, buf, bytes);
raw.into_iter().take(count).map(|p| [p[0], p[1], p[2]]).collect()
}
fn read_u32(device: &wgpu::Device, queue: &wgpu::Queue, buf: &wgpu::Buffer, bytes: u64, count: usize) -> Vec<u32> {
let raw: Vec<u32> = read_mapped(device, queue, buf, bytes);
raw.into_iter().take(count).collect()
}
/// Copy a storage buffer to a mappable buffer and block-read it (headless poll).
fn read_mapped<T: bytemuck::Pod>(device: &wgpu::Device, queue: &wgpu::Queue, buf: &wgpu::Buffer, bytes: u64) -> Vec<T> {
let readback = device.create_buffer(&wgpu::BufferDescriptor { label: Some("flight_extrude_readback"), size: bytes, 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(buf, 0, &readback, 0, bytes);
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<T> = bytemuck::cast_slice(&data).to_vec();
drop(data);
readback.unmap();
out
}
#[cfg(test)]
mod tests {
use super::super::tests::sample_footprints;
use super::super::{BuildingMarshal, geom_parity};
use super::ExtrudeGpu;
/// 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-extrude-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()
}
/// **GPU-PARITY PROOF (self-skips with a recorded sentinel when no adapter).**
/// The GPU-compute extrusion (`extrude.wgsl`) produces geometry identical to the
/// CPU reference [`BuildingMarshal::cpu_extrude`] (⇒ to the `extrude_prism_mesh`
/// oracle the GPU-less gate already pinned it to): same vertex count, exact index
/// buffer, and per-component position error at `f32` precision — proving the
/// compute shader is the same extrusion as the CPU path. The CPU-less gate
/// (`marshalled_extrusion_matches_cpu_oracle`) is the RED-when-broken correctness
/// proof that runs everywhere; this one confirms it on real silicon.
#[test]
fn gpu_extrusion_matches_cpu_reference() {
// 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 marshal = BuildingMarshal::from_footprints(&sample_footprints());
let cpu = marshal.cpu_extrude();
let Some((device, queue)) = compute_device() else {
eprintln!("[extrude_gpu] no compute adapter — skipping GPU extrusion parity (CPU oracle still asserted GPU-less)");
#[cfg(feature = "testmatrix")]
facett_core::testmatrix::emit(
"facett-map3d-flight",
"extrude_gpu_dispatch_parity",
true,
"gpu_absent skipped=true cpu_oracle_gated_elsewhere",
);
return;
};
let ext = ExtrudeGpu::new(&device);
let gpu = ext.extrude(&device, &queue, &marshal);
let rep = geom_parity(&gpu, &cpu);
#[cfg(feature = "testmatrix")]
facett_core::testmatrix::emit(
"facett-map3d-flight",
"extrude_gpu_dispatch_parity",
rep.vertex_count_match && rep.index_match && rep.max_pos_err < 1e-4 && rep.max_nrm_err < 1e-3,
&format!(
"gpu_present verts={} tris={} index_match={} max_pos_err={:.7} max_nrm_err={:.7} pos_tol=1e-4 nrm_tol=1e-3",
gpu.vertex_count(), rep.tri_count, rep.index_match, rep.max_pos_err, rep.max_nrm_err
),
);
assert!(rep.vertex_count_match, "GPU emits the same vertex count as the CPU reference");
assert!(rep.index_match, "GPU index buffer is identical to the CPU reference");
assert!(rep.max_pos_err < 1e-4, "GPU positions match CPU reference (worst {})", rep.max_pos_err);
assert!(rep.max_nrm_err < 1e-3, "GPU normals match CPU reference (worst {})", rep.max_nrm_err);
}
}
}
#[cfg(feature = "wgpu")]
pub use gpu::ExtrudeGpu;
#[cfg(test)]
mod tests {
use super::*;
use facett_map3d::extrude_prism_mesh;
/// A concave L-shaped footprint (exercises the ear-clip cap + many walls) plus a
/// simple quad, both risen — a realistic little skyline batch.
pub(crate) fn sample_footprints() -> Vec<Footprint> {
let l_shape = vec![(0.0, 0.0), (0.6, 0.0), (0.6, 0.3), (0.3, 0.3), (0.3, 0.6), (0.0, 0.6), (0.0, 0.0)];
let quad = vec![(-0.5, -0.5), (-0.2, -0.5), (-0.2, -0.2), (-0.5, -0.2)];
vec![Footprint::new(l_shape, 0.4), Footprint { ring: quad, height: 0.25, min_height: 0.0, building_type: 3 }]
}
/// INJECT-ASSERT (runs GPU-less): the WGSL exposes the named entry point + the
/// Rust `BuildingMeta` is the 40-byte, 4-byte-clean record the shader declares.
#[test]
fn shader_entry_point_and_meta_size() {
assert!(EXTRUDE_WGSL.contains("fn cs_extrude"), "compute entry present");
assert!(EXTRUDE_WGSL.contains("@workgroup_size(64)"), "workgroup size present");
assert_eq!(std::mem::size_of::<BuildingMeta>(), 40, "BuildingMeta is 10 × 4 bytes");
}
/// INJECT-ASSERT (GPU-less): the marshaller packs the buffers + prefix-summed bases
/// the way `BuildingMeta` promises — contiguous contour slices, `n−2` cap tris,
/// `6n`/`(4n−4)·3` output sizes, monotone non-overlapping output bases.
#[test]
fn marshaller_packs_consistent_buffers() {
let m = BuildingMarshal::from_footprints(&sample_footprints());
assert_eq!(m.meta.len(), 2, "both footprints marshalled");
let (mut exp_v, mut exp_i) = (0u32, 0u32);
let mut contour_cursor = 0u32;
for meta in &m.meta {
let n = meta.vertex_count;
assert!(n >= 3, "non-degenerate ring");
assert_eq!(meta.cap_count, n - 2, "clean ring → n−2 cap triangles");
assert_eq!(meta.vertex_start, contour_cursor, "contiguous contour slice");
contour_cursor += n;
assert_eq!(meta.out_vertex_base, exp_v, "monotone vertex base");
assert_eq!(meta.out_index_base, exp_i, "monotone index base");
exp_v += 6 * n;
exp_i += 12 * n - 12;
}
assert_eq!(m.contour.len() as u32, contour_cursor, "contour holds every ring vertex");
assert_eq!(m.total_vertices as u32, exp_v);
assert_eq!(m.total_indices as u32, exp_i);
}
/// **RED-WHEN-BROKEN GPU-LESS ORACLE GATE.** The marshaller + CPU reference
/// extrusion reproduces the landed CPU `extrude_prism_mesh` oracle **vertex-for-
/// vertex** for every footprint: same vertex count (`6n`), same triangle count
/// (`4n−4`), positions/normals bit-equal, indices equal (after rebasing the oracle
/// onto the building's output base), and the per-building AABB spans exactly the
/// footprint × `[min_height, height]`. A wrong offset / centroid / height / cap
/// winding in the marshaller or the shared extrusion op diverges this. Runs on any
/// box (no GPU) — the equivalence the GPU dispatch then confirms.
#[test]
fn marshalled_extrusion_matches_cpu_oracle() {
let footprints = sample_footprints();
let marshal = BuildingMarshal::from_footprints(&footprints);
let gpu_ref = marshal.cpu_extrude();
let mut worst_pos = 0.0f32;
let mut worst_nrm = 0.0f32;
for (fp, meta) in footprints.iter().zip(&marshal.meta) {
// The landed CPU oracle for this single footprint.
let oracle = extrude_prism_mesh(&fp.ring, fp.height, [200, 180, 160]);
let n = meta.vertex_count as usize;
assert_eq!(oracle.vertex_count(), 6 * n, "oracle emits 6n verts");
assert_eq!(oracle.triangle_count(), 4 * n - 4, "oracle emits 4n−4 tris");
// Slice this building out of the marshalled reference mesh.
let vb = meta.out_vertex_base as usize;
let ib = meta.out_index_base as usize;
let vcount = 6 * n;
let icount = (4 * n - 4) * 3;
let slice_pos = &gpu_ref.positions[vb..vb + vcount];
let slice_nrm = &gpu_ref.normals[vb..vb + vcount];
// Positions + normals are bit-identical to the oracle (same float math).
for (o, g) in oracle.positions.iter().zip(slice_pos) {
for k in 0..3 {
worst_pos = worst_pos.max((o[k] - g[k]).abs());
}
}
for (o, g) in oracle.normals.iter().zip(slice_nrm) {
for k in 0..3 {
worst_nrm = worst_nrm.max((o[k] - g[k]).abs());
}
}
// Indices equal once the oracle's local indices are rebased onto vb.
for (o, g) in oracle.indices.iter().zip(&gpu_ref.indices[ib..ib + icount]) {
assert_eq!(*o + vb as u32, *g, "index parity (rebased) for building");
}
// Per-building AABB: footprint bbox × [min_height, height].
let (lo, hi) = BuildingMarshal::building_aabb(&gpu_ref, meta);
let ring = open_cap_ring(&fp.ring);
let (mut fx0, mut fz0, mut fx1, mut fz1) = (f32::MAX, f32::MAX, f32::MIN, f32::MIN);
for (x, z) in ring {
fx0 = fx0.min(x);
fz0 = fz0.min(z);
fx1 = fx1.max(x);
fz1 = fz1.max(z);
}
assert!((lo[0] - fx0).abs() < 1e-5 && (hi[0] - fx1).abs() < 1e-5, "AABB x spans footprint");
assert!((lo[2] - fz0).abs() < 1e-5 && (hi[2] - fz1).abs() < 1e-5, "AABB z spans footprint");
assert!((lo[1] - fp.min_height).abs() < 1e-5, "AABB base = min_height");
assert!((hi[1] - fp.height).abs() < 1e-5, "AABB top = height");
}
#[cfg(feature = "testmatrix")]
facett_core::testmatrix::emit(
"facett-map3d-flight",
"extrude_gpu_marshal_oracle",
worst_pos < 1e-6 && worst_nrm < 1e-6,
&format!(
"buildings={} total_verts={} total_tris={} worst_pos={worst_pos:.9} worst_nrm={worst_nrm:.9} tol=1e-6",
marshal.meta.len(),
marshal.total_vertices,
marshal.total_indices / 3
),
);
assert!(worst_pos < 1e-6, "marshalled extrusion == CPU oracle positions (worst {worst_pos})");
assert!(worst_nrm < 1e-6, "marshalled extrusion == CPU oracle normals (worst {worst_nrm})");
}
/// The `min_height` seam (OSM `building:min_height`) is additive: a raised base
/// lifts the bottom cap without disturbing the roof (GPU-less).
#[test]
fn min_height_lifts_the_base() {
let ring = vec![(0.0, 0.0), (0.4, 0.0), (0.4, 0.4), (0.0, 0.4)];
let fp = Footprint { ring, height: 0.5, min_height: 0.2, building_type: 0 };
let marshal = BuildingMarshal::from_footprints(std::slice::from_ref(&fp));
let mesh = marshal.cpu_extrude();
let (lo, hi) = BuildingMarshal::building_aabb(&mesh, &marshal.meta[0]);
assert!((lo[1] - 0.2).abs() < 1e-6, "base sits at min_height");
assert!((hi[1] - 0.5).abs() < 1e-6, "roof unchanged at height");
}
}