diff --git a/core/dr-decode/src/lib.rs b/core/dr-decode/src/lib.rs
index 9a0f88d..c430b6e 100644
--- a/core/dr-decode/src/lib.rs
+++ b/core/dr-decode/src/lib.rs
@@ -481,11 +481,17 @@ pub fn decode(bytes: &[u8]) -> Result<RawImage, DecodeError> {
         data,
         cfa_pattern: cfa,
         black_level,
+        // The **maximum** of the per-channel saturation points, not the
+        // first. Taking one channel's value under-reports the others', and
+        // every sample above the level it is normalised against becomes a
+        // highlight the pipeline treats as brighter than white.
         white_level: image
             .whitelevel
             .0
-            .first()
-            .map(|v| *v as u16)
+            .iter()
+            .copied()
+            .max()
+            .map(|v| v as u16)
             .unwrap_or(u16::MAX),
         wb_coeffs,
         color_matrix,
diff --git a/core/dr-gpu/src/shaders/demosaic.wgsl b/core/dr-gpu/src/shaders/demosaic.wgsl
index 53dd18d..f86b0f9 100644
--- a/core/dr-gpu/src/shaders/demosaic.wgsl
+++ b/core/dr-gpu/src/shaders/demosaic.wgsl
@@ -102,7 +102,22 @@ fn sample(ix: i32, iy: i32) -> f32 {
     // result is clamped rather than allowed to wrap.
     let cell = (u32(cy) & 1u) * 2u + (u32(cx) & 1u);
     let value = (f32(raw_value) - params.black[cell]) * params.inv_range[cell];
-    return max(value, 0.0);
+
+    // **Clamped at the top as well, and that is what stops blown highlights
+    // going pink.** Sensors read above their declared white level — on a
+    // Canon 6D CR2 the data reaches 16383 against a white of 15070 — so a
+    // saturated pixel normalises to about 1.1 rather than 1.0.
+    //
+    // Left unclamped it survives the white balance, where red is multiplied
+    // by ~1.93 and blue by ~1.68 against green's 1.0, and then the camera
+    // matrix. Red and blue clip at the end of the pipeline; green, whose
+    // matrix row is far less positive-heavy, does not. Red and blue high with
+    // green low is magenta, and a clipped highlight came back pink.
+    //
+    // Clamping here makes a blown pixel saturate *neutrally*: all three
+    // channels reach 1.0 together and the highlight is white, which is what a
+    // blown highlight looks like and what every other developer produces.
+    return clamp(value, 0.0, 1.0);
 }
 
 @compute @workgroup_size(8, 8, 1)