facett-map3d-flight/src/scatter_gpu.rs
//! **GPU scattering-LUT precompute pass** (FLIGHT-SIM #7 / roadmap §2-GPU) — behind
//! the `wgpu` feature. Moves the [`crate::atmos::TransmittanceLut`] generation off
//! the CPU and onto the GPU: a compute pass ([`scatter_lut.wgsl`]) fills a storage
//! buffer of RGB transmittance, one invocation per `(altitude, μ)` texel, which this
//! module reads back **and** uploads into a sampled 2D `Rgba32Float` **texture** ready
//! for `sky.wgsl` to consume (the seam `atmospheric_sunsky` already documents).
//!
//! ## The CPU model stays the oracle (L2 — additive, no regression)
//! The compute shader is a *twin* of [`atmos::TransmittanceLut::precompute`], not a
//! replacement: the CPU path remains the fallback for hosts with no adapter (wasm /
//! the laptop tester / a deck pane) and the **reference** the GPU output is validated
//! against, texel-for-texel ([`GpuLut::max_abs_error_vs`]). Nothing in
//! `facett-map3d`'s landed 21-pass pipeline changes; this is a self-contained
//! precompute that produces a texture for a *later* sky-pass wiring step.
//!
//! ## Testability without a GPU (session constraint)
//! The RED-when-broken correctness check is the CPU oracle
//! (`atmos::gpu_lut_oracle_is_stable_and_physical`), which runs on the default,
//! adapter-free build. The GPU-parity test here self-skips with a recorded sentinel
//! when no compute adapter is present, so it is safe on a GPU-less box while still
//! proving equivalence wherever an adapter exists.
use crate::atmos::TransmittanceLut;
/// The compute entry point + LUT march. Included so a build validates the WGSL.
pub const SCATTER_LUT_WGSL: &str = include_str!("scatter_lut.wgsl");
/// Uniform params for the LUT compute pass: `[h_size, mu_size, steps, _pad]` (u32),
/// matching `Params` in `scatter_lut.wgsl`.
#[repr(C)]
#[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
struct LutParams {
dims: [u32; 4],
}
/// The GPU transmittance-LUT precompute pipeline. Build once ([`GpuScatterLut::new`]),
/// then [`GpuScatterLut::precompute`] as often as the atmosphere params change.
pub struct GpuScatterLut {
pipeline: wgpu::ComputePipeline,
bgl: wgpu::BindGroupLayout,
}
/// The result of a GPU precompute: the read-back RGB texels (row-major `h_size ×
/// mu_size`, the same layout as [`TransmittanceLut::texels`]) **plus** a sampled 2D
/// `Rgba32Float` texture holding them, ready for a shader to bind.
pub struct GpuLut {
/// Altitude axis size (rows).
pub h_size: usize,
/// μ axis size (columns).
pub mu_size: usize,
/// Row-major `h_size × mu_size` RGB transmittance read back from the GPU.
pub texels: Vec<[f32; 3]>,
/// The `Rgba32Float` LUT texture (`mu_size` wide × `h_size` tall) — the seam a
/// future `sky.wgsl` samples in place of the analytic gradient.
pub texture: wgpu::Texture,
/// A default view over [`GpuLut::texture`].
pub view: wgpu::TextureView,
}
/// The texel-for-texel comparison of a GPU LUT against the CPU oracle, split by
/// whether the ray points ABOVE the local horizon (μ ≥ 0). The sun-transmittance the
/// LUT is sampled for only ever reads μ ≥ 0 (a sun below the local horizon is
/// occluded → don't-care); the μ < 0 region has a genuine `sqrt`-cancellation
/// discontinuity at the horizon tangent (the near ray/ground root collapses to 0),
/// where CPU-f32 and GPU-f32 round to opposite sides of the branch. So parity is
/// judged on [`ParityReport::max_err_upper`] (must be tiny) plus a bound on how many /
/// where the below-horizon knife-edge texels diverge.
#[derive(Clone, Copy, Debug)]
pub struct ParityReport {
/// Max per-channel `|GPU − CPU|` over the μ ≥ 0 (above-horizon) texels — the
/// number the equivalence assertion gates on. `INFINITY` if the shapes differ.
pub max_err_upper: f32,
/// Max per-channel `|GPU − CPU|` over the whole grid (includes the knife-edge).
pub max_err_all: f32,
/// Count of texels whose error exceeds `1e-2` (the horizon-tangent divergences).
pub outliers: usize,
/// True iff every such outlier sits in the μ < 0 (below-horizon, occluded) band —
/// i.e. the divergence is confined to the degenerate tangent line, not a real bug.
pub outliers_all_below_horizon: bool,
}
impl GpuLut {
/// Compare this GPU LUT to a CPU [`TransmittanceLut`] of the same dimensions.
/// See [`ParityReport`]. Shapes differing is a hard failure (`max_err_* =
/// INFINITY`).
#[must_use]
pub fn parity_vs(&self, cpu: &TransmittanceLut) -> ParityReport {
if self.h_size != cpu.h_size || self.mu_size != cpu.mu_size || self.texels.len() != cpu.texels().len() {
return ParityReport {
max_err_upper: f32::INFINITY,
max_err_all: f32::INFINITY,
outliers: usize::MAX,
outliers_all_below_horizon: false,
};
}
let (mut max_err_upper, mut max_err_all, mut outliers, mut all_below) = (0.0f32, 0.0f32, 0usize, true);
for (i, (g, c)) in self.texels.iter().zip(cpu.texels()).enumerate() {
let mi = i % self.mu_size;
let mu = -1.0 + 2.0 * mi as f32 / (self.mu_size - 1) as f32;
let e = (0..3).map(|k| (g[k] - c[k]).abs()).fold(0.0, f32::max);
max_err_all = max_err_all.max(e);
if mu >= 0.0 {
max_err_upper = max_err_upper.max(e);
}
if e > 1e-2 {
outliers += 1;
if mu >= 0.0 {
all_below = false;
}
}
}
ParityReport { max_err_upper, max_err_all, outliers, outliers_all_below_horizon: all_below }
}
}
impl GpuScatterLut {
/// Build the compute pipeline from `scatter_lut.wgsl` (validates the WGSL on
/// build). Requires a device created with compute + storage support.
#[must_use]
pub fn new(device: &wgpu::Device) -> Self {
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("flight_scatter_lut"),
source: wgpu::ShaderSource::Wgsl(SCATTER_LUT_WGSL.into()),
});
let bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: Some("flight_scatter_lut_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: false },
has_dynamic_offset: false,
min_binding_size: None,
},
count: None,
},
],
});
let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: Some("flight_scatter_lut_pipeline"),
layout: Some(&device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("flight_scatter_lut_pll"),
bind_group_layouts: &[Some(&bgl)],
immediate_size: 0,
})),
module: &shader,
entry_point: Some("cs_transmittance"),
compilation_options: Default::default(),
cache: None,
});
Self { pipeline, bgl }
}
/// Dispatch the LUT precompute for `h_size × mu_size` cells with `steps` march
/// samples (the same knobs as [`TransmittanceLut::precompute`]), read the result
/// back, and upload it into a sampled `Rgba32Float` texture. Blocks on the readback
/// via `device.poll` (headless-friendly — no clock, deterministic).
#[must_use]
pub fn precompute(
&self,
device: &wgpu::Device,
queue: &wgpu::Queue,
h_size: usize,
mu_size: usize,
steps: usize,
) -> GpuLut {
let h_size = h_size.max(2);
let mu_size = mu_size.max(2);
let steps = steps.max(1);
let total = h_size * mu_size;
let out_bytes = (total * std::mem::size_of::<[f32; 4]>()) as u64;
let out_buf = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("flight_scatter_lut_out"),
size: out_bytes,
usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC,
mapped_at_creation: false,
});
let params = LutParams { dims: [h_size as u32, mu_size as u32, steps as u32, 0] };
let params_buf = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("flight_scatter_lut_params"),
size: std::mem::size_of::<LutParams>() 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_scatter_lut_bind"),
layout: &self.bgl,
entries: &[
wgpu::BindGroupEntry { binding: 0, resource: params_buf.as_entire_binding() },
wgpu::BindGroupEntry { binding: 1, 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_scatter_lut_pass"),
timestamp_writes: None,
});
cp.set_pipeline(&self.pipeline);
cp.set_bind_group(0, &bind, &[]);
cp.dispatch_workgroups((total as u32).div_ceil(64), 1, 1);
}
queue.submit(Some(enc.finish()));
// Read the storage buffer back to the CPU (headless-friendly blocking poll).
let readback = device.create_buffer(&wgpu::BufferDescriptor {
label: Some("flight_scatter_lut_readback"),
size: out_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(&out_buf, 0, &readback, 0, out_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 rgba: Vec<[f32; 4]> = bytemuck::cast_slice(&data).to_vec();
drop(data);
readback.unmap();
let texels: Vec<[f32; 3]> = rgba.iter().map(|p| [p[0], p[1], p[2]]).collect();
// Upload into a sampled Rgba32Float texture (mu_size wide × h_size tall).
let texture = device.create_texture(&wgpu::TextureDescriptor {
label: Some("flight_scatter_lut_tex"),
size: wgpu::Extent3d { width: mu_size as u32, height: h_size as u32, depth_or_array_layers: 1 },
mip_level_count: 1,
sample_count: facett_core::render::gpu::NO_MSAA_SAMPLES,
dimension: wgpu::TextureDimension::D2,
format: wgpu::TextureFormat::Rgba32Float,
usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::COPY_DST | wgpu::TextureUsages::COPY_SRC,
view_formats: &[],
});
queue.write_texture(
wgpu::TexelCopyTextureInfo {
texture: &texture,
mip_level: 0,
origin: wgpu::Origin3d::ZERO,
aspect: wgpu::TextureAspect::All,
},
bytemuck::cast_slice(&rgba),
wgpu::TexelCopyBufferLayout {
offset: 0,
bytes_per_row: Some(mu_size as u32 * std::mem::size_of::<[f32; 4]>() as u32),
rows_per_image: Some(h_size as u32),
},
wgpu::Extent3d { width: mu_size as u32, height: h_size as u32, depth_or_array_layers: 1 },
);
queue.submit(std::iter::empty());
let view = texture.create_view(&Default::default());
GpuLut { h_size, mu_size, texels, texture, view }
}
}
#[cfg(test)]
mod tests {
use super::*;
/// Headless compute device (real adapter limits so compute + storage are
/// available), guarded on compute support. `None` → the caller self-skips. Mirrors
/// the shared pattern in `facett-core::render::gpu::particles` (L5 reuse of the
/// established headless-device idiom).
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-scatter-lut-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: the WGSL exposes the entry point the pipeline names + the
/// uniform block is the size the Rust `LutParams` declares (runs GPU-less).
#[test]
fn shader_entry_point_and_uniform_size() {
assert!(SCATTER_LUT_WGSL.contains("fn cs_transmittance"), "compute entry present");
assert!(SCATTER_LUT_WGSL.contains("@workgroup_size(64)"), "workgroup size present");
assert_eq!(std::mem::size_of::<LutParams>(), 16, "params is one vec4<u32>");
}
/// **GPU-PARITY PROOF (self-skips with a recorded sentinel when no compute
/// adapter is present).** The GPU LUT precompute matches the CPU
/// [`TransmittanceLut`] oracle texel-for-texel over the sampled (above-horizon,
/// μ ≥ 0) region within a tight tolerance — proving `scatter_lut.wgsl`'s march is
/// the same function as `atmos.rs` (FC-7 parity) — while the only divergences (the
/// below-horizon horizon-tangent knife-edge, a genuine `sqrt`-cancellation
/// discontinuity in the occluded band) stay confined there and bounded. The
/// produced texture is `Rgba32Float`, `mu_size × h_size`. Emits a functional-status
/// row either way (gpu present → parity metrics; absent → the skip sentinel) so a
/// matrix row records what the tester's box could run.
#[test]
fn gpu_lut_matches_cpu_oracle_within_tolerance() {
// 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 (h, m, s) = (32usize, 128usize, 24usize);
let Some((device, queue)) = compute_device() else {
eprintln!("[scatter_gpu] no compute adapter — skipping GPU parity proof (CPU oracle still asserted GPU-less)");
#[cfg(feature = "testmatrix")]
facett_core::testmatrix::emit(
"facett-map3d-flight",
"scatter_lut_gpu_parity",
true,
"gpu_absent skipped=true cpu_oracle_gated_elsewhere",
);
return;
};
let gpu = GpuScatterLut::new(&device);
let lut = gpu.precompute(&device, &queue, h, m, s);
let cpu = TransmittanceLut::precompute(h, m, s);
// Shape + texture descriptor (the sampled GPU texture deliverable).
assert_eq!(lut.h_size, h);
assert_eq!(lut.mu_size, m);
assert_eq!(lut.texels.len(), h * m, "row-major h×m texels read back");
assert_eq!(lut.texture.format(), wgpu::TextureFormat::Rgba32Float, "LUT texture is Rgba32Float");
assert_eq!(lut.texture.width(), m as u32, "texture is mu_size wide");
assert_eq!(lut.texture.height(), h as u32, "texture is h_size tall");
// Parity: the GPU march equals the CPU oracle over the sampled region.
let rep = lut.parity_vs(&cpu);
#[cfg(feature = "testmatrix")]
facett_core::testmatrix::emit(
"facett-map3d-flight",
"scatter_lut_gpu_parity",
rep.max_err_upper < 5.0e-4 && rep.outliers_all_below_horizon && rep.outliers <= h + 2,
&format!(
"gpu_present max_err_upper={:.7} tol=5e-4 max_err_all={:.4} outliers={} all_below_horizon={} dims={h}x{m} content_hash={:#018x}",
rep.max_err_upper, rep.max_err_all, rep.outliers, rep.outliers_all_below_horizon, cpu.content_hash()
),
);
// RED-when-broken: a wrong constant / march diverges the whole grid, not just
// the tangent line, so max_err_upper explodes and/or outliers spill above μ=0.
assert!(rep.max_err_upper < 5.0e-4, "above-horizon parity: max_err_upper {} < 5e-4", rep.max_err_upper);
assert!(rep.outliers_all_below_horizon, "divergences confined below the horizon (occluded band)");
assert!(rep.outliers <= h + 2, "only the horizon-tangent knife-edge diverges ({} <= {})", rep.outliers, h + 2);
// Sanity: the readback is a real march, not all-zero / not saturated.
let overhead = lut.texels[m - 1]; // hi=0, μ=+1: high transmittance
assert!(overhead[0] > 0.5 && overhead[0] <= 1.0, "overhead texel is physical, got {}", overhead[0]);
}
}