From 7e1b64c18627d44c23dff1b64d90f14ce134fd4e Mon Sep 17 00:00:00 2001 From: Connor Johnstone Date: Fri, 16 Oct 2026 16:31:53 -0400 Subject: [PATCH] Saturation and blooming --- crates/resolved-cli/src/main.rs | 2 +- crates/resolved-cli/src/scene_file.rs | 26 +- crates/resolved/src/camera.rs | 38 +++ crates/resolved/src/fidelity.rs | 5 +- crates/resolved/src/gpu/sensor.rs | 158 +++++++---- crates/resolved/src/gpu/targets.rs | 18 ++ crates/resolved/src/lib.rs | 2 +- crates/resolved/src/render.rs | 10 +- crates/resolved/src/shaders/sensor.wgsl | 134 +++++++++- crates/resolved/tests/gpu/blooming.rs | 294 +++++++++++++++++++++ crates/resolved/tests/gpu/distortion.rs | 3 +- crates/resolved/tests/gpu/fixed_pattern.rs | 9 +- crates/resolved/tests/gpu/main.rs | 1 + crates/resolved/tests/gpu/noise.rs | 3 +- crates/resolved/tests/gpu/sensor.rs | 5 +- docs/optics.md | 51 +++- readme.md | 6 +- 17 files changed, 694 insertions(+), 71 deletions(-) create mode 100644 crates/resolved/tests/gpu/blooming.rs diff --git a/crates/resolved-cli/src/main.rs b/crates/resolved-cli/src/main.rs index 002a2ff..84fe347 100644 --- a/crates/resolved-cli/src/main.rs +++ b/crates/resolved-cli/src/main.rs @@ -60,7 +60,7 @@ struct RenderArgs { /// Camera fidelity tier: C0, or C1 (adds pixels that collect light over their area, /// for graded edges and continuous thin parts, lens distortion, lens blur, and the /// Sun, Moon and Earth where light time and aberration make them appear, detector - /// noise and fixed pattern, and at lighting L1 stars). + /// noise, fixed pattern and blooming, and at lighting L1 stars). #[arg(long, default_value = "C0")] camera: CameraTier, /// From camera C1: render each pixel as N x N finer ones (1 to 4; default 2), which diff --git a/crates/resolved-cli/src/scene_file.rs b/crates/resolved-cli/src/scene_file.rs index 4cb6548..0b76411 100644 --- a/crates/resolved-cli/src/scene_file.rs +++ b/crates/resolved-cli/src/scene_file.rs @@ -11,8 +11,8 @@ use resolved::bodies::Appearance; use resolved::nalgebra::{Quaternion, UnitQuaternion, Vector2, Vector3}; use resolved::scene::Velocities; use resolved::{ - BrownConrady, CameraConfig, LoadOptions, Material, Mesh, ModelId, Pose, Psf, Renderer, Scene, - Spectrum, Target, + Blooming, BrownConrady, CameraConfig, LoadOptions, Material, Mesh, ModelId, Pose, Psf, + Renderer, Scene, Spectrum, Target, }; use serde::Deserialize; @@ -82,6 +82,9 @@ pub struct CameraFile { pub hot_pixels: Option<[f64; 2]>, /// From camera C1: the seed of the fixed pattern. pub pattern_seed: Option, + /// From camera C1: where charge beyond the full well goes: `"none"`, + /// `{ column = efficiency }` or `{ neighbors = efficiency }`. + pub blooming: Option, /// The Gaussian blur's sigma in pixels, from camera C1: one for every channel, one /// per channel, or 0 for none. pub psf_sigma: Option, @@ -89,6 +92,15 @@ pub struct CameraFile { pub distortion: Option<[f64; 5]>, } +/// Blooming in a scene file. +#[derive(Clone, Copy, Debug, Deserialize)] +#[serde(rename_all = "lowercase")] +pub enum BloomingFile { + None, + Column(f64), + Neighbors(f64), +} + /// One sigma, or one per channel. #[derive(Debug, Deserialize)] #[serde(untagged)] @@ -191,6 +203,13 @@ impl SceneFile { if let Some(seed) = c.pattern_seed { camera.detector.pattern_seed = seed; } + if let Some(b) = c.blooming { + camera.detector.blooming = match b { + BloomingFile::None => Blooming::None, + BloomingFile::Column(efficiency) => Blooming::Column { efficiency }, + BloomingFile::Neighbors(efficiency) => Blooming::Neighbors { efficiency }, + }; + } if let Some([u, v]) = c.principal_point { camera.principal_point = Some(Vector2::new(u, v)); } @@ -471,7 +490,7 @@ mod tests { let file: SceneFile = toml::from_str(&format!( "{MINIMAL}\n[camera]\nwidth = 640\nexposure = 2e-3\nadc_bits = 10\nread_noise = 2.5\n\ dark_current = 0\nshot_noise = false\nprnu = 0.02\nhot_pixels = [1e-3, 20]\n\ - pattern_seed = 9\n" + pattern_seed = 9\nblooming = {{ column = 0.25 }}\n" )) .unwrap(); let camera = file.camera(); @@ -485,6 +504,7 @@ mod tests { ); assert_eq!((d.prnu, d.dsnu, d.pattern_seed), (0.02, 0.1, 9)); assert_eq!((d.hot_pixel_fraction, d.hot_pixel_multiple), (1e-3, 20.0)); + assert_eq!(d.blooming, Blooming::Column { efficiency: 0.25 }); for (line, psf) in [ ( diff --git a/crates/resolved/src/camera.rs b/crates/resolved/src/camera.rs index 4103187..1f7251d 100644 --- a/crates/resolved/src/camera.rs +++ b/crates/resolved/src/camera.rs @@ -54,6 +54,40 @@ pub struct Detector { pub hot_pixel_multiple: f64, /// The seed of the generated pattern: one camera, one pattern, whatever the noise. pub pattern_seed: u64, + /// Where charge beyond the full well goes. + pub blooming: Blooming, +} + +/// Where a pixel's charge beyond the full well goes, from camera tier C1. A fraction, +/// the efficiency, spreads to other pixels; the rest is drained away, as an +/// anti-blooming drain does. +#[derive(Clone, Copy, Debug, Default, PartialEq)] +pub enum Blooming { + /// All of it is lost: the pixel only clips. + #[default] + None, + /// Along the pixel's column, half each way, as in a CCD: each pixel it reaches fills + /// to the full well and passes the rest on, until it runs out or leaves the image. + Column { + /// The fraction of the excess that spreads, 0 to 1. + efficiency: f64, + }, + /// Into the 8 pixels around it in equal parts, as in a CMOS sensor; what overfills + /// them is lost. + Neighbors { + /// The fraction of the excess that spreads, 0 to 1. + efficiency: f64, + }, +} + +impl Blooming { + /// The fraction of the excess that spreads: 0 for none. + pub fn efficiency(&self) -> f64 { + match *self { + Self::None => 0.0, + Self::Column { efficiency } | Self::Neighbors { efficiency } => efficiency, + } + } } impl Detector { @@ -327,6 +361,7 @@ impl CameraConfig { hot_pixel_fraction: 1e-4, hot_pixel_multiple: 50.0, pattern_seed: 0, + blooming: Blooming::None, }, near_clip: 0.01, mount: Pose::new(Vector3::zeros(), UnitQuaternion::identity()), @@ -388,6 +423,9 @@ impl CameraConfig { if !(0.0..=1.0).contains(&d.hot_pixel_fraction) { return fail("the hot pixels' fraction must be between 0 and 1"); } + if !(0.0..=1.0).contains(&d.blooming.efficiency()) { + return fail("the blooming efficiency must be between 0 and 1"); + } Ok(()) } diff --git a/crates/resolved/src/fidelity.rs b/crates/resolved/src/fidelity.rs index 701b6be..b5145b2 100644 --- a/crates/resolved/src/fidelity.rs +++ b/crates/resolved/src/fidelity.rs @@ -50,8 +50,9 @@ pub enum CameraTier { /// ([`Scene::velocities`](crate::Scene::velocities)). Target meshes are covered at several /// points per rendered pixel, and rendered pixels on the edge of Earth, the Moon or /// the Sun cast a grid of rays, 4 x 4 and more for bodies only a few pixels across. - /// The detector adds its fixed pattern ([`FixedPattern`](crate::FixedPattern)) and - /// its shot, dark and read noise ([`Detector`](crate::camera::Detector)), seeded by + /// The detector adds its fixed pattern ([`FixedPattern`](crate::FixedPattern)), its + /// shot, dark and read noise and its blooming + /// ([`Detector`](crate::camera::Detector)), seeded by /// [`RendererOptions::noise_seed`](crate::RendererOptions::noise_seed) and the /// frame number. More of the camera joins this tier as it is modeled. C1, diff --git a/crates/resolved/src/gpu/sensor.rs b/crates/resolved/src/gpu/sensor.rs index 5356d5c..624bee8 100644 --- a/crates/resolved/src/gpu/sensor.rs +++ b/crates/resolved/src/gpu/sensor.rs @@ -10,7 +10,7 @@ use super::profile::PassTimestamps; use super::targets::FrameTargets; use super::warp::DistortionUniform; use super::GpuContext; -use crate::camera::CameraConfig; +use crate::camera::{Blooming, CameraConfig}; use crate::fixed_pattern::{FixedPattern, PATTERN_STREAM}; use crate::radiometry::Radiometry; @@ -36,6 +36,7 @@ pub struct SensorUniforms { key: [u32; 4], pattern: [f32; 4], pattern_key: [u32; 4], + bloom: [f32; 4], distortion: DistortionUniform, } @@ -76,6 +77,7 @@ impl SensorUniforms { PATTERN_STREAM, 0, ], + bloom: [d.blooming.efficiency() as f32, 0.0, 0.0, 0.0], distortion: DistortionUniform::new(k.distortion.as_ref().filter(|_| distorted)), } } @@ -88,10 +90,20 @@ impl SensorUniforms { const WORKGROUP: u32 = 256; +/// With blooming, the three steps the pass takes instead of one. +#[derive(Debug)] +struct BloomingSteps { + collect: wgpu::ComputePipeline, + /// Along columns only. + columns: Option, + read: wgpu::ComputePipeline, +} + /// The compiled sensor pipeline and its uniform buffer. #[derive(Debug)] pub struct SensorPass { pipeline: wgpu::ComputePipeline, + blooming: Option, uniforms: wgpu::Buffer, stars: bool, /// Bound for the stars' electrons without stars: the shader names the binding even @@ -103,9 +115,21 @@ pub struct SensorPass { impl SensorPass { /// Compiles the shader, adding the stars' electrons if `stars`, the detector's - /// noise if `noise`, and its fixed pattern from `pattern`, and allocates the uniform - /// buffer and uploads a loaded pattern. - pub fn new(ctx: &GpuContext, stars: bool, noise: bool, pattern: SensorPattern) -> Self { + /// noise if `noise`, its fixed pattern from `pattern` and its `blooming`, and + /// allocates the uniform buffer and uploads a loaded pattern. With blooming the + /// targets need a charge buffer ([`TargetLayout::charge`](super::targets::TargetLayout::charge)). + pub fn new( + ctx: &GpuContext, + stars: bool, + noise: bool, + pattern: SensorPattern, + blooming: Blooming, + ) -> Self { + let bloom = match blooming { + Blooming::None => 0, + Blooming::Column { .. } => 1, + Blooming::Neighbors { .. } => 2, + }; let kind = match pattern { SensorPattern::None => 0, SensorPattern::Generated => 1, @@ -119,7 +143,7 @@ impl SensorPass { [ &format!( "const STARS: bool = {stars};\nconst NOISE: bool = {noise};\n\ - const PATTERN: u32 = {kind}u;\n" + const PATTERN: u32 = {kind}u;\nconst BLOOMING: u32 = {bloom}u;\n" ), include_str!("../shaders/distortion.wgsl"), include_str!("../shaders/sensor.wgsl"), @@ -128,16 +152,23 @@ impl SensorPass { .into(), ), }); - let pipeline = ctx - .device - .create_compute_pipeline(&wgpu::ComputePipelineDescriptor { - label: Some("sensor"), - layout: None, - module: &module, - entry_point: Some("main"), - compilation_options: Default::default(), - cache: None, - }); + let pipeline = |entry: &str| { + ctx.device + .create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some(entry), + layout: None, + module: &module, + entry_point: Some(entry), + compilation_options: Default::default(), + cache: None, + }) + }; + let blooming = (bloom != 0).then(|| BloomingSteps { + collect: pipeline("collect_charge"), + columns: (bloom == 1).then(|| pipeline("bloom_columns")), + read: pipeline("read_out"), + }); + let pipeline = pipeline("main"); let uniforms = ctx.device.create_buffer(&wgpu::BufferDescriptor { label: Some("sensor uniforms"), size: std::mem::size_of::() as u64, @@ -146,6 +177,7 @@ impl SensorPass { }); Self { pipeline, + blooming, uniforms, stars, no_stars: ctx.device.create_buffer(&wgpu::BufferDescriptor { @@ -189,42 +221,74 @@ impl SensorPass { (true, None) => panic!("targets allocated without stars"), (false, _) => &self.no_stars, }; - let entries = [ - wgpu::BindGroupEntry { - binding: 0, - resource: wgpu::BindingResource::TextureView(&radiance), - }, - wgpu::BindGroupEntry { - binding: 1, - resource: targets.dn.as_entire_binding(), - }, - wgpu::BindGroupEntry { - binding: 2, - resource: self.uniforms.as_entire_binding(), - }, - wgpu::BindGroupEntry { - binding: 3, - resource: stars.as_entire_binding(), - }, - wgpu::BindGroupEntry { - binding: 4, - resource: self.pattern.as_entire_binding(), - }, + let charge = targets.charge.as_ref(); + if self.blooming.is_some() && charge.is_none() { + panic!("targets allocated without a charge buffer"); + } + let resources = [ + wgpu::BindingResource::TextureView(&radiance), + targets.dn.as_entire_binding(), + self.uniforms.as_entire_binding(), + stars.as_entire_binding(), + self.pattern.as_entire_binding(), ]; - let bind_group = ctx.device.create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("sensor"), - layout: &self.pipeline.get_bind_group_layout(0), - entries: &entries, - }); - // One invocation per u32 word, in rows of at most 65,535 workgroups. - let groups = (targets.dn.size() / 4).div_ceil(WORKGROUP as u64) as u32; - let (x, y) = (groups.min(65_535), groups.div_ceil(65_535)); + // Each step's pipeline binds only what its entry point uses. + let bind = |pipeline: &wgpu::ComputePipeline, bindings: &[u32]| { + let entries: Vec<_> = bindings + .iter() + .map(|&b| wgpu::BindGroupEntry { + binding: b, + resource: match b { + 5 | 6 => charge.expect("checked above")[b as usize - 5].as_entire_binding(), + _ => resources[b as usize].clone(), + }, + }) + .collect(); + ctx.device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("sensor"), + layout: &pipeline.get_bind_group_layout(0), + entries: &entries, + }) + }; + // In rows of at most 65,535 workgroups. + let rows = |groups: u64| { + let groups = groups as u32; + (groups.min(65_535), groups.div_ceil(65_535)) + }; + let words = (targets.dn.size() / 4).div_ceil(WORKGROUP as u64); + if let (Some(_), Some([_, overfull])) = (&self.blooming, charge) { + encoder.clear_buffer(overfull, 0, None); + } let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { label: Some("sensor"), timestamp_writes: timestamps.map(PassTimestamps::compute), }); - pass.set_pipeline(&self.pipeline); - pass.set_bind_group(0, &bind_group, &[]); - pass.dispatch_workgroups(x, y, 1); + match &self.blooming { + None => { + // One invocation per u32 word. + let (x, y) = rows(words); + pass.set_pipeline(&self.pipeline); + pass.set_bind_group(0, &bind(&self.pipeline, &[0, 1, 2, 3, 4]), &[]); + pass.dispatch_workgroups(x, y, 1); + } + Some(steps) => { + // One invocation per value, then per column and channel, then per word. + let values = targets.width as u64 * targets.height as u64 * targets.channels as u64; + let (x, y) = rows(values.div_ceil(WORKGROUP as u64)); + pass.set_pipeline(&steps.collect); + pass.set_bind_group(0, &bind(&steps.collect, &[0, 2, 3, 4, 5, 6]), &[]); + pass.dispatch_workgroups(x, y, 1); + if let Some(columns) = &steps.columns { + let lines = targets.width as u64 * targets.channels as u64; + pass.set_pipeline(columns); + pass.set_bind_group(0, &bind(columns, &[2, 5, 6]), &[]); + pass.dispatch_workgroups(lines.div_ceil(64) as u32, 1, 1); + } + let (x, y) = rows(words); + pass.set_pipeline(&steps.read); + pass.set_bind_group(0, &bind(&steps.read, &[1, 2, 5]), &[]); + pass.dispatch_workgroups(x, y, 1); + } + } } } diff --git a/crates/resolved/src/gpu/targets.rs b/crates/resolved/src/gpu/targets.rs index c896fe5..340383b 100644 --- a/crates/resolved/src/gpu/targets.rs +++ b/crates/resolved/src/gpu/targets.rs @@ -32,6 +32,8 @@ pub struct TargetLayout { pub undistorted: Option<[u32; 2]>, /// Whether stars are drawn, which needs a buffer of their electrons. pub stars: bool, + /// Whether the detector blooms, which needs a buffer of the charge before readout. + pub charge: bool, } impl TargetLayout { @@ -47,6 +49,7 @@ impl TargetLayout { margin: 0, undistorted: None, stars: false, + charge: false, } } @@ -106,6 +109,10 @@ pub struct FrameTargets { /// With stars: their electrons per pixel and channel, channels interleaved per pixel, /// in fixed point, 1/64 of an electron per unit, which the sensor pass adds. pub stars: Option, + /// With blooming: the electrons each pixel and channel holds before readout, f32, + /// channels interleaved per pixel; then a u32 flag per column and channel, set where + /// it holds charge past the full well. + pub charge: Option<[wgpu::Buffer; 2]>, /// The sensor pass's output: u16 DN, channels interleaved per pixel, rows from the /// top, packed two to a u32 word. Mappable where [`FrameTargets::dn_mappable`]. pub dn: wgpu::Buffer, @@ -202,6 +209,17 @@ impl FrameTargets { mapped_at_creation: false, }) }), + charge: layout.charge.then(|| { + let lines = layout.width as u64 * channels as u64; + [("charge", values * 4), ("overfull lines", lines * 4)].map(|(label, size)| { + ctx.device.create_buffer(&wgpu::BufferDescriptor { + label: Some(label), + size, + usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }) + }) + }), dn: ctx.device.create_buffer(&wgpu::BufferDescriptor { label: Some("dn"), size: values.div_ceil(2) * 4, diff --git a/crates/resolved/src/lib.rs b/crates/resolved/src/lib.rs index 095b4a1..f185c95 100644 --- a/crates/resolved/src/lib.rs +++ b/crates/resolved/src/lib.rs @@ -39,7 +39,7 @@ pub mod spectrum; pub mod stars; pub mod sun; -pub use camera::{BrownConrady, CameraConfig, Channel, Detector, Intrinsics, Psf}; +pub use camera::{Blooming, BrownConrady, CameraConfig, Channel, Detector, Intrinsics, Psf}; pub use earth::EarthTexture; pub use error::Error; pub use fidelity::{CameraTier, Fidelity, LightingTier}; diff --git a/crates/resolved/src/render.rs b/crates/resolved/src/render.rs index d189f85..6bda14f 100644 --- a/crates/resolved/src/render.rs +++ b/crates/resolved/src/render.rs @@ -12,7 +12,7 @@ use std::sync::mpsc; use nalgebra::Vector3; use crate::bodies::{Appearance, MOON_RADIUS}; -use crate::camera::{CameraConfig, Intrinsics, Psf}; +use crate::camera::{Blooming, CameraConfig, Intrinsics, Psf}; use std::sync::Arc; use crate::earth::EarthTexture; @@ -223,8 +223,15 @@ impl Renderer { .stars .clone() .filter(|_| area && options.fidelity.lighting >= LightingTier::L1); + // Charge beyond the full well spreads from camera C1. + let blooming = if area { + camera.detector.blooming + } else { + Blooming::None + }; let layout = TargetLayout { stars: stars.is_some(), + charge: blooming != Blooming::None, samples, scale, margin, @@ -347,6 +354,7 @@ impl Renderer { None if camera.detector.has_fixed_pattern() => SensorPattern::Generated, None => SensorPattern::None, }, + blooming, ), stars: stars.map(|catalog| { ( diff --git a/crates/resolved/src/shaders/sensor.wgsl b/crates/resolved/src/shaders/sensor.wgsl index 9e763cb..19efdb0 100644 --- a/crates/resolved/src/shaders/sensor.wgsl +++ b/crates/resolved/src/shaders/sensor.wgsl @@ -30,6 +30,8 @@ struct Sensor { pattern: vec4, // The pattern seed's low and high words, then the pattern's stream word. pattern_key: vec4, + // The fraction of the excess over the full well that spreads. + bloom: vec4, // The lens distortion (distortion.wgsl, included first). distortion: Distortion, } @@ -98,11 +100,11 @@ fn fixed_pattern(i: u32) -> vec2 { // Above this mean a Poisson count is drawn as a Gaussian of the same variance. const SMALL: f32 = 50.0; -fn value(i: u32) -> u32 { +// Value i's electrons before the full well: the light's (scaled by the fixed pattern) +// and with NOISE the dark current's, drawn with shot noise when it is on. Then, with +// NOISE, the read noise's electrons, from the same draw. +fn collect(i: u32) -> vec2 { let n = s.dims.z; - if i >= s.dims.x * s.dims.y * n { - return 0u; - } let pixel = i / n; let x = pixel % s.dims.x; let y = pixel / s.dims.x; @@ -131,7 +133,6 @@ fn value(i: u32) -> u32 { collected *= f.x; dark = f.y; } - var electrons = clamp(collected, 0.0, s.detector.y); if NOISE { let r = pcg4d(vec4(i, s.key.x, s.key.y, s.key.z)); // Two independent standard normals (Box-Muller). @@ -146,12 +147,36 @@ fn value(i: u32) -> u32 { count = poisson(mean, unit(r.z)); } } - electrons = clamp(count, 0.0, s.detector.y) + s.noise.y * radius * sin(angle); + return vec2(count, s.noise.y * radius * sin(angle)); } + return vec2(collected, 0.0); +} + +// Value i's read noise in electrons, as collect draws it. +fn read_noise(i: u32) -> f32 { + let r = pcg4d(vec4(i, s.key.x, s.key.y, s.key.z)); + let radius = sqrt(-2.0 * log(unit(r.x))); + return s.noise.y * radius * sin(6.2831853 * unit(r.y)); +} + +// Electrons read out to DN. +fn digitize(electrons: f32) -> u32 { let level = floor(electrons / s.detector.z + s.detector.w + 0.5); return u32(clamp(level, 0.0, f32(s.dims.w))); } +fn value(i: u32) -> u32 { + if i >= s.dims.x * s.dims.y * s.dims.z { + return 0u; + } + let c = collect(i); + var electrons = clamp(c.x, 0.0, s.detector.y); + if NOISE { + electrons += c.y; + } + return digitize(electrons); +} + @compute @workgroup_size(256) fn main(@builtin(global_invocation_id) id: vec3, @builtin(num_workgroups) groups: vec3) { let word = id.x + id.y * groups.x * 256u; @@ -160,3 +185,100 @@ fn main(@builtin(global_invocation_id) id: vec3, @builtin(num_workgroups) g } dn[word] = value(2u * word) | (value(2u * word + 1u) << 16u); } + +// With BLOOMING (declared by the includer: 0 for none, 1 along columns, 2 into the +// neighbors), the pass is three steps, with the charge between them in `charge`: +// collect_charge, then bloom_columns for 1, then read_out, which for 2 gathers the +// neighbors' spill as it reads. For 1, collect_charge marks each column and channel +// holding charge past the full well in `overfull`, cleared before the pass, and +// bloom_columns skips the others: in most frames, all of them. +@group(0) @binding(5) var charge: array; +@group(0) @binding(6) var overfull: array>; + +@compute @workgroup_size(256) +fn collect_charge(@builtin(global_invocation_id) id: vec3, @builtin(num_workgroups) groups: vec3) { + let i = id.x + id.y * groups.x * 256u; + if i >= s.dims.x * s.dims.y * s.dims.z { + return; + } + let c = max(collect(i).x, 0.0); + charge[i] = c; + if BLOOMING == 1u && c > s.detector.y { + let n = s.dims.z; + atomicStore(&overfull[(i / n) % s.dims.x * n + i % n], 1u); + } +} + +// One invocation per column and channel. Each pixel over the full well sends half the +// spreading fraction of its excess up and half down; going each way, a pixel takes +// what it has room for and passes the rest on. Up first, then down, which also clips +// the sources. +@compute @workgroup_size(64) +fn bloom_columns(@builtin(global_invocation_id) id: vec3) { + let n = s.dims.z; + if id.x >= s.dims.x * n || atomicLoad(&overfull[id.x]) == 0u { + return; + } + // Column id.x / n, channel id.x % n, from row 0. + let first = id.x; + let stride = s.dims.x * n; + let full = s.detector.y; + let half = 0.5 * s.bloom.x; + var carry = 0.0; + for (var y = i32(s.dims.y) - 1; y >= 0; y--) { + let i = first + u32(y) * stride; + var c = charge[i]; + let take = min(max(full - c, 0.0), carry); + c += take; + carry += half * max(c - full, 0.0) - take; + charge[i] = c; + } + carry = 0.0; + for (var y = 0u; y < s.dims.y; y++) { + let i = first + y * stride; + var c = charge[i]; + let take = min(max(full - c, 0.0), carry); + c += take; + carry += half * max(c - full, 0.0) - take; + charge[i] = min(c, full); + } +} + +fn read_value(i: u32) -> u32 { + if i >= s.dims.x * s.dims.y * s.dims.z { + return 0u; + } + var electrons = charge[i]; + if BLOOMING == 2u { + // An eighth of the spreading fraction of each neighbor's excess. + let n = s.dims.z; + let x = i32((i / n) % s.dims.x); + let y = i32((i / n) / s.dims.x); + var spill = 0.0; + for (var dy = -1; dy <= 1; dy++) { + for (var dx = -1; dx <= 1; dx++) { + let u = x + dx; + let v = y + dy; + if (dx != 0 || dy != 0) && u >= 0 && v >= 0 && u < i32(s.dims.x) && v < i32(s.dims.y) { + let j = (u32(v) * s.dims.x + u32(u)) * n + i % n; + spill += max(charge[j] - s.detector.y, 0.0); + } + } + } + electrons = min(electrons, s.detector.y) + s.bloom.x * spill / 8.0; + } + electrons = clamp(electrons, 0.0, s.detector.y); + if NOISE { + electrons += read_noise(i); + } + return digitize(electrons); +} + +@compute @workgroup_size(256) +fn read_out(@builtin(global_invocation_id) id: vec3, @builtin(num_workgroups) groups: vec3) { + let word = id.x + id.y * groups.x * 256u; + if word >= arrayLength(&dn) { + return; + } + dn[word] = read_value(2u * word) | (read_value(2u * word + 1u) << 16u); +} diff --git a/crates/resolved/tests/gpu/blooming.rs b/crates/resolved/tests/gpu/blooming.rs new file mode 100644 index 0000000..5502afd --- /dev/null +++ b/crates/resolved/tests/gpu/blooming.rs @@ -0,0 +1,294 @@ +//! Camera tier C1's blooming: charge beyond the full well spreading along a column or +//! into the neighbors, on the GPU and on the CPU (lavapipe). + +use crate::common; + +use resolved::camera::Detector; +use resolved::gpu::sensor::{SensorPass, SensorPattern, SensorUniforms}; +use resolved::gpu::targets::{FrameTargets, TargetLayout}; +use resolved::gpu::{AdapterChoice, GpuContext}; +use resolved::nalgebra::{UnitQuaternion, Vector2, Vector3}; +use resolved::radiometry::{Radiometry, AU}; +use resolved::{ + Blooming, CameraConfig, CameraTier, Fidelity, Pose, Renderer, RendererOptions, Scene, +}; + +const SIZE: u32 = 64; +const FULL: f64 = 1000.0; + +/// 64 x 64 pixels of a 1 m focal length (falloff off axis under 1e-6), reading whole +/// electrons (1 e-/DN, 16 bits, bias 100) with a full well of 1000 electrons and no +/// noise unless set. +fn camera(blooming: Blooming, noisy: bool) -> CameraConfig { + let quiet = CameraConfig::example().detector.noiseless(); + CameraConfig { + width: SIZE, + height: SIZE, + focal_length: 1.0, + detector: Detector { + gain: 1.0, + bias: 100.0, + full_well: FULL, + adc_bits: 16, + read_noise: if noisy { 3.0 } else { 0.0 }, + shot_noise: noisy, + blooming, + ..quiet + }, + ..CameraConfig::example() + } +} + +/// The value index of channel `k` at column `u`, row `v`. +fn at(u: u32, v: u32, k: u32) -> usize { + ((v * SIZE + u) * 3 + k) as usize +} + +/// The sensor pass over `electrons` per value (channels interleaved per pixel): the +/// DN less the bias. +fn sense(ctx: &GpuContext, config: &CameraConfig, electrons: &[f64]) -> Vec { + let layout = TargetLayout { + charge: config.detector.blooming != Blooming::None, + ..TargetLayout::new(SIZE, SIZE, 3) + }; + let targets = FrameTargets::new(ctx, &layout); + let radiometry = Radiometry::new(config); + let k = config.intrinsics(); + let radiance: Vec<[f32; 4]> = (0..(SIZE * SIZE) as usize) + .map(|p| { + let ray = k.ray(&Vector2::new( + (p as u32 % SIZE) as f64 + 0.5, + (p as u32 / SIZE) as f64 + 0.5, + )); + let per = radiometry.electrons(1.0, ray.z); + let e = |c: usize| (electrons[p * 3 + c] / per) as f32; + [e(0), e(1), e(2), 0.0] + }) + .collect(); + ctx.queue.write_texture( + wgpu::TexelCopyTextureInfo { + texture: &targets.radiance, + mip_level: 0, + origin: wgpu::Origin3d::ZERO, + aspect: wgpu::TextureAspect::All, + }, + bytemuck::cast_slice(&radiance), + wgpu::TexelCopyBufferLayout { + offset: 0, + bytes_per_row: Some(SIZE * 16), + rows_per_image: Some(SIZE), + }, + targets.radiance.size(), + ); + let d = config.detector; + let pass = SensorPass::new(ctx, false, d.is_noisy(), SensorPattern::None, d.blooming); + let mut uniforms = SensorUniforms::new(config, &radiometry, false); + uniforms.set_key(4, 0); + let mut encoder = ctx.device.create_command_encoder(&Default::default()); + pass.encode(ctx, &mut encoder, &targets, &uniforms, None); + ctx.queue.submit([encoder.finish()]); + targets + .read_dn(ctx) + .into_iter() + .map(|d| d as f64 - 100.0) + .collect() +} + +/// Column blooming on the CPU, as `bloom_columns` in `sensor.wgsl` does it, in f64: +/// `charge` is one column, top to bottom. +fn bloom_column(charge: &mut [f64], full: f64, efficiency: f64) { + let half = efficiency / 2.0; + let mut carry = 0.0; + for c in charge.iter_mut().rev() { + let take = (full - *c).max(0.0).min(carry); + *c += take; + carry += half * (*c - full).max(0.0) - take; + } + carry = 0.0; + for c in charge.iter_mut() { + let take = (full - *c).max(0.0).min(carry); + *c += take; + carry += half * (*c - full).max(0.0) - take; + *c = c.min(full); + } +} + +#[test] +fn an_overfull_pixel_spills_along_its_column() { + // On 400 electrons everywhere, a pixel of 10,500 (green) sends half of its 9,500 + // excess along its column, 2,375 each way: three pixels above fill to the full + // well, the fourth to 975, and the same below. One of 5,000 (red) in row 1 sends + // 1,000 each way: up, it fills row 0 and the other 400 leave the image; down, it + // fills row 2 and takes row 3 to 800. Every other value keeps its 400, and the + // column holds what it had plus the half that spread, less what left. + let config = camera(Blooming::Column { efficiency: 0.5 }, false); + let mut electrons = vec![400.0; (SIZE * SIZE * 3) as usize]; + electrons[at(20, 30, 1)] = 10_500.0; + electrons[at(40, 1, 0)] = 5_000.0; + let mut want = vec![400.0; electrons.len()]; + want[at(20, 30, 1)] = FULL; + for v in [27, 28, 29, 31, 32, 33] { + want[at(20, v, 1)] = FULL; + } + want[at(20, 26, 1)] = 975.0; + want[at(20, 34, 1)] = 975.0; + for v in [0, 1, 2] { + want[at(40, v, 0)] = FULL; + } + want[at(40, 3, 0)] = 800.0; + for ctx in common::contexts() { + let name = &ctx.info.name; + let got = sense(ctx, &config, &electrons); + assert_eq!(got, want, "{name}"); + let column: f64 = (0..SIZE).map(|v| got[at(20, v, 1)]).sum(); + assert_eq!(column, 400.0 * 63.0 + 10_500.0 - 0.5 * 9_500.0, "{name}"); + } +} + +#[test] +fn neighbors_share_the_spill_equally() { + // A pixel of 10,000 electrons on 400, spreading 40 % of its 9,000 excess: each of + // its 8 neighbors takes 450, to 850, and nothing farther changes. + let config = camera(Blooming::Neighbors { efficiency: 0.4 }, false); + let mut electrons = vec![400.0; (SIZE * SIZE * 3) as usize]; + electrons[at(30, 30, 2)] = 10_000.0; + let mut want = vec![400.0; electrons.len()]; + want[at(30, 30, 2)] = FULL; + for v in 29..=31 { + for u in 29..=31 { + if (u, v) != (30, 30) { + want[at(u, v, 2)] = 850.0; + } + } + } + for ctx in common::contexts() { + assert_eq!(sense(ctx, &config, &electrons), want, "{}", ctx.info.name); + } +} + +#[test] +fn blooming_changes_nothing_it_should_not() { + // With shot and read noise: below the full well, either kind of blooming reads + // exactly what clipping alone does, and at efficiency 0 so do pixels far past it. + let below: Vec = (0..(SIZE * SIZE * 3) as usize) + .map(|i| (i * 7919 % 900) as f64) + .collect(); + let mut over = below.clone(); + for i in (0..over.len()).step_by(37) { + over[i] = 1e5; + } + for ctx in common::contexts() { + let name = &ctx.info.name; + let clip = |e: &[f64]| sense(ctx, &camera(Blooming::None, true), e); + let (clip_below, clip_over) = (clip(&below), clip(&over)); + for blooming in [ + Blooming::Column { efficiency: 0.5 }, + Blooming::Neighbors { efficiency: 0.5 }, + ] { + let got = sense(ctx, &camera(blooming, true), &below); + assert_eq!(got, clip_below, "{name}: {blooming:?} below the full well"); + } + for blooming in [ + Blooming::Column { efficiency: 0.0 }, + Blooming::Neighbors { efficiency: 0.0 }, + ] { + let got = sense(ctx, &camera(blooming, true), &over); + assert_eq!(got, clip_over, "{name}: {blooming:?}"); + } + } +} + +#[test] +fn the_sun_streaks_as_far_as_its_charge_reaches() { + // The Sun in a 256 x 256 frame of the example camera's optics, about 26,000 full + // wells in its core, spreading 2e-5 of its excess along each column: about 20 + // pixels of streak each way from the middle of the disk, shorter toward its edges. + // Each column matches column blooming on the CPU applied to the electrons the + // radiance gives, to 1 DN, and the middle column's streak is the length the CPU + // gives it. + let _devices = common::device_lock(); + let config = CameraConfig { + width: 256, + height: 256, + detector: Detector { + blooming: Blooming::Column { efficiency: 2e-5 }, + ..CameraConfig::example().detector.noiseless() + }, + ..CameraConfig::example() + }; + let d = config.detector; + let k = config.intrinsics(); + let position = Vector3::new(4.2e7, 0.0, 0.0); + let toward = Vector3::new(0.0, 1.0, 0.0); + let observer = Pose::looking_at(position, position + toward, Vector3::z()).unwrap(); + let scene = Scene { + earth_attitude: UnitQuaternion::identity(), + sun: position + toward * AU, + moon: Vector3::new(0.0, 0.0, -3.844e8), + moon_attitude: UnitQuaternion::identity(), + targets: Vec::new(), + velocities: Default::default(), + }; + for adapter in [AdapterChoice::Gpu, AdapterChoice::Cpu] { + let options = RendererOptions { + adapter, + fidelity: Fidelity { + camera: CameraTier::C1, + ..Default::default() + }, + read_radiance: true, + ..Default::default() + }; + let mut r = Renderer::new(config.clone(), options).unwrap(); + let name = r.adapter_info().name.clone(); + let frame = r.submit(&scene, &observer).unwrap().wait().unwrap(); + let radiance = frame.radiance.as_ref().unwrap(); + let radiometry = Radiometry::new(&config); + let electrons = |u: u32, v: u32, c: usize| { + let ray = k.ray(&Vector2::new(u as f64 + 0.5, v as f64 + 0.5)); + radiometry.electrons(radiance[(v * 256 + u) as usize][c] as f64, ray.z) + }; + let mut worst = 0.0f64; + let mut middle = 0; + for u in 0..256 { + let mut column: Vec = (0..256).map(|v| electrons(u, v, 1)).collect(); + let sources = column.iter().filter(|&&e| e > d.full_well).count(); + bloom_column(&mut column, d.full_well, 2e-5); + for (v, e) in column.iter().enumerate() { + let got = frame.dn_at(u, v as u32, 1) as f64; + worst = worst.max((got - d.dn(*e) as f64).abs()); + } + if u == 128 { + let full = column.iter().filter(|&&e| e >= d.full_well).count(); + middle = full - sources; + let seen = (0..256) + .filter(|&v| frame.dn_at(u, v, 1) as f64 >= d.dn(d.full_well) as f64) + .count(); + assert_eq!(seen, full, "{name}: middle column"); + } + } + assert!(worst <= 1.0, "{name}: off by {worst} DN"); + assert!((30..=50).contains(&middle), "{name}: streak of {middle} px"); + } +} + +#[test] +fn column_blooming_conserves_charge() { + // Many sources, some adjacent: spreading everything (efficiency 1) along columns + // keeps every electron that does not leave the image, to f32 rounding, which a + // frame with the sources well inside shows by its total. + let config = camera(Blooming::Column { efficiency: 1.0 }, false); + let mut electrons = vec![100.0; (SIZE * SIZE * 3) as usize]; + for (u, v) in [(10, 20), (10, 21), (11, 25), (30, 32), (50, 40)] { + electrons[at(u, v, 1)] = 3_000.0; + } + let total: f64 = electrons.iter().sum(); + for ctx in common::contexts() { + let got: f64 = sense(ctx, &config, &electrons).iter().sum(); + assert!( + (got - total).abs() < 1.0, + "{}: {got} of {total}", + ctx.info.name + ); + } +} diff --git a/crates/resolved/tests/gpu/distortion.rs b/crates/resolved/tests/gpu/distortion.rs index 1a7f60f..f5420bc 100644 --- a/crates/resolved/tests/gpu/distortion.rs +++ b/crates/resolved/tests/gpu/distortion.rs @@ -9,6 +9,7 @@ use resolved::gpu::targets::{FrameTargets, TargetLayout}; use resolved::gpu::AdapterChoice; use resolved::nalgebra::{UnitQuaternion, Vector2, Vector3}; use resolved::radiometry::{Radiometry, AU}; +use resolved::Blooming; use resolved::{ BrownConrady, CameraConfig, CameraTier, Fidelity, Frame, Pose, Renderer, RendererOptions, Scene, Target, @@ -302,7 +303,7 @@ fn the_sensor_uses_the_undistorted_line_of_sight() { targets.radiance.size(), ); let mut encoder = ctx.device.create_command_encoder(&Default::default()); - SensorPass::new(ctx, false, false, SensorPattern::None).encode( + SensorPass::new(ctx, false, false, SensorPattern::None, Blooming::None).encode( ctx, &mut encoder, &targets, diff --git a/crates/resolved/tests/gpu/fixed_pattern.rs b/crates/resolved/tests/gpu/fixed_pattern.rs index a8d5dc5..b6b3008 100644 --- a/crates/resolved/tests/gpu/fixed_pattern.rs +++ b/crates/resolved/tests/gpu/fixed_pattern.rs @@ -9,6 +9,7 @@ use resolved::gpu::sensor::{SensorPass, SensorPattern, SensorUniforms}; use resolved::gpu::targets::{FrameTargets, TargetLayout}; use resolved::gpu::GpuContext; use resolved::radiometry::Radiometry; +use resolved::Blooming; use resolved::{CameraConfig, FixedPattern}; const SIZE: u32 = 256; @@ -65,7 +66,13 @@ fn flat( }, targets.radiance.size(), ); - let pass = SensorPass::new(ctx, false, config.detector.is_noisy(), pattern); + let pass = SensorPass::new( + ctx, + false, + config.detector.is_noisy(), + pattern, + Blooming::None, + ); let mut uniforms = SensorUniforms::new(config, &radiometry, false); uniforms.set_key(seed, frame); let mut encoder = ctx.device.create_command_encoder(&Default::default()); diff --git a/crates/resolved/tests/gpu/main.rs b/crates/resolved/tests/gpu/main.rs index 2f9b208..3ee1796 100644 --- a/crates/resolved/tests/gpu/main.rs +++ b/crates/resolved/tests/gpu/main.rs @@ -6,6 +6,7 @@ mod common; mod analytic; mod apparent; +mod blooming; mod bodies; mod distortion; mod earth_texture; diff --git a/crates/resolved/tests/gpu/noise.rs b/crates/resolved/tests/gpu/noise.rs index 457fc63..af2fcee 100644 --- a/crates/resolved/tests/gpu/noise.rs +++ b/crates/resolved/tests/gpu/noise.rs @@ -13,6 +13,7 @@ use resolved::gpu::{AdapterChoice, GpuContext}; use resolved::nalgebra::{UnitQuaternion, Vector2, Vector3}; use resolved::radiometry::{Radiometry, AU}; use resolved::stars::CatalogStar; +use resolved::Blooming; use resolved::{ CameraConfig, CameraTier, EarthTexture, Fidelity, LightingTier, Pose, Renderer, RendererOptions, Scene, StarCatalog, StarPhotometry, @@ -84,7 +85,7 @@ fn flat( }, targets.radiance.size(), ); - let pass = SensorPass::new(ctx, false, true, SensorPattern::None); + let pass = SensorPass::new(ctx, false, true, SensorPattern::None, Blooming::None); let mut uniforms = SensorUniforms::new(config, &radiometry, false); uniforms.set_key(seed, frame); let mut encoder = ctx.device.create_command_encoder(&Default::default()); diff --git a/crates/resolved/tests/gpu/sensor.rs b/crates/resolved/tests/gpu/sensor.rs index 79d46dc..1098026 100644 --- a/crates/resolved/tests/gpu/sensor.rs +++ b/crates/resolved/tests/gpu/sensor.rs @@ -10,6 +10,7 @@ use resolved::gpu::targets::{FrameTargets, TargetLayout}; use resolved::gpu::GpuContext; use resolved::nalgebra::{UnitQuaternion, Vector2, Vector3}; use resolved::radiometry::{Radiometry, AU}; +use resolved::Blooming; use resolved::{CameraConfig, Pose, Scene, Spectrum}; /// Runs the sensor pass over `radiance` (example camera, rows from the top) and @@ -32,7 +33,7 @@ fn sense(ctx: &GpuContext, name: &str, radiance: &[[f32; 4]]) -> Vec { }, targets.radiance.size(), ); - let pass = SensorPass::new(ctx, false, false, SensorPattern::None); + let pass = SensorPass::new(ctx, false, false, SensorPattern::None, Blooming::None); let uniforms = SensorUniforms::new(&config, &Radiometry::new(&config), false); let mut encoder = ctx.device.create_command_encoder(&Default::default()); pass.encode(ctx, &mut encoder, &targets, &uniforms, None); @@ -164,7 +165,7 @@ fn sunlit_earth_through_the_whole_chain() { let targets = FrameTargets::new(ctx, &TargetLayout::new(config.width, config.height, 3)); let mut encoder = ctx.device.create_command_encoder(&Default::default()); BodiesPass::new(ctx, 1, false).encode(ctx, &mut encoder, &targets, &bodies, None); - SensorPass::new(ctx, false, false, SensorPattern::None).encode( + SensorPass::new(ctx, false, false, SensorPattern::None, Blooming::None).encode( ctx, &mut encoder, &targets, diff --git a/docs/optics.md b/docs/optics.md index 4ffc2c2..d88e216 100644 --- a/docs/optics.md +++ b/docs/optics.md @@ -362,7 +362,7 @@ Code: `camera.rs` (`Detector::dn`, `Detector::variance`); `gpu/sensor.rs`, - **Full well.** The maximum number of electrons a pixel holds. Beyond it the pixel saturates and reads its maximum value whatever the light. Brighter scenes do not look brighter. In real sensors, excess charge can - also spill into neighboring pixels (blooming, not modeled yet). + also spill into neighboring pixels (blooming, below). - **Gain.** Electrons per DN. With a gain of 5, every 5 electrons add 1 to the output. - **Bias.** A fixed offset added to every pixel, so that noise around zero @@ -561,6 +561,54 @@ The pattern costs nothing measurable on an RTX 3060; on lavapipe it adds about 2.3 ms at 1920 x 1080 and 6.3 ms at 3840 x 2160. A loaded pattern takes 8 bytes per value, 199 MB at 3840 x 2160 in RGB. +### Blooming + +Code: `Blooming`; `sensor.wgsl` (`BLOOMING`: `collect_charge`, +`bloom_columns`, `read_out`). + +A saturated pixel's extra charge does not vanish. Where it goes depends on +the sensor, and matters whenever something in or near the frame is far past +the full well: the Sun at the example camera's 0.1 ms is about 26,000 full +wells at the center of its disk, and a sunlit specular part can be hundreds. + +- **CCDs** move charge down columns to read it out, and the channels that + hold it are only walled off from each other across the columns. Excess + charge spills along the column, up and down, filling each pixel it + reaches before passing on: the familiar vertical streak through the Sun + or a glint, which can cross a whole frame and hide a target in it. +- **CMOS sensors** read each pixel on its own, and most have an + anti-blooming drain that carries excess charge away. What escapes it + spills into the pixels around. + +From camera tier C1, `Detector::blooming` chooses: `None` (the default: the +pixel clips and the excess is lost), `Column { efficiency }` or +`Neighbors { efficiency }`, where the efficiency is the fraction of the +excess that spreads and the rest drains. Along a column, half goes up and +half down; each pixel it reaches takes what it has room for and passes the +rest on, until it runs out or leaves the frame. Into the neighbors, each of +the 8 takes an eighth, and what overfills them drains. + +Blooming happens to the charge, so it comes after the shot noise (and the +fixed pattern) and before the read noise and the ADC. With blooming on, the +sensor pass works in three steps rather than one: it writes each value's +charge to a buffer, spreads it, and reads it out. Along columns the spread +is a walk up and down each column, one at a time; the first step marks the +columns that hold any excess, and the walk skips the rest, which in most +frames is all of them. The CPU reference in the tests does the same walk in +f64. + +Measured, on the GPU and lavapipe: a single overfull pixel on an even +background fills the pixels it should along its column to the electron, +charge that leaves the frame is lost, and nothing else changes; the +neighbors share equally; below the full well, or at efficiency 0, blooming +changes no value. The Sun at 2e-5 efficiency, in a 256 x 256 frame of the +example optics, streaks 22 pixels each way from the middle of its disk, +matching the CPU column by column to 0 DN. Column blooming costs 0.23 ms per +1920 x 1080 frame on an RTX 3060 with nothing past the full well and 0.75 ms +with the Sun in the frame (1.0 and 2.2 ms at 3840 x 2160); into the +neighbors, 0.24 ms (1.0 ms). On lavapipe, about 3 ms and 10 to 12 ms at +1920 x 1080. + ## How shapes become pixels Two techniques turn 3D shapes into an image. The renderer uses one for the @@ -1609,7 +1657,6 @@ is a run of the script with `--gaia 16`, at twice the size. - Earthshine on the Moon's night side - The diffraction pattern's rings, and a lens's measured point spread function -- Saturation spilling into neighbors (blooming) - Non-Lambertian materials (BRDFs), glints, and the full Moon's edge as bright as its center - Clouds on Earth diff --git a/readme.md b/readme.md index 6ffaead..69254d4 100644 --- a/readme.md +++ b/readme.md @@ -206,10 +206,10 @@ pixels that collect light over their area, lens distortion (Brown-Conrady), lens blur (a Gaussian PSF) applied to the scene rendered 2 x 2 finer than the pixels, the Sun, Moon and Earth at their apparent positions (light time and aberration), detector noise (shot, dark current and read noise, seeded -and reproducible), and the detector's fixed pattern (PRNU, DSNU and hot -pixels). Planned, in order: +and reproducible), the detector's fixed pattern (PRNU, DSNU and hot pixels), +and blooming along columns or into neighbors. Planned, in order: -- Camera: saturation and blooming. +- Camera: the C1 checkpoint, then diffraction and measured PSFs (C2). - Lighting: BRDF and spectral materials, atmospheric scattering, clouds, lunar terrain. - Portable ray tracing, path tracing, and label outputs for ML.