Saturation and blooming

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

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<u64>,
+    /// From camera C1: where charge beyond the full well goes: `"none"`,
+    /// `{ column = efficiency }` or `{ neighbors = efficiency }`.
+    pub blooming: Option<BloomingFile>,
     /// 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<Sigma>,
@@ -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 [
             (

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(())
     }
 

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,

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<wgpu::ComputePipeline>,
+    read: wgpu::ComputePipeline,
+}
+
 /// The compiled sensor pipeline and its uniform buffer.
 #[derive(Debug)]
 pub struct SensorPass {
     pipeline: wgpu::ComputePipeline,
+    blooming: Option<BloomingSteps>,
     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::<SensorUniforms>() 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);
+            }
+        }
     }
 }

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<wgpu::Buffer>,
+    /// 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,

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};

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| {
                 (

crates/resolved/src/shaders/sensor.wgsl

@@ -30,6 +30,8 @@ struct Sensor {
     pattern: vec4<f32>,
     // The pattern seed's low and high words, then the pattern's stream word.
     pattern_key: vec4<u32>,
+    // The fraction of the excess over the full well that spreads.
+    bloom: vec4<f32>,
     // The lens distortion (distortion.wgsl, included first).
     distortion: Distortion,
 }
@@ -98,11 +100,11 @@ fn fixed_pattern(i: u32) -> vec2<f32> {
 // 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<f32> {
     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<u32>(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<f32>(count, s.noise.y * radius * sin(angle));
     }
+    return vec2<f32>(collected, 0.0);
+}
+
+// Value i's read noise in electrons, as collect draws it.
+fn read_noise(i: u32) -> f32 {
+    let r = pcg4d(vec4<u32>(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<u32>, @builtin(num_workgroups) groups: vec3<u32>) {
     let word = id.x + id.y * groups.x * 256u;
@@ -160,3 +185,100 @@ fn main(@builtin(global_invocation_id) id: vec3<u32>, @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<storage, read_write> charge: array<f32>;
+@group(0) @binding(6) var<storage, read_write> overfull: array<atomic<u32>>;
+
+@compute @workgroup_size(256)
+fn collect_charge(@builtin(global_invocation_id) id: vec3<u32>, @builtin(num_workgroups) groups: vec3<u32>) {
+    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<u32>) {
+    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<u32>, @builtin(num_workgroups) groups: vec3<u32>) {
+    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);
+}

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<f64> {
+    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<f64> = (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<f64> = (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
+        );
+    }
+}

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,

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());

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;

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());

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<u16> {
         },
         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,