From 5a6676ad219c5f19d7f6ecaabc720b84c4ef06d3 Mon Sep 17 00:00:00 2001
From: GF <greg.furletti@gmail.com>
Date: Mon, 13 Jul 2026 00:18:55 -0400
Subject: [PATCH 1/7] perf(native): fuse inverse MCT level shift

---
 crates/j2k-native/src/j2c/decode.rs       |   4 +-
 crates/j2k-native/src/j2c/decode/store.rs |   9 +-
 crates/j2k-native/src/j2c/mct.rs          | 222 +++++++++++++++++++---
 3 files changed, 206 insertions(+), 29 deletions(-)

diff --git a/crates/j2k-native/src/j2c/decode.rs b/crates/j2k-native/src/j2c/decode.rs
index c32e24ec..b4f9c74a 100644
--- a/crates/j2k-native/src/j2c/decode.rs
+++ b/crates/j2k-native/src/j2c/decode.rs
@@ -46,7 +46,7 @@ mod tier1;
 pub(crate) use self::allocation::DecodeAllocationBudget;
 use self::direct_plan::collect_classic_code_block_data;
 pub(crate) use self::direct_plan::{build_direct_color_plan, build_direct_grayscale_plan};
-use self::store::{apply_sign_shift, component_unsigned_level_shift, store};
+use self::store::{apply_sign_shift_after_mct, component_unsigned_level_shift, store};
 use self::subband::{code_block_required_by_index, decode_component_tile_bit_planes};
 #[cfg(all(test, feature = "parallel"))]
 use self::subband::{
@@ -138,7 +138,7 @@ pub(crate) fn decode<'a>(
     if tiles[0].mct {
         let stage_start = profile::profile_now(profile_enabled);
         mct::apply_inverse(tile_ctx, &tiles[0].component_infos, header, ht_decoder)?;
-        apply_sign_shift(tile_ctx, &header.component_infos);
+        apply_sign_shift_after_mct(tile_ctx, &header.component_infos);
         profile_timings.mct_us += profile::elapsed_us(stage_start);
     }
 
diff --git a/crates/j2k-native/src/j2c/decode/store.rs b/crates/j2k-native/src/j2c/decode/store.rs
index 610c3739..db79962f 100644
--- a/crates/j2k-native/src/j2c/decode/store.rs
+++ b/crates/j2k-native/src/j2c/decode/store.rs
@@ -5,12 +5,15 @@ use super::{
     J2kStoreComponentJob, OutputRegion, ResolutionTile, Result, Tile, TileDecodeContext,
 };
 
-pub(super) fn apply_sign_shift(
+pub(super) fn apply_sign_shift_after_mct(
     tile_ctx: &mut TileDecodeContext,
     component_infos: &[ComponentInfo],
 ) {
-    for (channel_data, component_info) in
-        tile_ctx.channel_data.iter_mut().zip(component_infos.iter())
+    for (channel_data, component_info) in tile_ctx
+        .channel_data
+        .iter_mut()
+        .zip(component_infos)
+        .skip(3)
     {
         if let Some(samples) = channel_data.integer_container.as_mut() {
             let addend = component_unsigned_level_shift_i64(component_info);
diff --git a/crates/j2k-native/src/j2c/mct.rs b/crates/j2k-native/src/j2c/mct.rs
index ea9359da..94b029b1 100644
--- a/crates/j2k-native/src/j2c/mct.rs
+++ b/crates/j2k-native/src/j2c/mct.rs
@@ -38,11 +38,27 @@ pub(crate) fn apply_inverse(
         bail!(ColorError::Mct);
     }
 
+    let addends = [
+        unsigned_level_shift(&component_infos[0]),
+        unsigned_level_shift(&component_infos[1]),
+        unsigned_level_shift(&component_infos[2]),
+    ];
+
     if s0.integer_container.is_some()
         || s1.integer_container.is_some()
         || s2.integer_container.is_some()
     {
-        return apply_inverse_i64(transform, s0, s1, s2);
+        return apply_inverse_i64(
+            transform,
+            s0,
+            s1,
+            s2,
+            [
+                unsigned_level_shift_i64(&component_infos[0]),
+                unsigned_level_shift_i64(&component_infos[1]),
+                unsigned_level_shift_i64(&component_infos[2]),
+            ],
+        );
     }
 
     let handled = if let Some(backend) = backend.as_deref_mut() {
@@ -51,9 +67,9 @@ pub(crate) fn apply_inverse(
             plane0: &mut s0.container,
             plane1: &mut s1.container,
             plane2: &mut s2.container,
-            addend0: unsigned_level_shift(&component_infos[0]),
-            addend1: unsigned_level_shift(&component_infos[1]),
-            addend2: unsigned_level_shift(&component_infos[2]),
+            addend0: addends[0],
+            addend1: addends[1],
+            addend2: addends[2],
         })?
     } else {
         false
@@ -65,6 +81,7 @@ pub(crate) fn apply_inverse(
             &mut s0.container,
             &mut s1.container,
             &mut s2.container,
+            addends,
         );
     }
 
@@ -76,6 +93,7 @@ fn apply_inverse_i64(
     s0: &mut super::ComponentData,
     s1: &mut super::ComponentData,
     s2: &mut super::ComponentData,
+    addends: [i64; 3],
 ) -> Result<()> {
     if transform != WaveletTransform::Reversible53 {
         bail!(ColorError::Mct);
@@ -97,9 +115,9 @@ fn apply_inverse_i64(
         let src1 = *y1;
         let src2 = *y2;
         let green = src0 - floor_div_i64(src2 + src1, 4);
-        *y0 = src2 + green;
-        *y1 = green;
-        *y2 = src1 + green;
+        *y0 = src2 + green + addends[0];
+        *y1 = green + addends[1];
+        *y2 = src1 + green + addends[2];
     }
 
     Ok(())
@@ -117,6 +135,14 @@ fn unsigned_level_shift(component_info: &ComponentInfo) -> f32 {
     }
 }
 
+fn unsigned_level_shift_i64(component_info: &ComponentInfo) -> i64 {
+    if component_info.size_info.signed {
+        0
+    } else {
+        1_i64 << (component_info.size_info.precision - 1)
+    }
+}
+
 #[expect(
     clippy::inline_always,
     reason = "this scalar primitive is intentionally inlined into the reversible color-transform hot loop"
@@ -133,8 +159,14 @@ fn floor_div_i64(numerator: i64, denominator: i64) -> i64 {
     }
 }
 
-fn apply_inner(transform: WaveletTransform, s0: &mut [f32], s1: &mut [f32], s2: &mut [f32]) {
-    dispatch!(Level::new(), simd => apply_inner_impl(simd, transform, s0, s1, s2));
+fn apply_inner(
+    transform: WaveletTransform,
+    s0: &mut [f32],
+    s1: &mut [f32],
+    s2: &mut [f32],
+    addends: [f32; 3],
+) {
+    dispatch!(Level::new(), simd => apply_inner_impl(simd, transform, s0, s1, s2, addends));
 }
 
 #[expect(
@@ -148,10 +180,18 @@ fn apply_inner_impl<S: Simd>(
     s0: &mut [f32],
     s1: &mut [f32],
     s2: &mut [f32],
+    addends: [f32; 3],
 ) {
     match transform {
         // Irreversible MCT, specified in G.3.
         WaveletTransform::Irreversible97 => {
+            let red_from_chroma = f32x8::splat(simd, mct::ICT_INV_R_CR);
+            let green_from_red_chroma = f32x8::splat(simd, mct::ICT_INV_G_CR);
+            let green_from_blue_chroma = f32x8::splat(simd, mct::ICT_INV_G_CB);
+            let blue_from_chroma = f32x8::splat(simd, mct::ICT_INV_B_CB);
+            let red_level = f32x8::splat(simd, addends[0]);
+            let green_level = f32x8::splat(simd, addends[1]);
+            let blue_level = f32x8::splat(simd, addends[2]);
             let mut s0_chunks = s0.chunks_exact_mut(8);
             let mut s1_chunks = s1.chunks_exact_mut(8);
             let mut s2_chunks = s2.chunks_exact_mut(8);
@@ -164,12 +204,12 @@ fn apply_inner_impl<S: Simd>(
                 let y_1 = f32x8::from_slice(simd, y1);
                 let y_2 = f32x8::from_slice(simd, y2);
 
-                let i0 = y_2.mul_add(f32x8::splat(simd, mct::ICT_INV_R_CR), y_0);
+                let i0 = y_2.mul_add(red_from_chroma, y_0) + red_level;
                 let i1 = y_2.mul_add(
-                    f32x8::splat(simd, mct::ICT_INV_G_CR),
-                    y_1.mul_add(f32x8::splat(simd, mct::ICT_INV_G_CB), y_0),
-                );
-                let i2 = y_1.mul_add(f32x8::splat(simd, mct::ICT_INV_B_CB), y_0);
+                    green_from_red_chroma,
+                    y_1.mul_add(green_from_blue_chroma, y_0),
+                ) + green_level;
+                let i2 = y_1.mul_add(blue_from_chroma, y_0) + blue_level;
 
                 i0.store(y0);
                 i1.store(y1);
@@ -184,13 +224,17 @@ fn apply_inner_impl<S: Simd>(
                 let src0 = *y0;
                 let src1 = *y1;
                 let src2 = *y2;
-                *y0 = src0 + mct::ICT_INV_R_CR * src2;
-                *y1 = src0 + mct::ICT_INV_G_CB * src1 + mct::ICT_INV_G_CR * src2;
-                *y2 = src0 + mct::ICT_INV_B_CB * src1;
+                *y0 = src0 + mct::ICT_INV_R_CR * src2 + addends[0];
+                *y1 = src0 + mct::ICT_INV_G_CB * src1 + mct::ICT_INV_G_CR * src2 + addends[1];
+                *y2 = src0 + mct::ICT_INV_B_CB * src1 + addends[2];
             }
         }
         // Reversible MCT, specified in G.2.
         WaveletTransform::Reversible53 => {
+            let quarter = f32x8::splat(simd, mct::RCT_QUARTER);
+            let red_level = f32x8::splat(simd, addends[0]);
+            let green_level = f32x8::splat(simd, addends[1]);
+            let blue_level = f32x8::splat(simd, addends[2]);
             let mut s0_chunks = s0.chunks_exact_mut(8);
             let mut s1_chunks = s1.chunks_exact_mut(8);
             let mut s2_chunks = s2.chunks_exact_mut(8);
@@ -203,12 +247,12 @@ fn apply_inner_impl<S: Simd>(
                 let y_1 = f32x8::from_slice(simd, y1);
                 let y_2 = f32x8::from_slice(simd, y2);
 
-                let i1 = y_0 - ((y_2 + y_1) * mct::RCT_QUARTER).floor();
-                let i0 = y_2 + i1;
-                let i2 = y_1 + i1;
+                let i1 = y_0 - ((y_2 + y_1) * quarter).floor();
+                let i0 = y_2 + i1 + red_level;
+                let i2 = y_1 + i1 + blue_level;
 
                 i0.store(y0);
-                i1.store(y1);
+                (i1 + green_level).store(y1);
                 i2.store(y2);
             }
             for ((y0, y1), y2) in s0_chunks
@@ -221,10 +265,140 @@ fn apply_inner_impl<S: Simd>(
                 let src1 = *y1;
                 let src2 = *y2;
                 let i1 = src0 - floor_f32((src2 + src1) * mct::RCT_QUARTER);
-                *y0 = src2 + i1;
-                *y1 = i1;
-                *y2 = src1 + i1;
+                *y0 = src2 + i1 + addends[0];
+                *y1 = i1 + addends[1];
+                *y2 = src1 + i1 + addends[2];
             }
         }
     }
 }
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+    use std::time::{Duration, Instant};
+
+    #[test]
+    fn inverse_mct_applies_mixed_addends_to_simd_chunks_and_scalar_tail() {
+        let source0 = [-15.0, -8.0, -1.0, 0.0, 1.0, 7.0, 13.0, 21.0, 34.0];
+        let source1 = [9.0, -7.0, 5.0, -3.0, 1.0, 2.0, -4.0, 6.0, -8.0];
+        let source2 = [-6.0, 4.0, -2.0, 0.0, 2.0, -4.0, 6.0, -8.0, 10.0];
+        let addends = [128.0, 0.0, 2048.0];
+
+        for transform in [
+            WaveletTransform::Reversible53,
+            WaveletTransform::Irreversible97,
+        ] {
+            let mut plane0 = source0;
+            let mut plane1 = source1;
+            let mut plane2 = source2;
+            apply_inner(transform, &mut plane0, &mut plane1, &mut plane2, addends);
+
+            for index in 0..source0.len() {
+                let (expected0, expected1, expected2) = match transform {
+                    WaveletTransform::Reversible53 => {
+                        let green = source0[index]
+                            - floor_f32((source2[index] + source1[index]) * mct::RCT_QUARTER);
+                        (source2[index] + green, green, source1[index] + green)
+                    }
+                    WaveletTransform::Irreversible97 => (
+                        source0[index] + mct::ICT_INV_R_CR * source2[index],
+                        source0[index]
+                            + mct::ICT_INV_G_CB * source1[index]
+                            + mct::ICT_INV_G_CR * source2[index],
+                        source0[index] + mct::ICT_INV_B_CB * source1[index],
+                    ),
+                };
+                let expected = [
+                    expected0 + addends[0],
+                    expected1 + addends[1],
+                    expected2 + addends[2],
+                ];
+                let actual = [plane0[index], plane1[index], plane2[index]];
+                if transform == WaveletTransform::Reversible53 {
+                    assert_eq!(actual.map(f32::to_bits), expected.map(f32::to_bits));
+                } else {
+                    for (actual, expected) in actual.into_iter().zip(expected) {
+                        assert!((actual - expected).abs() <= 0.000_25);
+                    }
+                }
+            }
+        }
+    }
+
+    #[test]
+    #[ignore = "performance guard harness; run explicitly with --ignored --nocapture"]
+    fn inverse_mct_shift_fusion_perf_guard() {
+        const LEN: usize = 512 * 512;
+        const SAMPLES: usize = 21;
+        let source0 = (0..LEN)
+            .map(|index| {
+                f32::from(u16::try_from(index % 251).expect("bounded test sample")) - 125.0
+            })
+            .collect::<Vec<_>>();
+        let source1 = (0..LEN)
+            .map(|index| f32::from(u16::try_from(index % 127).expect("bounded test sample")) - 63.0)
+            .collect::<Vec<_>>();
+        let source2 = (0..LEN)
+            .map(|index| f32::from(u16::try_from(index % 61).expect("bounded test sample")) - 30.0)
+            .collect::<Vec<_>>();
+        let mut plane0 = source0.clone();
+        let mut plane1 = source1.clone();
+        let mut plane2 = source2.clone();
+        let addends = [128.0, 128.0, 128.0];
+        let mut fused = Vec::with_capacity(SAMPLES);
+        let mut separate = Vec::with_capacity(SAMPLES);
+
+        for _ in 0..SAMPLES {
+            plane0.copy_from_slice(&source0);
+            plane1.copy_from_slice(&source1);
+            plane2.copy_from_slice(&source2);
+            let started = Instant::now();
+            apply_inner(
+                WaveletTransform::Irreversible97,
+                &mut plane0,
+                &mut plane1,
+                &mut plane2,
+                addends,
+            );
+            std::hint::black_box((&plane0, &plane1, &plane2));
+            fused.push(started.elapsed());
+
+            plane0.copy_from_slice(&source0);
+            plane1.copy_from_slice(&source1);
+            plane2.copy_from_slice(&source2);
+            let started = Instant::now();
+            apply_inner(
+                WaveletTransform::Irreversible97,
+                &mut plane0,
+                &mut plane1,
+                &mut plane2,
+                [0.0; 3],
+            );
+            for plane in [&mut plane0, &mut plane1, &mut plane2] {
+                for sample in plane {
+                    *sample += 128.0;
+                }
+            }
+            std::hint::black_box((&plane0, &plane1, &plane2));
+            separate.push(started.elapsed());
+        }
+
+        let fused = median(fused);
+        let separate = median(separate);
+        eprintln!(
+            "j2k_native_inverse_mct_shift_perf len={LEN} fused_us={} separate_us={}",
+            fused.as_micros(),
+            separate.as_micros()
+        );
+        assert!(
+            fused.as_nanos().saturating_mul(100) <= separate.as_nanos().saturating_mul(95),
+            "fused inverse MCT/sign shift must improve its targeted median by at least 5%"
+        );
+    }
+
+    fn median(mut samples: Vec<Duration>) -> Duration {
+        samples.sort_unstable();
+        samples[samples.len() / 2]
+    }
+}

From 248ece64e35e9c3a94fdfb8094fb5a148fc6d014 Mon Sep 17 00:00:00 2001
From: GF <greg.furletti@gmail.com>
Date: Mon, 13 Jul 2026 00:20:36 -0400
Subject: [PATCH 2/7] fix(metal): honor inverse MCT addends once

---
 .../src/compute/decode_dispatch/mct.rs        | 12 ++-----
 crates/j2k-metal/src/mct.rs                   | 36 +++++++++++++++++++
 2 files changed, 39 insertions(+), 9 deletions(-)

diff --git a/crates/j2k-metal/src/compute/decode_dispatch/mct.rs b/crates/j2k-metal/src/compute/decode_dispatch/mct.rs
index f72fe09d..ee5c2467 100644
--- a/crates/j2k-metal/src/compute/decode_dispatch/mct.rs
+++ b/crates/j2k-metal/src/compute/decode_dispatch/mct.rs
@@ -89,15 +89,9 @@ pub(crate) fn decode_inverse_mct(job: J2kInverseMctJob<'_>) -> Result<Vec<Buffer
         let plane0_host = checked_buffer_slice::<f32>(&plane0_buffer, len, "inverse MCT plane 0")?;
         let plane1_host = checked_buffer_slice::<f32>(&plane1_buffer, len, "inverse MCT plane 1")?;
         let plane2_host = checked_buffer_slice::<f32>(&plane2_buffer, len, "inverse MCT plane 2")?;
-        for (dst, sample) in plane0.iter_mut().zip(plane0_host.iter().copied()) {
-            *dst = sample - addend0;
-        }
-        for (dst, sample) in plane1.iter_mut().zip(plane1_host.iter().copied()) {
-            *dst = sample - addend1;
-        }
-        for (dst, sample) in plane2.iter_mut().zip(plane2_host.iter().copied()) {
-            *dst = sample - addend2;
-        }
+        plane0.copy_from_slice(&plane0_host);
+        plane1.copy_from_slice(&plane1_host);
+        plane2.copy_from_slice(&plane2_host);
         Ok(vec![plane0_buffer, plane1_buffer, plane2_buffer])
     })
 }
diff --git a/crates/j2k-metal/src/mct.rs b/crates/j2k-metal/src/mct.rs
index 900a820c..530585a3 100644
--- a/crates/j2k-metal/src/mct.rs
+++ b/crates/j2k-metal/src/mct.rs
@@ -53,6 +53,7 @@ mod tests {
     use super::MetalMctDecoder;
     use j2k_native::{
         encode, DecodeSettings, DecoderContext, EncodeOptions, HtCodeBlockDecoder, Image,
+        J2kInverseMctJob, J2kWaveletTransform,
     };
 
     #[cfg(target_os = "macos")]
@@ -80,6 +81,41 @@ mod tests {
         encode(&pixels, 2, 2, 3, 8, false, &options).expect("encode irreversible rgb8")
     }
 
+    #[test]
+    fn metal_inverse_mct_job_applies_addends_exactly_once() {
+        #[cfg(target_os = "macos")]
+        if !should_run_metal_runtime() {
+            return;
+        }
+
+        let mut plane0 = vec![0.0; 9];
+        let mut plane1 = vec![0.0; 9];
+        let mut plane2 = vec![0.0; 9];
+        let mut decoder = MetalMctDecoder::default();
+        let handled = decoder
+            .decode_inverse_mct(J2kInverseMctJob {
+                transform: J2kWaveletTransform::Reversible53,
+                plane0: &mut plane0,
+                plane1: &mut plane1,
+                plane2: &mut plane2,
+                addend0: 17.0,
+                addend1: 31.0,
+                addend2: 47.0,
+            })
+            .expect("Metal inverse MCT job");
+
+        #[cfg(target_os = "macos")]
+        assert!(handled);
+        #[cfg(not(target_os = "macos"))]
+        assert!(!handled);
+        #[cfg(target_os = "macos")]
+        {
+            assert_eq!(plane0, vec![17.0; 9]);
+            assert_eq!(plane1, vec![31.0; 9]);
+            assert_eq!(plane2, vec![47.0; 9]);
+        }
+    }
+
     #[test]
     fn metal_mct_decoder_matches_native_decode() {
         #[cfg(target_os = "macos")]

From 6a6590c5878ab74674f895636f249f6725bfbbc1 Mon Sep 17 00:00:00 2001
From: GF <greg.furletti@gmail.com>
Date: Mon, 13 Jul 2026 00:57:23 -0400
Subject: [PATCH 3/7] perf(metal): stage irreversible 9/7 IDWT

---
 crates/j2k-metal/src/compute/abi.rs           |  11 +
 .../j2k-metal/src/compute/decode_dispatch.rs  |  18 +-
 .../src/compute/decode_dispatch/idwt.rs       | 168 +++++++++++--
 crates/j2k-metal/src/compute/runtime.rs       |  14 +-
 .../j2k-metal/src/compute/symbol_inventory.rs |   5 +-
 crates/j2k-metal/src/idwt.metal               | 223 ++++++------------
 crates/j2k-metal/src/idwt.rs                  | 206 ++++++++++++++++
 .../j2k-metal/tests/metal_decode_benchmark.rs |  17 ++
 8 files changed, 474 insertions(+), 188 deletions(-)

diff --git a/crates/j2k-metal/src/compute/abi.rs b/crates/j2k-metal/src/compute/abi.rs
index 4d8ba54b..50798757 100644
--- a/crates/j2k-metal/src/compute/abi.rs
+++ b/crates/j2k-metal/src/compute/abi.rs
@@ -263,6 +263,16 @@ pub(crate) struct J2kIdwtStatus {
     pub(crate) reserved1: u32,
 }
 
+#[cfg(target_os = "macos")]
+#[repr(C)]
+#[derive(Clone, Copy)]
+pub(crate) struct J2kIdwt97StepParams {
+    pub(crate) coefficient: f32,
+    pub(crate) parity: u32,
+    pub(crate) _reserved0: u32,
+    pub(crate) _reserved1: u32,
+}
+
 #[cfg(target_os = "macos")]
 pub(crate) const J2K_MCT_STATUS_OK: u32 = 0;
 #[cfg(target_os = "macos")]
@@ -988,6 +998,7 @@ impl_gpu_readback_abi!(
     J2kIdwtSingleDecompositionParams,
     J2kRepeatedIdwtSingleDecompositionParams,
     J2kIdwtStatus,
+    J2kIdwt97StepParams,
     J2kInverseMctParams,
     J2kForwardRctParams,
     J2kForwardIctParams,
diff --git a/crates/j2k-metal/src/compute/decode_dispatch.rs b/crates/j2k-metal/src/compute/decode_dispatch.rs
index 250433a1..abd6aaf9 100644
--- a/crates/j2k-metal/src/compute/decode_dispatch.rs
+++ b/crates/j2k-metal/src/compute/decode_dispatch.rs
@@ -5,14 +5,14 @@ use super::{
     copied_slice_buffer, decode_classic_status_error, decode_idwt_status_error,
     decode_mct_status_error, dispatch_1d_pipeline, dispatch_2d_pipeline, dispatch_3d_pipeline,
     dispatch_ht_cleanup_batched_in_command_buffer, dispatch_ht_cleanup_batched_in_encoder,
-    dispatch_ht_cleanup_repeated_batched_in_command_buffer, dispatch_single_thread,
-    hybrid_stage_signpost, j2k_u32_param, label_compute_encoder, new_command_buffer,
-    new_compute_command_encoder, new_shared_buffer, prepared_ht_buffer, size_of,
-    take_classic_coefficients_scratch_buffer, take_classic_states_scratch_buffer, with_runtime,
-    zeroed_shared_buffer, Buffer, CommandBufferRef, ComputeCommandEncoderRef,
-    DirectIdwtCommandBuffers, DirectScratchBuffer, DirectStatusCheck, Error, HtCodeBlockDecodeJob,
-    HtRepeatedCleanupDispatch, J2kClassicCleanupBatchJob, J2kClassicRepeatedBatchParams,
-    J2kClassicSegment, J2kClassicStatus, J2kGrayStoreParams, J2kHtCleanupBatchJob,
+    dispatch_ht_cleanup_repeated_batched_in_command_buffer, hybrid_stage_signpost, j2k_u32_param,
+    label_compute_encoder, new_command_buffer, new_compute_command_encoder, new_shared_buffer,
+    prepared_ht_buffer, size_of, take_classic_coefficients_scratch_buffer,
+    take_classic_states_scratch_buffer, with_runtime, zeroed_shared_buffer, Buffer,
+    CommandBufferRef, ComputeCommandEncoderRef, DirectIdwtCommandBuffers, DirectScratchBuffer,
+    DirectStatusCheck, Error, HtCodeBlockDecodeJob, HtRepeatedCleanupDispatch,
+    J2kClassicCleanupBatchJob, J2kClassicRepeatedBatchParams, J2kClassicSegment, J2kClassicStatus,
+    J2kGrayStoreParams, J2kHtCleanupBatchJob, J2kIdwt97StepParams,
     J2kIdwtSingleDecompositionParams, J2kIdwtStatus, J2kInverseMctJob, J2kInverseMctParams,
     J2kMctStatus, J2kRepeatedGrayStoreParams, J2kRepeatedIdwtSingleDecompositionParams,
     J2kRepeatedStoreParams, J2kSingleDecompositionIdwtJob, J2kStoreComponentJob, J2kStoreParams,
@@ -60,6 +60,8 @@ pub(in crate::compute) use self::ht_subband::{
     encode_repeated_ht_sub_band_to_buffer_in_command_buffer, ht_batch_output_word_count,
     ht_output_word_count, required_ht_output_len,
 };
+#[cfg(test)]
+pub(crate) use self::idwt::decode_irreversible97_staged_single_decomposition_idwt;
 pub(crate) use self::idwt::{
     decode_irreversible97_single_decomposition_idwt, decode_reversible53_single_decomposition_idwt,
 };
diff --git a/crates/j2k-metal/src/compute/decode_dispatch/idwt.rs b/crates/j2k-metal/src/compute/decode_dispatch/idwt.rs
index 6a26655e..aa6b135e 100644
--- a/crates/j2k-metal/src/compute/decode_dispatch/idwt.rs
+++ b/crates/j2k-metal/src/compute/decode_dispatch/idwt.rs
@@ -2,14 +2,16 @@
 
 use super::{
     checked_buffer_read, checked_buffer_slice, commit_and_wait_metal, copied_slice_buffer,
-    decode_idwt_status_error, dispatch_2d_pipeline, dispatch_3d_pipeline, dispatch_single_thread,
-    hybrid_stage_signpost, label_compute_encoder, new_command_buffer, new_compute_command_encoder,
-    size_of, with_runtime, zeroed_shared_buffer, Buffer, CommandBufferRef,
-    ComputeCommandEncoderRef, DirectIdwtCommandBuffers, DirectStatusCheck, Error,
+    decode_idwt_status_error, dispatch_2d_pipeline, dispatch_3d_pipeline, hybrid_stage_signpost,
+    label_compute_encoder, new_command_buffer, new_compute_command_encoder, size_of, with_runtime,
+    zeroed_shared_buffer, Buffer, CommandBufferRef, ComputeCommandEncoderRef,
+    DirectIdwtCommandBuffers, DirectStatusCheck, Error, J2kIdwt97StepParams,
     J2kIdwtSingleDecompositionParams, J2kIdwtStatus, J2kRepeatedIdwtSingleDecompositionParams,
     J2kSingleDecompositionIdwtJob, MTLSize, MetalRuntime, J2K_IDWT_STATUS_OK,
     SIGNPOST_DECODE_HYBRID_IDWT_COMMAND_ENCODE,
 };
+#[cfg(target_os = "macos")]
+use j2k_codec_math::dwt;
 
 #[cfg(target_os = "macos")]
 #[expect(
@@ -373,6 +375,14 @@ pub(in crate::compute) fn dispatch_reversible53_repeated_buffers_in_command_buff
 pub(crate) fn decode_irreversible97_single_decomposition_idwt(
     job: J2kSingleDecompositionIdwtJob<'_>,
     output: &mut [f32],
+) -> Result<(), Error> {
+    decode_irreversible97_staged_single_decomposition_idwt(job, output)
+}
+
+#[cfg(target_os = "macos")]
+pub(crate) fn decode_irreversible97_staged_single_decomposition_idwt(
+    job: J2kSingleDecompositionIdwtJob<'_>,
+    output: &mut [f32],
 ) -> Result<(), Error> {
     with_runtime(|runtime| {
         let required_len = job.rect.width() as usize * job.rect.height() as usize;
@@ -416,19 +426,26 @@ pub(crate) fn decode_irreversible97_single_decomposition_idwt(
 
         let command_buffer = new_command_buffer(&runtime.queue)?;
         let encoder = new_compute_command_encoder(&command_buffer)?;
-        encoder.set_compute_pipeline_state(&runtime.idwt_irreversible97_single_decomposition);
-        encoder.set_buffer(0, Some(&ll), 0);
-        encoder.set_buffer(1, Some(&hl), 0);
-        encoder.set_buffer(2, Some(&lh), 0);
-        encoder.set_buffer(3, Some(&hh), 0);
-        encoder.set_buffer(4, Some(&decoded), 0);
-        encoder.set_bytes(
-            5,
-            size_of::<J2kIdwtSingleDecompositionParams>() as u64,
-            (&raw const params).cast(),
+        dispatch_irreversible97_single_decomposition_buffers_in_encoder_with_status(
+            &encoder,
+            SingleIdwtDispatch {
+                runtime,
+                sub_bands: IdwtSubBandBuffers {
+                    ll: &ll,
+                    ll_offset: 0,
+                    hl: &hl,
+                    hl_offset: 0,
+                    lh: &lh,
+                    lh_offset: 0,
+                    hh: &hh,
+                    hh_offset: 0,
+                },
+                params,
+                decoded: &decoded,
+                decoded_offset: 0,
+            },
+            &status_buffer,
         );
-        encoder.set_buffer(6, Some(&status_buffer), 0);
-        dispatch_single_thread(&encoder);
         encoder.end_encoding();
         commit_and_wait_metal(&command_buffer)?;
 
@@ -481,7 +498,7 @@ pub(in crate::compute) fn dispatch_irreversible97_single_decomposition_buffers_i
 pub(in crate::compute) fn dispatch_irreversible97_single_decomposition_buffers_in_encoder_with_status(
     encoder: &ComputeCommandEncoderRef,
     dispatch: SingleIdwtDispatch<'_>,
-    status_buffer: &Buffer,
+    _status_buffer: &Buffer,
 ) {
     let SingleIdwtDispatch {
         runtime,
@@ -500,7 +517,7 @@ pub(in crate::compute) fn dispatch_irreversible97_single_decomposition_buffers_i
         hh,
         hh_offset,
     } = sub_bands;
-    encoder.set_compute_pipeline_state(&runtime.idwt_irreversible97_single_decomposition);
+    encoder.set_compute_pipeline_state(&runtime.idwt_interleave);
     encoder.set_buffer(0, Some(ll), ll_offset as u64);
     encoder.set_buffer(1, Some(hl), hl_offset as u64);
     encoder.set_buffer(2, Some(lh), lh_offset as u64);
@@ -511,6 +528,117 @@ pub(in crate::compute) fn dispatch_irreversible97_single_decomposition_buffers_i
         size_of::<J2kIdwtSingleDecompositionParams>() as u64,
         (&raw const params).cast(),
     );
-    encoder.set_buffer(6, Some(status_buffer), 0);
-    dispatch_single_thread(encoder);
+    dispatch_2d_pipeline(
+        encoder,
+        &runtime.idwt_interleave,
+        (params.width, params.height),
+    );
+    encoder.memory_barrier_with_resources(&[decoded]);
+
+    dispatch_irreversible97_stages(encoder, runtime, decoded, decoded_offset, params);
+}
+
+#[cfg(target_os = "macos")]
+fn dispatch_irreversible97_stages(
+    encoder: &ComputeCommandEncoderRef,
+    runtime: &MetalRuntime,
+    decoded: &Buffer,
+    decoded_offset: usize,
+    params: J2kIdwtSingleDecompositionParams,
+) {
+    encoder.set_compute_pipeline_state(&runtime.idwt_irreversible97_horizontal_scale);
+    encoder.set_buffer(0, Some(decoded), decoded_offset as u64);
+    encoder.set_bytes(
+        1,
+        size_of::<J2kIdwtSingleDecompositionParams>() as u64,
+        (&raw const params).cast(),
+    );
+    dispatch_2d_pipeline(
+        encoder,
+        &runtime.idwt_irreversible97_horizontal_scale,
+        (params.width, params.height),
+    );
+    encoder.memory_barrier_with_resources(&[decoded]);
+
+    let first_even_x = (params.x0 + params.output_x) & 1;
+    let first_odd_x = 1 - first_even_x;
+    encoder.set_compute_pipeline_state(&runtime.idwt_irreversible97_horizontal_step);
+    encoder.set_buffer(0, Some(decoded), decoded_offset as u64);
+    encoder.set_bytes(
+        1,
+        size_of::<J2kIdwtSingleDecompositionParams>() as u64,
+        (&raw const params).cast(),
+    );
+    for (coefficient, parity) in [
+        (dwt::IDWT97_NEG_DELTA_F32, first_even_x),
+        (dwt::IDWT97_NEG_GAMMA_F32, first_odd_x),
+        (dwt::IDWT97_NEG_BETA_F32, first_even_x),
+        (dwt::IDWT97_NEG_ALPHA_F32, first_odd_x),
+    ] {
+        let step = J2kIdwt97StepParams {
+            coefficient,
+            parity,
+            _reserved0: 0,
+            _reserved1: 0,
+        };
+        encoder.set_bytes(
+            2,
+            size_of::<J2kIdwt97StepParams>() as u64,
+            (&raw const step).cast(),
+        );
+        dispatch_2d_pipeline(
+            encoder,
+            &runtime.idwt_irreversible97_horizontal_step,
+            (params.width, params.height),
+        );
+        encoder.memory_barrier_with_resources(&[decoded]);
+    }
+
+    encoder.set_compute_pipeline_state(&runtime.idwt_irreversible97_vertical_scale);
+    encoder.set_buffer(0, Some(decoded), decoded_offset as u64);
+    encoder.set_bytes(
+        1,
+        size_of::<J2kIdwtSingleDecompositionParams>() as u64,
+        (&raw const params).cast(),
+    );
+    dispatch_2d_pipeline(
+        encoder,
+        &runtime.idwt_irreversible97_vertical_scale,
+        (params.width, params.height),
+    );
+    encoder.memory_barrier_with_resources(&[decoded]);
+
+    let first_even_y = (params.y0 + params.output_y) & 1;
+    let first_odd_y = 1 - first_even_y;
+    encoder.set_compute_pipeline_state(&runtime.idwt_irreversible97_vertical_step);
+    encoder.set_buffer(0, Some(decoded), decoded_offset as u64);
+    encoder.set_bytes(
+        1,
+        size_of::<J2kIdwtSingleDecompositionParams>() as u64,
+        (&raw const params).cast(),
+    );
+    for (coefficient, parity) in [
+        (dwt::IDWT97_NEG_DELTA_F32, first_even_y),
+        (dwt::IDWT97_NEG_GAMMA_F32, first_odd_y),
+        (dwt::IDWT97_NEG_BETA_F32, first_even_y),
+        (dwt::IDWT97_NEG_ALPHA_F32, first_odd_y),
+    ] {
+        let step = J2kIdwt97StepParams {
+            coefficient,
+            parity,
+            _reserved0: 0,
+            _reserved1: 0,
+        };
+        encoder.set_bytes(
+            2,
+            size_of::<J2kIdwt97StepParams>() as u64,
+            (&raw const step).cast(),
+        );
+        dispatch_2d_pipeline(
+            encoder,
+            &runtime.idwt_irreversible97_vertical_step,
+            (params.width, params.height),
+        );
+        encoder.memory_barrier_with_resources(&[decoded]);
+    }
 }
diff --git a/crates/j2k-metal/src/compute/runtime.rs b/crates/j2k-metal/src/compute/runtime.rs
index 432f8bda..de0289d5 100644
--- a/crates/j2k-metal/src/compute/runtime.rs
+++ b/crates/j2k-metal/src/compute/runtime.rs
@@ -62,7 +62,10 @@ pub(crate) struct MetalRuntime {
     pub(super) idwt_interleave_batched: ComputePipelineState,
     pub(super) idwt_reversible53_horizontal_batched: ComputePipelineState,
     pub(super) idwt_reversible53_vertical_batched: ComputePipelineState,
-    pub(super) idwt_irreversible97_single_decomposition: ComputePipelineState,
+    pub(super) idwt_irreversible97_horizontal_scale: ComputePipelineState,
+    pub(super) idwt_irreversible97_vertical_scale: ComputePipelineState,
+    pub(super) idwt_irreversible97_horizontal_step: ComputePipelineState,
+    pub(super) idwt_irreversible97_vertical_step: ComputePipelineState,
     pub(super) fdwt53_horizontal: ComputePipelineState,
     pub(super) fdwt53_vertical: ComputePipelineState,
     pub(super) fdwt53_horizontal_batched: ComputePipelineState,
@@ -184,9 +187,14 @@ impl MetalRuntime {
             idwt_reversible53_vertical_batched: pipeline(
                 "j2k_idwt_reversible53_vertical_pass_batched",
             )?,
-            idwt_irreversible97_single_decomposition: pipeline(
-                "j2k_idwt_irreversible97_single_decomposition",
+            idwt_irreversible97_horizontal_scale: pipeline(
+                "j2k_idwt_irreversible97_horizontal_scale",
             )?,
+            idwt_irreversible97_vertical_scale: pipeline("j2k_idwt_irreversible97_vertical_scale")?,
+            idwt_irreversible97_horizontal_step: pipeline(
+                "j2k_idwt_irreversible97_horizontal_step",
+            )?,
+            idwt_irreversible97_vertical_step: pipeline("j2k_idwt_irreversible97_vertical_step")?,
             fdwt53_horizontal: pipeline("j2k_forward_dwt53_horizontal")?,
             fdwt53_vertical: pipeline("j2k_forward_dwt53_vertical")?,
             fdwt53_horizontal_batched: pipeline("j2k_forward_dwt53_horizontal_batched")?,
diff --git a/crates/j2k-metal/src/compute/symbol_inventory.rs b/crates/j2k-metal/src/compute/symbol_inventory.rs
index da416561..e862f37e 100644
--- a/crates/j2k-metal/src/compute/symbol_inventory.rs
+++ b/crates/j2k-metal/src/compute/symbol_inventory.rs
@@ -19,7 +19,8 @@ macro_rules! wire_compute_symbols {
             J2kForwardDwt53BatchedParams, J2kForwardDwt53Params, J2kForwardIctParams,
             J2kForwardRctParams, J2kGrayStoreParams, J2kHtCleanupBatchJob, J2kHtCleanupParams,
             J2kHtEncodeBatchJob, J2kHtEncodeParams, J2kHtEncodeStatus, J2kHtRepeatedBatchParams,
-            J2kHtStatus, J2kIdwtSingleDecompositionParams, J2kIdwtStatus, J2kInverseMctParams,
+            J2kHtStatus, J2kIdwt97StepParams, J2kIdwtSingleDecompositionParams, J2kIdwtStatus,
+            J2kInverseMctParams,
             J2kLosslessCodestreamAssemblyParams, J2kLosslessCoefficientJob,
             J2kLosslessDeinterleaveParams, J2kMctRgb8PackParams, J2kMctStatus, J2kPackParams,
             J2kPacketBlock, J2kPacketDescriptor, J2kPacketEncodeParams, J2kPacketEncodeStatus,
@@ -138,6 +139,8 @@ macro_rules! wire_compute_symbols {
             classic_batch_uses_plain_fast_path, classic_repeated_uses_plain_fast_path,
             repeated_gray_store_is_contiguous_full_surface,
         };
+        #[cfg(all(target_os = "macos", test))]
+        pub(crate) use crate::compute::decode_dispatch::decode_irreversible97_staged_single_decomposition_idwt;
         #[cfg(target_os = "macos")]
         pub(crate) use crate::compute::decode_dispatch::{
             decode_inverse_mct, decode_irreversible97_single_decomposition_idwt,
diff --git a/crates/j2k-metal/src/idwt.metal b/crates/j2k-metal/src/idwt.metal
index 030c411a..196e69df 100644
--- a/crates/j2k-metal/src/idwt.metal
+++ b/crates/j2k-metal/src/idwt.metal
@@ -62,6 +62,13 @@ struct J2kIdwtStatus {
     uint reserved1;
 };
 
+struct J2kIdwt97StepParams {
+    float coefficient;
+    uint parity;
+    uint reserved0;
+    uint reserved1;
+};
+
 constant uint J2K_IDWT_STATUS_OK = 0;
 constant uint J2K_IDWT_STATUS_FAIL = 1;
 
@@ -428,186 +435,90 @@ kernel void j2k_idwt_reversible53_vertical_pass_batched(
     }
 }
 
-kernel void j2k_idwt_irreversible97_single_decomposition(
-    device const float *ll [[buffer(0)]],
-    device const float *hl [[buffer(1)]],
-    device const float *lh [[buffer(2)]],
-    device const float *hh [[buffer(3)]],
-    device float *out [[buffer(4)]],
-    constant J2kIdwtSingleDecompositionParams &params [[buffer(5)]],
-    device J2kIdwtStatus *status [[buffer(6)]],
-    uint gid [[thread_position_in_grid]]
+kernel void j2k_idwt_irreversible97_horizontal_scale(
+    device float *out [[buffer(0)]],
+    constant J2kIdwtSingleDecompositionParams &params [[buffer(1)]],
+    uint2 gid [[thread_position_in_grid]]
 ) {
-    if (gid != 0u) {
-        return;
-    }
-
-    status->code = J2K_IDWT_STATUS_OK;
-    status->detail = 0u;
-    status->reserved0 = 0u;
-    status->reserved1 = 0u;
-
-    if (params.width == 0u || params.height == 0u) {
-        status->code = J2K_IDWT_STATUS_FAIL;
-        status->detail = 1u;
+    if (gid.x >= params.width || gid.y >= params.height) {
         return;
     }
 
-    const float NEG_ALPHA = CODEC_MATH_IDWT97_NEG_ALPHA;
-    const float NEG_BETA = CODEC_MATH_IDWT97_NEG_BETA;
-    const float NEG_GAMMA = CODEC_MATH_IDWT97_NEG_GAMMA;
-    const float NEG_DELTA = CODEC_MATH_IDWT97_NEG_DELTA;
     const float KAPPA = CODEC_MATH_DWT97_KAPPA;
     const float INV_KAPPA = CODEC_MATH_DWT97_INV_KAPPA;
-
-    const uint low_x_parity = params.x0 & 1u;
-    const uint low_y_parity = params.y0 & 1u;
-
-    for (uint local_y = 0u; local_y < params.height; ++local_y) {
-        const uint global_y = params.y0 + params.output_y + local_y;
-        const bool low_y = (global_y & 1u) == low_y_parity;
-        const uint full_band_y = low_y ? low_index(global_y, params.y0) : high_index(global_y, params.y0);
-
-        for (uint local_x = 0u; local_x < params.width; ++local_x) {
-            const uint global_x = params.x0 + params.output_x + local_x;
-            const bool low_x = (global_x & 1u) == low_x_parity;
-            const uint full_band_x = low_x ? low_index(global_x, params.x0) : high_index(global_x, params.x0);
-            const uint out_idx = local_y * params.width + local_x;
-
-            if (low_y && low_x) {
-                const uint band_x = full_band_x - params.ll_x;
-                const uint band_y = full_band_y - params.ll_y;
-                out[out_idx] = (band_x < params.ll_width && band_y < params.ll_height)
-                    ? ll[band_y * params.ll_width + band_x]
-                    : 0.0f;
-            } else if (low_y) {
-                const uint band_x = full_band_x - params.hl_x;
-                const uint band_y = full_band_y - params.hl_y;
-                out[out_idx] = (band_x < params.hl_width && band_y < params.hl_height)
-                    ? hl[band_y * params.hl_width + band_x]
-                    : 0.0f;
-            } else if (low_x) {
-                const uint band_x = full_band_x - params.lh_x;
-                const uint band_y = full_band_y - params.lh_y;
-                out[out_idx] = (band_x < params.lh_width && band_y < params.lh_height)
-                    ? lh[band_y * params.lh_width + band_x]
-                    : 0.0f;
-            } else {
-                const uint band_x = full_band_x - params.hh_x;
-                const uint band_y = full_band_y - params.hh_y;
-                out[out_idx] = (band_x < params.hh_width && band_y < params.hh_height)
-                    ? hh[band_y * params.hh_width + band_x]
-                    : 0.0f;
-            }
-        }
-    }
+    float sample = out[gid.y * params.width + gid.x];
 
     if (params.width == 1u) {
         if (((params.x0 + params.output_x) & 1u) != 0u) {
-            for (uint row = 0u; row < params.height; ++row) {
-                out[row * params.width] *= 0.5f;
-            }
+            sample *= 0.5f;
         }
     } else {
         const uint first_even_x = (params.x0 + params.output_x) & 1u;
-        const uint first_odd_x = 1u - first_even_x;
-        const float k0 = first_even_x == 0u ? KAPPA : INV_KAPPA;
-        const float k1 = first_even_x == 0u ? INV_KAPPA : KAPPA;
-
-        for (uint row = 0u; row < params.height; ++row) {
-            device float *row_ptr = out + row * params.width;
-
-            for (uint x = 0u; x + 1u < params.width; x += 2u) {
-                row_ptr[x] *= k0;
-                row_ptr[x + 1u] *= k1;
-            }
-            if ((params.width & 1u) != 0u) {
-                row_ptr[params.width - 1u] *= k0;
-            }
-
-            irreversible97_horizontal_step(row_ptr, params.width, first_even_x, NEG_DELTA);
-            irreversible97_horizontal_step(row_ptr, params.width, first_odd_x, NEG_GAMMA);
-            irreversible97_horizontal_step(row_ptr, params.width, first_even_x, NEG_BETA);
-            irreversible97_horizontal_step(row_ptr, params.width, first_odd_x, NEG_ALPHA);
-        }
+        sample *= (gid.x & 1u) == first_even_x ? KAPPA : INV_KAPPA;
     }
 
-    if (params.height == 1u) {
-        if (((params.y0 + params.output_y) & 1u) != 0u) {
-            for (uint col = 0u; col < params.width; ++col) {
-                out[col] *= 0.5f;
-            }
-        }
+    out[gid.y * params.width + gid.x] = sample;
+}
+
+kernel void j2k_idwt_irreversible97_vertical_scale(
+    device float *out [[buffer(0)]],
+    constant J2kIdwtSingleDecompositionParams &params [[buffer(1)]],
+    uint2 gid [[thread_position_in_grid]]
+) {
+    if (gid.x >= params.width || gid.y >= params.height) {
         return;
     }
 
-    const uint first_even_y = (params.y0 + params.output_y) & 1u;
-    const uint first_odd_y = 1u - first_even_y;
-    const float k0 = first_even_y == 0u ? KAPPA : INV_KAPPA;
-    const float k1 = first_even_y == 0u ? INV_KAPPA : KAPPA;
+    const float KAPPA = CODEC_MATH_DWT97_KAPPA;
+    const float INV_KAPPA = CODEC_MATH_DWT97_INV_KAPPA;
+    float sample = out[gid.y * params.width + gid.x];
 
-    for (uint row = 0u; row + 1u < params.height; row += 2u) {
-        for (uint col = 0u; col < params.width; ++col) {
-            out[row * params.width + col] *= k0;
-            out[(row + 1u) * params.width + col] *= k1;
-        }
-    }
-    if ((params.height & 1u) != 0u) {
-        const uint row = params.height - 1u;
-        for (uint col = 0u; col < params.width; ++col) {
-            out[row * params.width + col] *= k0;
+    if (params.height == 1u) {
+        if (((params.y0 + params.output_y) & 1u) != 0u) {
+            sample *= 0.5f;
         }
+    } else {
+        const uint first_even_y = (params.y0 + params.output_y) & 1u;
+        sample *= (gid.y & 1u) == first_even_y ? KAPPA : INV_KAPPA;
     }
 
-    for (uint row = first_even_y; row < params.height; row += 2u) {
-        const uint row_above = periodic_symmetric_extension_left_u32(row, 1u);
-        const uint row_below = periodic_symmetric_extension_right_u32(row, 1u, params.height);
-        for (uint col = 0u; col < params.width; ++col) {
-            const uint idx = row * params.width + col;
-            out[idx] = fma(
-                out[row_above * params.width + col] + out[row_below * params.width + col],
-                NEG_DELTA,
-                out[idx]
-            );
-        }
-    }
+    out[gid.y * params.width + gid.x] = sample;
+}
 
-    for (uint row = first_odd_y; row < params.height; row += 2u) {
-        const uint row_above = periodic_symmetric_extension_left_u32(row, 1u);
-        const uint row_below = periodic_symmetric_extension_right_u32(row, 1u, params.height);
-        for (uint col = 0u; col < params.width; ++col) {
-            const uint idx = row * params.width + col;
-            out[idx] = fma(
-                out[row_above * params.width + col] + out[row_below * params.width + col],
-                NEG_GAMMA,
-                out[idx]
-            );
-        }
+kernel void j2k_idwt_irreversible97_horizontal_step(
+    device float *out [[buffer(0)]],
+    constant J2kIdwtSingleDecompositionParams &params [[buffer(1)]],
+    constant J2kIdwt97StepParams &step [[buffer(2)]],
+    uint2 gid [[thread_position_in_grid]]
+) {
+    if (gid.x >= params.width || gid.y >= params.height || params.width <= 1u
+        || (gid.x & 1u) != step.parity) {
+        return;
     }
 
-    for (uint row = first_even_y; row < params.height; row += 2u) {
-        const uint row_above = periodic_symmetric_extension_left_u32(row, 1u);
-        const uint row_below = periodic_symmetric_extension_right_u32(row, 1u, params.height);
-        for (uint col = 0u; col < params.width; ++col) {
-            const uint idx = row * params.width + col;
-            out[idx] = fma(
-                out[row_above * params.width + col] + out[row_below * params.width + col],
-                NEG_BETA,
-                out[idx]
-            );
-        }
-    }
+    const uint left = periodic_symmetric_extension_left_u32(gid.x, 1u);
+    const uint right = periodic_symmetric_extension_right_u32(gid.x, 1u, params.width);
+    const uint idx = gid.y * params.width + gid.x;
+    out[idx] = fma(out[gid.y * params.width + left] + out[gid.y * params.width + right],
+                   step.coefficient,
+                   out[idx]);
+}
 
-    for (uint row = first_odd_y; row < params.height; row += 2u) {
-        const uint row_above = periodic_symmetric_extension_left_u32(row, 1u);
-        const uint row_below = periodic_symmetric_extension_right_u32(row, 1u, params.height);
-        for (uint col = 0u; col < params.width; ++col) {
-            const uint idx = row * params.width + col;
-            out[idx] = fma(
-                out[row_above * params.width + col] + out[row_below * params.width + col],
-                NEG_ALPHA,
-                out[idx]
-            );
-        }
+kernel void j2k_idwt_irreversible97_vertical_step(
+    device float *out [[buffer(0)]],
+    constant J2kIdwtSingleDecompositionParams &params [[buffer(1)]],
+    constant J2kIdwt97StepParams &step [[buffer(2)]],
+    uint2 gid [[thread_position_in_grid]]
+) {
+    if (gid.x >= params.width || gid.y >= params.height || params.height <= 1u
+        || (gid.y & 1u) != step.parity) {
+        return;
     }
+
+    const uint above = periodic_symmetric_extension_left_u32(gid.y, 1u);
+    const uint below = periodic_symmetric_extension_right_u32(gid.y, 1u, params.height);
+    const uint idx = gid.y * params.width + gid.x;
+    out[idx] = fma(out[above * params.width + gid.x] + out[below * params.width + gid.x],
+                   step.coefficient,
+                   out[idx]);
 }
diff --git a/crates/j2k-metal/src/idwt.rs b/crates/j2k-metal/src/idwt.rs
index da24821d..43d69193 100644
--- a/crates/j2k-metal/src/idwt.rs
+++ b/crates/j2k-metal/src/idwt.rs
@@ -217,6 +217,212 @@ mod tests {
         );
     }
 
+    #[cfg(target_os = "macos")]
+    #[test]
+    fn staged_irreversible_idwt_matches_native_odd_geometry() {
+        const EXPECTED_BITS: [u32; 15] = [
+            3_243_516_307,
+            1_088_535_889,
+            1_086_781_832,
+            3_237_068_116,
+            1_072_899_983,
+            1_090_811_482,
+            3_205_449_560,
+            3_240_528_043,
+            1_057_965_919,
+            1_095_801_595,
+            3_240_069_383,
+            1_090_703_406,
+            1_072_530_023,
+            3_238_711_343,
+            1_091_758_981,
+        ];
+        if !should_run_metal_runtime() {
+            return;
+        }
+        let rect = j2k_native::J2kRect {
+            x0: 0,
+            y0: 0,
+            x1: 5,
+            y1: 3,
+        };
+        let ll = [0.5, -1.25, 2.0, 3.5, -4.25, 5.75];
+        let hl = [6.5, -7.0, 8.25, -9.5];
+        let lh = [10.0, -11.5, 12.75];
+        let hh = [-13.0, 14.25];
+        let job = j2k_native::J2kSingleDecompositionIdwtJob {
+            rect,
+            transform: j2k_native::J2kWaveletTransform::Irreversible97,
+            ll: j2k_native::J2kIdwtBand {
+                rect: j2k_native::J2kRect {
+                    x0: 0,
+                    y0: 0,
+                    x1: 3,
+                    y1: 2,
+                },
+                coefficients: &ll,
+            },
+            hl: j2k_native::J2kIdwtBand {
+                rect: j2k_native::J2kRect {
+                    x0: 0,
+                    y0: 0,
+                    x1: 2,
+                    y1: 2,
+                },
+                coefficients: &hl,
+            },
+            lh: j2k_native::J2kIdwtBand {
+                rect: j2k_native::J2kRect {
+                    x0: 0,
+                    y0: 0,
+                    x1: 3,
+                    y1: 1,
+                },
+                coefficients: &lh,
+            },
+            hh: j2k_native::J2kIdwtBand {
+                rect: j2k_native::J2kRect {
+                    x0: 0,
+                    y0: 0,
+                    x1: 2,
+                    y1: 1,
+                },
+                coefficients: &hh,
+            },
+        };
+        let mut actual = vec![0.0; EXPECTED_BITS.len()];
+
+        crate::compute::decode_irreversible97_staged_single_decomposition_idwt(job, &mut actual)
+            .expect("staged irreversible Metal IDWT");
+
+        for (index, (actual, expected_bits)) in actual.iter().zip(EXPECTED_BITS).enumerate() {
+            let expected = f32::from_bits(expected_bits);
+            assert!(
+                (actual - expected).abs() <= 2.0e-5,
+                "sample {index}: expected {expected}, got {actual}"
+            );
+        }
+    }
+
+    #[cfg(target_os = "macos")]
+    #[test]
+    fn staged_irreversible_idwt_preserves_degenerate_origin_scaling() {
+        if !should_run_metal_runtime() {
+            return;
+        }
+
+        let coefficient = [8.0];
+        for (x0, y0, expected) in [(0, 0, 8.0_f32), (1, 0, 4.0), (0, 1, 4.0), (1, 1, 2.0)] {
+            let rect = j2k_native::J2kRect {
+                x0,
+                y0,
+                x1: x0 + 1,
+                y1: y0 + 1,
+            };
+            let band = j2k_native::J2kIdwtBand {
+                rect: j2k_native::J2kRect {
+                    x0: 0,
+                    y0: 0,
+                    x1: 1,
+                    y1: 1,
+                },
+                coefficients: &coefficient,
+            };
+            let job = j2k_native::J2kSingleDecompositionIdwtJob {
+                rect,
+                transform: j2k_native::J2kWaveletTransform::Irreversible97,
+                ll: band,
+                hl: band,
+                lh: band,
+                hh: band,
+            };
+            let mut actual = [0.0];
+
+            crate::compute::decode_irreversible97_staged_single_decomposition_idwt(
+                job,
+                &mut actual,
+            )
+            .expect("degenerate staged irreversible Metal IDWT");
+
+            assert_eq!(
+                actual[0].to_bits(),
+                expected.to_bits(),
+                "origin ({x0}, {y0})"
+            );
+        }
+    }
+
+    #[cfg(target_os = "macos")]
+    #[test]
+    #[ignore = "performance guard harness; run explicitly with --ignored --nocapture"]
+    fn metal_irreversible_idwt_perf_guard() {
+        const WIDTH: u32 = 1023;
+        const HEIGHT: u32 = 767;
+        const ITERS: usize = 11;
+        if !should_run_metal_runtime() {
+            return;
+        }
+        let low_width = WIDTH.div_ceil(2);
+        let low_height = HEIGHT.div_ceil(2);
+        let high_width = WIDTH / 2;
+        let high_height = HEIGHT / 2;
+        let make_band = |width: u32, height: u32, seed: u32| {
+            (0..width * height)
+                .map(|index| {
+                    let value = index.wrapping_mul(37).wrapping_add(seed * 101) % 4093;
+                    (f32::from(u16::try_from(value).expect("pattern value fits u16")) - 2046.0)
+                        * 0.03125
+                })
+                .collect::<Vec<_>>()
+        };
+        let ll = make_band(low_width, low_height, 1);
+        let hl = make_band(high_width, low_height, 2);
+        let lh = make_band(low_width, high_height, 3);
+        let hh = make_band(high_width, high_height, 4);
+        let rect = j2k_native::J2kRect {
+            x0: 0,
+            y0: 0,
+            x1: WIDTH,
+            y1: HEIGHT,
+        };
+        let band = |x1, y1, coefficients| j2k_native::J2kIdwtBand {
+            rect: j2k_native::J2kRect {
+                x0: 0,
+                y0: 0,
+                x1,
+                y1,
+            },
+            coefficients,
+        };
+        let job = j2k_native::J2kSingleDecompositionIdwtJob {
+            rect,
+            transform: j2k_native::J2kWaveletTransform::Irreversible97,
+            ll: band(low_width, low_height, &ll),
+            hl: band(high_width, low_height, &hl),
+            lh: band(low_width, high_height, &lh),
+            hh: band(high_width, high_height, &hh),
+        };
+        let mut output = vec![0.0; WIDTH as usize * HEIGHT as usize];
+        crate::compute::decode_irreversible97_staged_single_decomposition_idwt(job, &mut output)
+            .expect("warm irreversible Metal IDWT");
+        let mut samples = Vec::with_capacity(ITERS);
+        for _ in 0..ITERS {
+            let started = std::time::Instant::now();
+            crate::compute::decode_irreversible97_staged_single_decomposition_idwt(
+                job,
+                &mut output,
+            )
+            .expect("measured irreversible Metal IDWT");
+            samples.push(started.elapsed());
+        }
+        samples.sort_unstable();
+        let median = samples[ITERS / 2];
+        println!(
+            "j2k_metal_idwt97_perf mode=staged size={WIDTH}x{HEIGHT} median_ms={:.3}",
+            median.as_secs_f64() * 1_000.0
+        );
+    }
+
     struct CpuOnlyCodeBlockDecoder;
 
     impl HtCodeBlockDecoder for CpuOnlyCodeBlockDecoder {}
diff --git a/crates/j2k-metal/tests/metal_decode_benchmark.rs b/crates/j2k-metal/tests/metal_decode_benchmark.rs
index 92b7f0fe..56402cc7 100644
--- a/crates/j2k-metal/tests/metal_decode_benchmark.rs
+++ b/crates/j2k-metal/tests/metal_decode_benchmark.rs
@@ -178,6 +178,14 @@ fn generated_decode_cases() -> Vec<DecodeBenchCase> {
             container: "raw-codestream".to_string(),
             source: "generated".to_string(),
         });
+        cases.push(DecodeBenchCase {
+            id: format!("generated_classic_irreversible97_gray8_{dim}"),
+            bytes: encode_classic_irreversible(&gray, dim, dim, 1),
+            fmt: PixelFormat::Gray8,
+            codec: "j2k-irreversible97".to_string(),
+            container: "raw-codestream".to_string(),
+            source: "generated".to_string(),
+        });
         cases.push(DecodeBenchCase {
             id: format!("generated_htj2k_gray8_{dim}"),
             bytes: encode_ht(&gray, dim, dim, 1),
@@ -204,6 +212,15 @@ fn encode_classic(pixels: &[u8], width: u32, height: u32, components: u16) -> Ve
     encode(pixels, width, height, components, 8, false, &options).expect("encode classic fixture")
 }
 
+fn encode_classic_irreversible(pixels: &[u8], width: u32, height: u32, components: u16) -> Vec<u8> {
+    let options = EncodeOptions {
+        reversible: false,
+        ..encode_options()
+    };
+    encode(pixels, width, height, components, 8, false, &options)
+        .expect("encode irreversible classic fixture")
+}
+
 fn encode_ht(pixels: &[u8], width: u32, height: u32, components: u16) -> Vec<u8> {
     let options = encode_options();
     encode_htj2k(pixels, width, height, components, 8, false, &options)

From c35b7d43cfe1f6da6a1b6abc725e472bd6c71622 Mon Sep 17 00:00:00 2001
From: GF <greg.furletti@gmail.com>
Date: Mon, 13 Jul 2026 19:53:15 -0400
Subject: [PATCH 4/7] test(metal): add irreversible IDWT GPU capture harness

---
 crates/j2k-metal/src/idwt.rs | 227 +++++++++++++++++++++++++++--------
 1 file changed, 179 insertions(+), 48 deletions(-)

diff --git a/crates/j2k-metal/src/idwt.rs b/crates/j2k-metal/src/idwt.rs
index 43d69193..35c40949 100644
--- a/crates/j2k-metal/src/idwt.rs
+++ b/crates/j2k-metal/src/idwt.rs
@@ -352,64 +352,188 @@ mod tests {
         }
     }
 
+    #[cfg(target_os = "macos")]
+    const IRREVERSIBLE_IDWT_PERF_WIDTH: u32 = 1023;
+    #[cfg(target_os = "macos")]
+    const IRREVERSIBLE_IDWT_PERF_HEIGHT: u32 = 767;
+
+    #[cfg(target_os = "macos")]
+    struct IrreversibleIdwtPerfFixture {
+        rect: j2k_native::J2kRect,
+        ll: Vec<f32>,
+        hl: Vec<f32>,
+        lh: Vec<f32>,
+        hh: Vec<f32>,
+    }
+
+    #[cfg(target_os = "macos")]
+    impl IrreversibleIdwtPerfFixture {
+        fn new() -> Self {
+            let low_width = IRREVERSIBLE_IDWT_PERF_WIDTH.div_ceil(2);
+            let low_height = IRREVERSIBLE_IDWT_PERF_HEIGHT.div_ceil(2);
+            let high_width = IRREVERSIBLE_IDWT_PERF_WIDTH / 2;
+            let high_height = IRREVERSIBLE_IDWT_PERF_HEIGHT / 2;
+            let make_band = |width: u32, height: u32, seed: u32| {
+                (0..width * height)
+                    .map(|index| {
+                        let value = index.wrapping_mul(37).wrapping_add(seed * 101) % 4093;
+                        (f32::from(u16::try_from(value).expect("pattern value fits u16")) - 2046.0)
+                            * 0.03125
+                    })
+                    .collect::<Vec<_>>()
+            };
+            Self {
+                rect: j2k_native::J2kRect {
+                    x0: 0,
+                    y0: 0,
+                    x1: IRREVERSIBLE_IDWT_PERF_WIDTH,
+                    y1: IRREVERSIBLE_IDWT_PERF_HEIGHT,
+                },
+                ll: make_band(low_width, low_height, 1),
+                hl: make_band(high_width, low_height, 2),
+                lh: make_band(low_width, high_height, 3),
+                hh: make_band(high_width, high_height, 4),
+            }
+        }
+
+        fn job(&self) -> j2k_native::J2kSingleDecompositionIdwtJob<'_> {
+            let low_width = IRREVERSIBLE_IDWT_PERF_WIDTH.div_ceil(2);
+            let low_height = IRREVERSIBLE_IDWT_PERF_HEIGHT.div_ceil(2);
+            let high_width = IRREVERSIBLE_IDWT_PERF_WIDTH / 2;
+            let high_height = IRREVERSIBLE_IDWT_PERF_HEIGHT / 2;
+            let band = |x1, y1, coefficients| j2k_native::J2kIdwtBand {
+                rect: j2k_native::J2kRect {
+                    x0: 0,
+                    y0: 0,
+                    x1,
+                    y1,
+                },
+                coefficients,
+            };
+            j2k_native::J2kSingleDecompositionIdwtJob {
+                rect: self.rect,
+                transform: j2k_native::J2kWaveletTransform::Irreversible97,
+                ll: band(low_width, low_height, &self.ll),
+                hl: band(high_width, low_height, &self.hl),
+                lh: band(low_width, high_height, &self.lh),
+                hh: band(high_width, high_height, &self.hh),
+            }
+        }
+
+        fn output(&self) -> Vec<f32> {
+            vec![
+                0.0;
+                IRREVERSIBLE_IDWT_PERF_WIDTH as usize * IRREVERSIBLE_IDWT_PERF_HEIGHT as usize
+            ]
+        }
+    }
+
+    #[cfg(target_os = "macos")]
+    struct MetalCaptureGuard<'a> {
+        manager: &'a metal::CaptureManagerRef,
+    }
+
+    #[cfg(target_os = "macos")]
+    impl Drop for MetalCaptureGuard<'_> {
+        fn drop(&mut self) {
+            if self.manager.is_capturing() {
+                self.manager.stop_capture();
+            }
+        }
+    }
+
+    #[cfg(target_os = "macos")]
+    #[test]
+    #[ignore = "GPU capture harness; run explicitly with --ignored --nocapture"]
+    fn metal_irreversible_idwt_gpu_capture() {
+        use metal::{CaptureDescriptor, CaptureManager, MTLCaptureDestination};
+
+        if !should_run_metal_runtime() {
+            return;
+        }
+        assert_eq!(
+            std::env::var("MTL_CAPTURE_ENABLED").as_deref(),
+            Ok("1"),
+            "set MTL_CAPTURE_ENABLED=1 to enable the Metal capture API"
+        );
+        let trace_path = std::path::PathBuf::from(
+            std::env::var_os("J2K_METAL_CAPTURE_PATH")
+                .expect("set J2K_METAL_CAPTURE_PATH to an absolute .gputrace output path"),
+        );
+        assert!(
+            trace_path.is_absolute(),
+            "J2K_METAL_CAPTURE_PATH must be absolute"
+        );
+        assert_eq!(
+            trace_path.extension().and_then(std::ffi::OsStr::to_str),
+            Some("gputrace"),
+            "J2K_METAL_CAPTURE_PATH must end in .gputrace"
+        );
+        assert!(
+            !trace_path.exists(),
+            "refusing to overwrite existing GPU trace {}",
+            trace_path.display()
+        );
+
+        let fixture = IrreversibleIdwtPerfFixture::new();
+        let mut output = fixture.output();
+        let device = j2k_metal_support::system_default_device().expect("Metal capture device");
+        crate::compute::with_isolated_runtime_for_device_for_test(&device, || {
+            crate::compute::decode_irreversible97_staged_single_decomposition_idwt(
+                fixture.job(),
+                &mut output,
+            )?;
+
+            let manager = CaptureManager::shared();
+            assert!(
+                manager.supports_destination(MTLCaptureDestination::GpuTraceDocument),
+                "Metal GPU trace documents are unavailable on this host"
+            );
+            let descriptor = CaptureDescriptor::new();
+            descriptor.set_capture_device(&device);
+            descriptor.set_destination(MTLCaptureDestination::GpuTraceDocument);
+            descriptor.set_output_url(&trace_path);
+            manager
+                .start_capture(&descriptor)
+                .map_err(|message| crate::Error::MetalRuntime { message })?;
+            let capture = MetalCaptureGuard { manager };
+            let result = crate::compute::decode_irreversible97_staged_single_decomposition_idwt(
+                fixture.job(),
+                &mut output,
+            );
+            drop(capture);
+            result
+        })
+        .expect("captured irreversible Metal IDWT");
+
+        assert!(
+            trace_path.is_dir(),
+            "Metal capture did not create GPU trace package {}",
+            trace_path.display()
+        );
+        println!("j2k_metal_idwt97_capture path={}", trace_path.display());
+    }
+
     #[cfg(target_os = "macos")]
     #[test]
     #[ignore = "performance guard harness; run explicitly with --ignored --nocapture"]
     fn metal_irreversible_idwt_perf_guard() {
-        const WIDTH: u32 = 1023;
-        const HEIGHT: u32 = 767;
         const ITERS: usize = 11;
         if !should_run_metal_runtime() {
             return;
         }
-        let low_width = WIDTH.div_ceil(2);
-        let low_height = HEIGHT.div_ceil(2);
-        let high_width = WIDTH / 2;
-        let high_height = HEIGHT / 2;
-        let make_band = |width: u32, height: u32, seed: u32| {
-            (0..width * height)
-                .map(|index| {
-                    let value = index.wrapping_mul(37).wrapping_add(seed * 101) % 4093;
-                    (f32::from(u16::try_from(value).expect("pattern value fits u16")) - 2046.0)
-                        * 0.03125
-                })
-                .collect::<Vec<_>>()
-        };
-        let ll = make_band(low_width, low_height, 1);
-        let hl = make_band(high_width, low_height, 2);
-        let lh = make_band(low_width, high_height, 3);
-        let hh = make_band(high_width, high_height, 4);
-        let rect = j2k_native::J2kRect {
-            x0: 0,
-            y0: 0,
-            x1: WIDTH,
-            y1: HEIGHT,
-        };
-        let band = |x1, y1, coefficients| j2k_native::J2kIdwtBand {
-            rect: j2k_native::J2kRect {
-                x0: 0,
-                y0: 0,
-                x1,
-                y1,
-            },
-            coefficients,
-        };
-        let job = j2k_native::J2kSingleDecompositionIdwtJob {
-            rect,
-            transform: j2k_native::J2kWaveletTransform::Irreversible97,
-            ll: band(low_width, low_height, &ll),
-            hl: band(high_width, low_height, &hl),
-            lh: band(low_width, high_height, &lh),
-            hh: band(high_width, high_height, &hh),
-        };
-        let mut output = vec![0.0; WIDTH as usize * HEIGHT as usize];
-        crate::compute::decode_irreversible97_staged_single_decomposition_idwt(job, &mut output)
-            .expect("warm irreversible Metal IDWT");
+        let fixture = IrreversibleIdwtPerfFixture::new();
+        let mut output = fixture.output();
+        crate::compute::decode_irreversible97_staged_single_decomposition_idwt(
+            fixture.job(),
+            &mut output,
+        )
+        .expect("warm irreversible Metal IDWT");
         let mut samples = Vec::with_capacity(ITERS);
         for _ in 0..ITERS {
             let started = std::time::Instant::now();
             crate::compute::decode_irreversible97_staged_single_decomposition_idwt(
-                job,
+                fixture.job(),
                 &mut output,
             )
             .expect("measured irreversible Metal IDWT");
@@ -417,9 +541,16 @@ mod tests {
         }
         samples.sort_unstable();
         let median = samples[ITERS / 2];
+        let p25 = samples[ITERS / 4];
+        let p75 = samples[ITERS * 3 / 4];
         println!(
-            "j2k_metal_idwt97_perf mode=staged size={WIDTH}x{HEIGHT} median_ms={:.3}",
-            median.as_secs_f64() * 1_000.0
+            "j2k_metal_idwt97_perf mode=staged size={}x{} iterations={ITERS} median_ms={:.3} p25_ms={:.3} p75_ms={:.3} iqr_ms={:.3}",
+            IRREVERSIBLE_IDWT_PERF_WIDTH,
+            IRREVERSIBLE_IDWT_PERF_HEIGHT,
+            median.as_secs_f64() * 1_000.0,
+            p25.as_secs_f64() * 1_000.0,
+            p75.as_secs_f64() * 1_000.0,
+            (p75 - p25).as_secs_f64() * 1_000.0
         );
     }
 

From d2ba40c527e858ebfb07516b55854e6dd59118fe Mon Sep 17 00:00:00 2001
From: GF <greg.furletti@gmail.com>
Date: Mon, 13 Jul 2026 20:26:31 -0400
Subject: [PATCH 5/7] test(metal): harden IDWT profiling harness

---
 crates/j2k-metal/src/idwt.rs | 12 ++++++++----
 1 file changed, 8 insertions(+), 4 deletions(-)

diff --git a/crates/j2k-metal/src/idwt.rs b/crates/j2k-metal/src/idwt.rs
index 35c40949..87413cde 100644
--- a/crates/j2k-metal/src/idwt.rs
+++ b/crates/j2k-metal/src/idwt.rs
@@ -420,7 +420,7 @@ mod tests {
             }
         }
 
-        fn output(&self) -> Vec<f32> {
+        fn output() -> Vec<f32> {
             vec![
                 0.0;
                 IRREVERSIBLE_IDWT_PERF_WIDTH as usize * IRREVERSIBLE_IDWT_PERF_HEIGHT as usize
@@ -451,6 +451,10 @@ mod tests {
         if !should_run_metal_runtime() {
             return;
         }
+        if std::env::var_os("CARGO_LLVM_COV_TARGET_DIR").is_some() {
+            println!("skipping GPU capture while the Metal coverage lane is instrumented");
+            return;
+        }
         assert_eq!(
             std::env::var("MTL_CAPTURE_ENABLED").as_deref(),
             Ok("1"),
@@ -476,7 +480,7 @@ mod tests {
         );
 
         let fixture = IrreversibleIdwtPerfFixture::new();
-        let mut output = fixture.output();
+        let mut output = IrreversibleIdwtPerfFixture::output();
         let device = j2k_metal_support::system_default_device().expect("Metal capture device");
         crate::compute::with_isolated_runtime_for_device_for_test(&device, || {
             crate::compute::decode_irreversible97_staged_single_decomposition_idwt(
@@ -523,7 +527,7 @@ mod tests {
             return;
         }
         let fixture = IrreversibleIdwtPerfFixture::new();
-        let mut output = fixture.output();
+        let mut output = IrreversibleIdwtPerfFixture::output();
         crate::compute::decode_irreversible97_staged_single_decomposition_idwt(
             fixture.job(),
             &mut output,
@@ -550,7 +554,7 @@ mod tests {
             median.as_secs_f64() * 1_000.0,
             p25.as_secs_f64() * 1_000.0,
             p75.as_secs_f64() * 1_000.0,
-            (p75 - p25).as_secs_f64() * 1_000.0
+            p75.saturating_sub(p25).as_secs_f64() * 1_000.0
         );
     }
 

From 5a73a38be0628cb3b6d8f7c18e5a2b0eb01a64c5 Mon Sep 17 00:00:00 2001
From: GF <greg.furletti@gmail.com>
Date: Mon, 13 Jul 2026 20:26:37 -0400
Subject: [PATCH 6/7] docs(metal): record rejected IDWT fusion experiment

---
 docs/benchmark-evidence.md | 90 ++++++++++++++++++++++++++++++++++++++
 1 file changed, 90 insertions(+)

diff --git a/docs/benchmark-evidence.md b/docs/benchmark-evidence.md
index ec855686..d5a2639e 100644
--- a/docs/benchmark-evidence.md
+++ b/docs/benchmark-evidence.md
@@ -240,6 +240,96 @@ cargo test -p j2k-cuda-runtime --lib \
 The validation host also needed its linker search path to include the local
 LLVM/libffi runtime used by cuda-oxide.
 
+## Metal Irreversible 9/7 Heatmap Experiment - 2026-07-13
+
+This local experiment used Apple's GPU performance heat maps to test whether
+the staged irreversible 9/7 IDWT should be replaced by full-axis
+threadgroup-memory kernels. The prototype was rejected and removed because it
+failed the targeted performance gate. The capture and dispersion harness at
+`c35b7d43` remains as permanent diagnostic infrastructure.
+
+Environment:
+
+- Machine: MacBook Pro `Mac16,8`, Apple M4 Pro, 48 GB
+- OS: macOS 26.5 build `25F71`
+- Xcode: 26.6 build `17F113`
+- Metal toolchain: `com.apple.dt.toolchain.Metal.32023.883`, component build
+  `17F109`
+- Baseline and retained source: `c35b7d43`
+- Profile: Cargo `release-bench`, isolated baseline/candidate target
+  directories, with no concurrent Cargo, CUDA, Metal, Criterion, or heavy
+  build workload
+
+The captured 1023 x 767 staged transform exposed the runtime-generated Metal
+source in Xcode. Profile GPU Trace reported 214.21 us total GPU time, one
+command buffer, one compute encoder, 11 dispatches, and 12.01 MiB of buffers.
+The four horizontal lifting dispatches accounted for 37.02% of GPU time, the
+four vertical lifting dispatches 36.43%, vertical scale 9.12%, horizontal
+scale 8.87%, and interleave 8.55%. The scale and lifting passes therefore
+accounted for 91.44% of captured GPU time.
+
+The trace reported 73.87% Occupancy Manager target, an 87.12% Compute Shader
+Launch limiter, and a 32.09% Last Level Cache limiter with 31.94% utilization.
+The staged pipelines used 12-14 allocated registers and spilled zero bytes.
+Those counters supported prototyping dispatch fusion; they did not establish
+that fusion would be faster.
+
+The prototype used one threadgroup per row or column, dynamic threadgroup
+memory, and separate horizontal and vertical kernels. Forced staged/fused
+tests matched within the existing irreversible tolerance across degenerate,
+odd/even, origin-parity, and 1023 x 767 cases. A pipeline-relative sweep then
+tested 1, 2, 4, 8, 16, and 32 SIMD groups per threadgroup:
+
+| SIMD groups | Threads on M4 Pro | Exploratory median | IQR |
+| ---: | ---: | ---: | ---: |
+| 1 | 32 | 1.658 ms | 0.227 ms |
+| 2 | 64 | 1.563 ms | 0.360 ms |
+| 4 | 128 | 1.437 ms | 0.298 ms |
+| 8 | 256 | 1.667 ms | 0.256 ms |
+| 16 | 512 | 1.686 ms | 0.207 ms |
+| 32 | 1024 | 1.745 ms | 0.701 ms |
+
+Because the 128-thread exploratory row appeared promising, it was remeasured
+in five alternating baseline/candidate processes with eleven warmed samples
+per process. Each cell reports process median / IQR:
+
+| Process | Staged baseline | 128-thread fused prototype |
+| ---: | ---: | ---: |
+| 1 | 1.422 / 0.381 ms | 1.641 / 0.300 ms |
+| 2 | 1.467 / 0.262 ms | 1.649 / 0.557 ms |
+| 3 | 1.629 / 0.261 ms | 1.677 / 0.288 ms |
+| 4 | 1.383 / 0.354 ms | 1.459 / 0.378 ms |
+| 5 | 1.324 / 0.410 ms | 1.646 / 0.373 ms |
+
+The median of process medians regressed from 1.422 ms to 1.646 ms (+15.75%).
+The initial maximum-thread configuration also regressed, from 1.550 ms to
+1.936 ms (+24.90%). Both miss the required 5% improvement by a wide margin.
+The kernel prototype and its temporary routing/tests were therefore removed.
+The end-to-end 512/1024 decode rows were not run because the targeted gate had
+already required rejection; no routing or performance claim changed.
+
+Changed-line coverage could not be reported. The canonical command
+`cargo xtask coverage metal --base HEAD^` failed before compiling coverage
+tests because the repository's Metal lane supplied `cargo-llvm-cov 0.8.7`
+with mutually incompatible `--no-report` and `--no-clean` options. The manual
+capture body is also intentionally skipped when coverage instrumentation is
+active so a coverage run cannot create a GPU trace. Package tests and strict
+Clippy passed, but this experiment therefore has no changed-line percentage.
+
+The permanent guard can be rerun with:
+
+```bash
+J2K_REQUIRE_METAL_RUNTIME=1 CARGO_TARGET_DIR=<isolated-target> \
+  cargo test -p j2k-metal metal_irreversible_idwt_perf_guard \
+  --profile release-bench --locked -- \
+  --ignored --nocapture --test-threads=1
+```
+
+The ignored GPU-capture test additionally requires `MTL_CAPTURE_ENABLED=1`
+and an absolute, nonexistent `.gputrace` path in
+`J2K_METAL_CAPTURE_PATH`. Captured timings are diagnostic only and are not
+used for acceptance.
+
 ## Metal Status
 
 Metal acceleration is selective. Public claims should say Metal-accelerated

From 21467fd8d8d612f8ed3ba3b9ec42181865eeb7c9 Mon Sep 17 00:00:00 2001
From: GF <greg.furletti@gmail.com>
Date: Tue, 21 Jul 2026 16:39:49 -0400
Subject: [PATCH 7/7] chore: remove superseded Metal capture evidence

---
 crates/j2k-metal/src/idwt.rs | 90 ------------------------------------
 docs/benchmark-evidence.md   | 90 ------------------------------------
 2 files changed, 180 deletions(-)

diff --git a/crates/j2k-metal/src/idwt.rs b/crates/j2k-metal/src/idwt.rs
index 87413cde..406d2d10 100644
--- a/crates/j2k-metal/src/idwt.rs
+++ b/crates/j2k-metal/src/idwt.rs
@@ -428,96 +428,6 @@ mod tests {
         }
     }
 
-    #[cfg(target_os = "macos")]
-    struct MetalCaptureGuard<'a> {
-        manager: &'a metal::CaptureManagerRef,
-    }
-
-    #[cfg(target_os = "macos")]
-    impl Drop for MetalCaptureGuard<'_> {
-        fn drop(&mut self) {
-            if self.manager.is_capturing() {
-                self.manager.stop_capture();
-            }
-        }
-    }
-
-    #[cfg(target_os = "macos")]
-    #[test]
-    #[ignore = "GPU capture harness; run explicitly with --ignored --nocapture"]
-    fn metal_irreversible_idwt_gpu_capture() {
-        use metal::{CaptureDescriptor, CaptureManager, MTLCaptureDestination};
-
-        if !should_run_metal_runtime() {
-            return;
-        }
-        if std::env::var_os("CARGO_LLVM_COV_TARGET_DIR").is_some() {
-            println!("skipping GPU capture while the Metal coverage lane is instrumented");
-            return;
-        }
-        assert_eq!(
-            std::env::var("MTL_CAPTURE_ENABLED").as_deref(),
-            Ok("1"),
-            "set MTL_CAPTURE_ENABLED=1 to enable the Metal capture API"
-        );
-        let trace_path = std::path::PathBuf::from(
-            std::env::var_os("J2K_METAL_CAPTURE_PATH")
-                .expect("set J2K_METAL_CAPTURE_PATH to an absolute .gputrace output path"),
-        );
-        assert!(
-            trace_path.is_absolute(),
-            "J2K_METAL_CAPTURE_PATH must be absolute"
-        );
-        assert_eq!(
-            trace_path.extension().and_then(std::ffi::OsStr::to_str),
-            Some("gputrace"),
-            "J2K_METAL_CAPTURE_PATH must end in .gputrace"
-        );
-        assert!(
-            !trace_path.exists(),
-            "refusing to overwrite existing GPU trace {}",
-            trace_path.display()
-        );
-
-        let fixture = IrreversibleIdwtPerfFixture::new();
-        let mut output = IrreversibleIdwtPerfFixture::output();
-        let device = j2k_metal_support::system_default_device().expect("Metal capture device");
-        crate::compute::with_isolated_runtime_for_device_for_test(&device, || {
-            crate::compute::decode_irreversible97_staged_single_decomposition_idwt(
-                fixture.job(),
-                &mut output,
-            )?;
-
-            let manager = CaptureManager::shared();
-            assert!(
-                manager.supports_destination(MTLCaptureDestination::GpuTraceDocument),
-                "Metal GPU trace documents are unavailable on this host"
-            );
-            let descriptor = CaptureDescriptor::new();
-            descriptor.set_capture_device(&device);
-            descriptor.set_destination(MTLCaptureDestination::GpuTraceDocument);
-            descriptor.set_output_url(&trace_path);
-            manager
-                .start_capture(&descriptor)
-                .map_err(|message| crate::Error::MetalRuntime { message })?;
-            let capture = MetalCaptureGuard { manager };
-            let result = crate::compute::decode_irreversible97_staged_single_decomposition_idwt(
-                fixture.job(),
-                &mut output,
-            );
-            drop(capture);
-            result
-        })
-        .expect("captured irreversible Metal IDWT");
-
-        assert!(
-            trace_path.is_dir(),
-            "Metal capture did not create GPU trace package {}",
-            trace_path.display()
-        );
-        println!("j2k_metal_idwt97_capture path={}", trace_path.display());
-    }
-
     #[cfg(target_os = "macos")]
     #[test]
     #[ignore = "performance guard harness; run explicitly with --ignored --nocapture"]
diff --git a/docs/benchmark-evidence.md b/docs/benchmark-evidence.md
index d5a2639e..ec855686 100644
--- a/docs/benchmark-evidence.md
+++ b/docs/benchmark-evidence.md
@@ -240,96 +240,6 @@ cargo test -p j2k-cuda-runtime --lib \
 The validation host also needed its linker search path to include the local
 LLVM/libffi runtime used by cuda-oxide.
 
-## Metal Irreversible 9/7 Heatmap Experiment - 2026-07-13
-
-This local experiment used Apple's GPU performance heat maps to test whether
-the staged irreversible 9/7 IDWT should be replaced by full-axis
-threadgroup-memory kernels. The prototype was rejected and removed because it
-failed the targeted performance gate. The capture and dispersion harness at
-`c35b7d43` remains as permanent diagnostic infrastructure.
-
-Environment:
-
-- Machine: MacBook Pro `Mac16,8`, Apple M4 Pro, 48 GB
-- OS: macOS 26.5 build `25F71`
-- Xcode: 26.6 build `17F113`
-- Metal toolchain: `com.apple.dt.toolchain.Metal.32023.883`, component build
-  `17F109`
-- Baseline and retained source: `c35b7d43`
-- Profile: Cargo `release-bench`, isolated baseline/candidate target
-  directories, with no concurrent Cargo, CUDA, Metal, Criterion, or heavy
-  build workload
-
-The captured 1023 x 767 staged transform exposed the runtime-generated Metal
-source in Xcode. Profile GPU Trace reported 214.21 us total GPU time, one
-command buffer, one compute encoder, 11 dispatches, and 12.01 MiB of buffers.
-The four horizontal lifting dispatches accounted for 37.02% of GPU time, the
-four vertical lifting dispatches 36.43%, vertical scale 9.12%, horizontal
-scale 8.87%, and interleave 8.55%. The scale and lifting passes therefore
-accounted for 91.44% of captured GPU time.
-
-The trace reported 73.87% Occupancy Manager target, an 87.12% Compute Shader
-Launch limiter, and a 32.09% Last Level Cache limiter with 31.94% utilization.
-The staged pipelines used 12-14 allocated registers and spilled zero bytes.
-Those counters supported prototyping dispatch fusion; they did not establish
-that fusion would be faster.
-
-The prototype used one threadgroup per row or column, dynamic threadgroup
-memory, and separate horizontal and vertical kernels. Forced staged/fused
-tests matched within the existing irreversible tolerance across degenerate,
-odd/even, origin-parity, and 1023 x 767 cases. A pipeline-relative sweep then
-tested 1, 2, 4, 8, 16, and 32 SIMD groups per threadgroup:
-
-| SIMD groups | Threads on M4 Pro | Exploratory median | IQR |
-| ---: | ---: | ---: | ---: |
-| 1 | 32 | 1.658 ms | 0.227 ms |
-| 2 | 64 | 1.563 ms | 0.360 ms |
-| 4 | 128 | 1.437 ms | 0.298 ms |
-| 8 | 256 | 1.667 ms | 0.256 ms |
-| 16 | 512 | 1.686 ms | 0.207 ms |
-| 32 | 1024 | 1.745 ms | 0.701 ms |
-
-Because the 128-thread exploratory row appeared promising, it was remeasured
-in five alternating baseline/candidate processes with eleven warmed samples
-per process. Each cell reports process median / IQR:
-
-| Process | Staged baseline | 128-thread fused prototype |
-| ---: | ---: | ---: |
-| 1 | 1.422 / 0.381 ms | 1.641 / 0.300 ms |
-| 2 | 1.467 / 0.262 ms | 1.649 / 0.557 ms |
-| 3 | 1.629 / 0.261 ms | 1.677 / 0.288 ms |
-| 4 | 1.383 / 0.354 ms | 1.459 / 0.378 ms |
-| 5 | 1.324 / 0.410 ms | 1.646 / 0.373 ms |
-
-The median of process medians regressed from 1.422 ms to 1.646 ms (+15.75%).
-The initial maximum-thread configuration also regressed, from 1.550 ms to
-1.936 ms (+24.90%). Both miss the required 5% improvement by a wide margin.
-The kernel prototype and its temporary routing/tests were therefore removed.
-The end-to-end 512/1024 decode rows were not run because the targeted gate had
-already required rejection; no routing or performance claim changed.
-
-Changed-line coverage could not be reported. The canonical command
-`cargo xtask coverage metal --base HEAD^` failed before compiling coverage
-tests because the repository's Metal lane supplied `cargo-llvm-cov 0.8.7`
-with mutually incompatible `--no-report` and `--no-clean` options. The manual
-capture body is also intentionally skipped when coverage instrumentation is
-active so a coverage run cannot create a GPU trace. Package tests and strict
-Clippy passed, but this experiment therefore has no changed-line percentage.
-
-The permanent guard can be rerun with:
-
-```bash
-J2K_REQUIRE_METAL_RUNTIME=1 CARGO_TARGET_DIR=<isolated-target> \
-  cargo test -p j2k-metal metal_irreversible_idwt_perf_guard \
-  --profile release-bench --locked -- \
-  --ignored --nocapture --test-threads=1
-```
-
-The ignored GPU-capture test additionally requires `MTL_CAPTURE_ENABLED=1`
-and an absolute, nonexistent `.gputrace` path in
-`J2K_METAL_CAPTURE_PATH`. Captured timings are diagnostic only and are not
-used for acceptance.
-
 ## Metal Status
 
 Metal acceleration is selective. Public claims should say Metal-accelerated