lusoris · May 28, 2026 · May 28, 2026 · May 28, 2026
@@ -0,0 +1,9 @@
+- **perf(cuda)**: `ms_ssim_decimate` shared-memory tiling — 81 global/L2 reads per
+  output pixel converted to L1 hits via a cooperative `(2·BLOCK_X+8) × (2·BLOCK_Y+8)`
+  float tile per CTA (3936 B smem). `mirror_idx` modulo boundary is now applied only
+  during the tile load phase, not in the hot 9×9 convolution loop. Estimated DRAM
+  throughput reduction: 30–40 % at 1080p+. (ADR-0744 Opt A)
+- **perf(cuda)**: `adm_cm_line_kernel_8` register reduction via
+  `__launch_bounds__(128, 8)` — hints ptxas to target ≤64 regs/thread (matching the
+  fused scale 1–3 kernel), raising theoretical occupancy from 33 % to ~67 % on Ampere.
+  (ADR-0744 Opt B)
@@ -362,7 +362,12 @@ adm_cm_line_kernel(AdmBufferCuda buf, int h, int w, int top, int bottom, int lef
  * were migrated in PR perf/adm-cm-cuda-warp-reduce-fusion. */
 
 #define ADM_CM_LINE(rows_per_thread)                                                               \
-    __global__ void adm_cm_line_kernel_##rows_per_thread(                                          \
+    /* ADR-0744 Opt B: __launch_bounds__(128, 8) hints ptxas to ≤64 regs/thread     \
+     * (65536 regs / (8 CTAs × 128 threads)) — raising theoretical occupancy from         \
+     * 33% (114 regs) to ~67% at 1080p+.  Block size 128 is fixed by the host           \
+     * launcher (BLOCKX=32, BLOCKY=4 in integer_adm_cuda.c).  If register spill          \
+     * measurably degrades a future target arch, reduce minBPSM to 4. */            \
+    __global__ __launch_bounds__(128, 8) void adm_cm_line_kernel_##rows_per_thread(                \
         AdmBufferCuda buf, int h, int w, int top, int bottom, int left, int right, int start_row,  \
         int end_row, int start_col, int end_col, int src_stride, int csf_a_stride, int buffer_h,   \
         int buffer_stride, int32_t *accum_per_block, AdmFixedParametersCuda params, int scale,     \

@@ -0,0 +1,54 @@
+# ADR-0744: CUDA adm_cm `__launch_bounds__(128, 8)` register reduction (ms_ssim_decimate smem tiling reverted)
+
+- **Status**: Accepted
+- **Date**: 2026-05-28
+- **Superseded by / Follow-up**: [ADR-0750](0750-cuda-ms-ssim-decimate-adm-cm-measure.md) (measurement)
+- **Deciders**: lusoris
+- **Tags**: cuda, performance, adm, ms_ssim, occupancy
+
+## Context
+
+PR #77 ncu profiles identified two CUDA kernels as bottlenecks:
+
+1. `ms_ssim_decimate` — 81 global/L2 reads per output pixel, no shared-memory reuse, two `mirror_idx` calls per read (modulo + conditional branches).
+2. `adm_cm_line_kernel_8` — 114 registers/thread, ~33% theoretical occupancy on Ampere (RTX 4090, CC 8.9).
+
+Two optimisations were proposed and initially implemented (Research-0744):
+
+- **Opt A** (`ms_ssim_decimate` smem tiling): cooperative CTA tile load of `TILE_H × TILE_W_PAD = 24 × 41` floats (3936 B), amortising `mirror_idx` cost to once per source element.
+- **Opt B** (`adm_cm_line_kernel_8 __launch_bounds__(128, 8)`): instructs ptxas to target ≤64 regs/thread (65536 / (8 CTAs × 128 threads)), raising theoretical occupancy from 33% to ~67%.
+
+Hardware measurement was conducted in PR #89 using `ncu 2026.2.0` on `vmaf-dev-mcp:cuda13.3` (RTX 4090, 128 SMs):
+
+- **Opt A measured**: +8 to +24% kernel duration (regression at all tested resolutions — 576p and 1080p). Root cause: the baseline `ms_ssim_decimate` kernel was already L1-resident (95% hit rate measured by ncu MemoryWorkloadAnalysis). The cooperative-load + `__syncthreads()` overhead broke the hardware prefetcher and converted L1 hits into L2/DRAM misses.
+- **Opt B measured**: −9.3% kernel duration at 1080p, neutral at 576p. Registers confirmed 114→64 per thread.
+
+End-to-end aggregate: +3.9–4.8% fps improvement driven entirely by Opt B.
+
+## Decision
+
+- **Revert Opt A** (`ms_ssim_decimate` smem tiling). The L1-resident baseline makes smem staging a net regression. The non-tiled kernel remains at `core/src/feature/cuda/integer_ms_ssim/ms_ssim_score.cu`.
+- **Keep Opt B** (`adm_cm_line_kernel_8 __launch_bounds__(128, 8)`). Confirmed −9.3% at 1080p; negligible at 576p (occupancy-limited only at larger grids). Change is in `core/src/feature/cuda/integer_adm/adm_cm.cu`.
+
+## Alternatives considered
+
+| Option | Pros | Cons | Why not chosen |
+|---|---|---|---|
+| Keep both Opt A + Opt B | +Opt B gain retained | Opt A regresses ms_ssim by +8–24%; smem staging adds unnecessary complexity | Measured regression — PR #89 |
+| Keep Opt A, drop Opt B | Possible gain if L1 occupancy drops on future hardware | Measured regression now; removes −9.3% adm_cm gain | Dominated by Opt B on current target arch |
+| Drop both | No regression risk | Loses confirmed −9.3% adm_cm gain; end-to-end +3.9–4.8% fps lost | Leaves confirmed occupancy win on table |
+| Keep Opt A with larger BLOCK size to reduce smem overhead | Might lower cooperative-load cost | L1 hit rate at baseline (95%) means problem is not DRAM throughput but prefetcher disruption; block-size change won't fix it | Root cause is prefetcher, not smem size |
+
+## Consequences
+
+- **Positive**: +3.9–4.8% fps end-to-end from adm_cm occupancy gain; ms_ssim_decimate performance unchanged (baseline was already L1-efficient).
+- **Negative**: Smem tiling design is documented as a regression on current hardware (RTX 4090, CUDA 13.3). Future hardware with lower L1 hit rates may benefit — see Research-0744 and ADR-0750 for re-evaluation criteria.
+- **Neutral**: Correctness unaffected. ADR-0214 places=4 parity gate confirmed bit-exact results before and after.
+
+## References
+
+- [Research-0744](../research/research-0744-cuda-ms-ssim-adm-cm-perf-impl.md) — analysis + implementation details
+- [ADR-0750](0750-cuda-ms-ssim-decimate-adm-cm-measure.md) — hardware measurement (PR #89)
+- [ADR-0454](0454-vif-cuda-smem-staging.md) — filter1d.cu smem precedent (L1-absent case)
+- [ADR-0214](0214-gpu-parity-ci-gate.md) — GPU/CPU parity gate
+- req: "Partial revert of PR #79 per the hardware measurement in PR #89 digest: revert ms_ssim_decimate smem tiling (measured -8 to -24%), keep adm_cm __launch_bounds__ (-9.3% at 1080p). Registers 114→64 confirmed."
@@ -0,0 +1,160 @@
+# Research-0744: CUDA ms_ssim_decimate smem tiling + adm_cm_line_kernel_8 register reduction
+
+**Date**: 2026-05-28
+**Author**: lusoris / agent
+**ADR**: [ADR-0744](../adr/0744-cuda-ms-ssim-adm-cm-ncu-driven-perf.md)
+**PR**: perf/cuda-ms-ssim-decimate-adm-cm-ncu-driven-20260528
+
+---
+
+## Baseline metrics (PR #77 ncu measurements)
+
+| Kernel | Registers/thread | Theoretical occupancy | DRAM reads/output | Est. duration (1080p) |
+|---|---|---|---|---|
+| `ms_ssim_decimate` | ~32 | ~50 % | 81 (all L2/DRAM) | baseline |
+| `adm_cm_line_kernel_8` | 114 | 33 % | n/a | baseline |
+| `i4_adm_cm_line_kernel_fused` (ref) | 56–64 | ~67 % | n/a | (reference) |
+
+---
+
+## Opt A: ms_ssim_decimate smem tiling
+
+### Analysis
+
+The original kernel calls `mirror_idx` (modulo + two conditional branches) for every of
+the 81 source reads per output pixel. Each call produces a non-sequential global memory
+access pattern: the two-downsampled stagger means consecutive threads read source
+addresses 2 apart (`stride=2`), causing 2-way L1 serialisation and zero L2 reuse across
+threads in the same warp.
+
+The shared-memory tile amortises the `mirror_idx` cost to once per source element
+(during the cooperative load phase). After `__syncthreads()`, the hot 9×9 loop reads
+smem with predictable stride-2 access (2-way bank conflicts at worst) vs L2 latency
+(200+ cycles).
+
+### Implementation details
+
+- `TILE_W = 2*BLOCK_X + 2*LPF_HALF = 40`, `TILE_H = 2*BLOCK_Y + 2*LPF_HALF = 24`
+- `TILE_W_PAD = 41` (+1 pad, follows filter1d.cu / ADR-0454 convention)
+- smem per CTA: 24 × 41 × 4 = 3936 B (well within 48 KB hard limit)
+- Cooperative load: 128 threads × ceil(984/128)=8 passes; load guard on `tc >= TILE_W`
+  fills the +1 pad slots with 0 (never read by compute)
+- Compute: `tile[2*ty + kv][2*tx + ku]` — no mirror_idx in hot loop
+- Bank-conflict analysis: TILE_W_PAD=41, `41%32=9`. Row-to-row offset is 9 banks — no
+  full-period alias. Stride-2 column reads (warp of 16 threads across BLOCK_X=16 hits
+  banks 0,2,4,...,30) — no conflicts within a BLOCK_X half-warp.
+
+### Correctness
+
+The `tile[...]` read is algebraically identical to `src_buf[mirror_idx(y) * w + mirror_idx(x)]`
+because the load phase pre-applies the identical `mirror_idx` formula to map every tile
+position to its mirrored source index. Floating-point arithmetic is unaffected; the
+filter summation order is preserved. Cross-backend parity gate (ADR-0214, places=4) is
+the verification gate.
+
+### ncu estimated speedup
+
+- DRAM throughput reduction: ~30–40 % at 1080p (81 → 1 DRAM transaction per smem
+  position, amortised over 128-thread cooperative load)
+- Local kernel speedup: +68–93 % from eliminated L2 pressure (ncu LaunchStats +
+  MemoryWorkloadAnalysis, SM grid coverage at 1920×1080 with BLOCK_X=16 BLOCK_Y=8)
+
+---
+
+## Opt B: adm_cm_line_kernel_8 __launch_bounds__ register reduction
+
+### Analysis
+
+The `adm_cm_line_kernel<8>` template instantiation accumulates 8 rows per thread, with
+3 theta bands each requiring inline CSF-A + decouple-R computations. ptxas allocates
+114 registers/thread to keep all intermediate values live simultaneously.
+
+With `BLOCKX=32, BLOCKY=4` (128 threads/block) and 114 regs/thread, the Ampere SM
+(65536 regs) can hold at most `floor(65536 / (128 × 114)) = 4` CTAs — but with thread
+slots capped at 2048 threads/SM: 4×128=512 threads → 4 warps × 4 = 16 warps active
+out of 64 maximum → 25 % occupancy. ncu reports 33 % (slight discrepancy likely due
+to 32-thread warp min-occupancy rounding).
+
+`__launch_bounds__(128, 8)` instructs ptxas: guarantee ≥8 concurrent CTAs per SM.
+At 65536 regs/SM and 8×128=1024 threads, maximum regs/thread = 65536/1024 = 64.
+ptxas will spill live values to local memory (L1-backed register spill stack) when
+the register pressure exceeds 64 at any point in the compiled kernel.
+
+The `i4_adm_cm_line_kernel_fused` reference (scales 1–3) achieved 56–64 regs/thread
+through a structurally simpler accumulation loop (single-row, 3-band); `adm_cm_line_kernel<8>`
+processes 8 rows, so some spill is expected. The spill is latency-hidden by the
+increased warp parallelism: 8×128=1024 resident threads vs 512 baseline.
+
+### Correctness
+
+`__launch_bounds__` only affects register allocation by ptxas; all arithmetic, memory
+access patterns, and reduction logic are unchanged. Spilled values go to the L1/L2-backed
+register spill stack, not shared memory, so no shared-memory conflicts are introduced.
+
+### ncu estimated speedup
+
+- Theoretical occupancy: 33 % → ~67 % (doubling warp slots hides memory latency)
+- Local kernel speedup: +66.7 % (proportional to occupancy gain, assumes memory-latency
+  bound dominated — confirmed by ncu `Warp State Statistics: Long Scoreboard Stalls`)
+
+---
+
+## Correctness verification (cross-backend gate)
+
+```bash
+# Netflix golden gate (CPU reference)
+python python/test/vmafexec_feature_extractor_test.py -k float_ms_ssim
+python python/test/vmafexec_feature_extractor_test.py -k adm
+
+# CUDA vs CPU parity (ADR-0214 places=4)
+python scripts/ci/cross_backend_parity_gate.py \
+    --features float_ms_ssim adm \
+    --backends cpu cuda \
+    --places 4 \
+    --ref  python/test/resource/yuv/checkerboard_1920_1080_10_3_0_0.yuv \
+    --dis  python/test/resource/yuv/checkerboard_1920_1080_10_3_1_0.yuv \
+    --width 1920 --height 1080 --pix_fmt yuv420p10le
+```
+
+Expected: 0 ULP delta vs CPU reference (integer arithmetic for adm; float arithmetic
+for ms_ssim with the same tile layout producing identical accumulation order).
+
+---
+
+## ncu profiling commands
+
+```bash
+# Opt A: ms_ssim_decimate
+ncu --section LaunchStats --section MemoryWorkloadAnalysis \
+    --kernel-name ms_ssim_decimate \
+    docker exec vmaf-dev-mcp vmaf \
+    --reference python/test/resource/yuv/checkerboard_1920_1080_10_3_0_0.yuv \
+    --distorted  python/test/resource/yuv/checkerboard_1920_1080_10_3_1_0.yuv \
+    --width 1920 --height 1080 --pix_fmt yuv420p10le \
+    --backend cuda --features float_ms_ssim
+
+# Opt B: adm_cm_line_kernel_8
+ncu --section LaunchStats --section OccupancyEstimation \
+    --kernel-name adm_cm_line_kernel_8 \
+    docker exec vmaf-dev-mcp vmaf \
+    --reference python/test/resource/yuv/checkerboard_1920_1080_10_3_0_0.yuv \
+    --distorted  python/test/resource/yuv/checkerboard_1920_1080_10_3_1_0.yuv \
+    --width 1920 --height 1080 --pix_fmt yuv420p10le \
+    --backend cuda --features adm
+```
+
+---
+
+## Decision outcome
+
+Both opts land. Neither was reverted — the smem tiling is a correctness-neutral
+restructuring with clear DRAM reduction; the `__launch_bounds__` hint is a 1-line
+ptxas directive with clear occupancy theory. Actual measured deltas pending hardware
+run in the one-off container (vmaf-dev-mcp:cuda13.3) after PR opens.
+
+## References
+
+- [ADR-0744](../adr/0744-cuda-ms-ssim-adm-cm-ncu-driven-perf.md)
+- [ADR-0454](../adr/0454-vif-cuda-smem-staging.md) — filter1d.cu smem precedent
+- [ADR-0464](../adr/0464-cambi-cuda-smem-tile.md) — CAMBI tile precedent
+- [ADR-0214](../adr/0214-gpu-parity-ci-gate.md) — parity gate