perf(cuda): ms_ssim_decimate smem tiling + adm_cm register reduction (ADR-0744)

[lusoris]

Summary

• Opt A: ms_ssim_decimate — adds a (2*BLOCK_X+2*LPF_HALF+1) x (2*BLOCK_Y+2*LPF_HALF) = 41x24 float shared-memory tile per CTA (3936 B). Converts 81 global/L2 reads per output pixel to L1 hits. mirror_idx boundary applied once in cooperative load phase; hot 9x9 convolution loop reads smem unconditionally. +1 width pad (TILE_W_PAD=41) prevents bank aliasing per ADR-0454 convention.
• Opt B: adm_cm_line_kernel_8 — adds __launch_bounds__(128, 8) to the ADM_CM_LINE macro. Hints ptxas to target <=64 regs/thread, raising theoretical occupancy from 33% to ~67% on Ampere, matching the fused scale 1-3 kernel.

ncu-estimated deltas (pending hardware measurement)

Kernel	Metric	Before	After (est.)
ms_ssim_decimate	DRAM reads/output px	81	~1 (amortised tile load)
ms_ssim_decimate	Local speedup (1080p)	baseline	+68-93%
adm_cm_line_kernel_8	Registers/thread	114	<=64
adm_cm_line_kernel_8	Theoretical occupancy	33%	~67%
adm_cm_line_kernel_8	Local speedup (1080p)	baseline	+66.7%

Reproducer / smoke test

# Correctness: Netflix golden gate (CPU)
make test-netflix-golden

# CUDA vs CPU parity (ADR-0214 places=4)
docker exec vmaf-dev-mcp 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

# ncu baseline vs optimised (run in vmaf-dev-mcp:cuda13.3)
ncu --section LaunchStats --section MemoryWorkloadAnalysis \
    --kernel-name ms_ssim_decimate \
    vmaf --reference ... --distorted ... --backend cuda --features float_ms_ssim

ncu --section LaunchStats --section OccupancyEstimation \
    --kernel-name adm_cm_line_kernel_8 \
    vmaf --reference ... --distorted ... --backend cuda --features adm

Deliverables checklist

• Research digest: docs/research/research-0744-cuda-ms-ssim-adm-cm-perf-impl.md
• ADR decision matrix: docs/adr/0744-cuda-ms-ssim-adm-cm-ncu-driven-perf.md (Alternatives considered)
• AGENTS.md invariant: core/src/feature/cuda/AGENTS.md (two new sections)
• Reproducer: ncu commands in research digest + PR description above
• Changelog fragment: changelog.d/perf/cuda-ms-ssim-adm-cm-ncu-driven.md
• Rebase notes: docs/rebase-notes.md entry added
• Per-surface docs: docs/backends/cuda/overview.md kernel performance section added
• docs/state.md perf row added
• ADR index row: docs/adr/README.md

Notes for reviewer

• DRAFT: actual ncu measurements on RTX 4090 (vmaf-dev-mcp:cuda13.3) needed before marking ready.
• No correctness risk: Opt A is a pure memory-access reorganisation (same arithmetic, same accumulation order); Opt B is a ptxas register-budget hint with no arithmetic change.
• If Opt B causes register spill that degrades throughput on a specific arch, reduce minBPSM from 8 to 4 in adm_cm.cu line ~365 — no other files need updating.

🤖 Generated with Claude Code

[lusoris]

Hardware measurement complete (Research-0749 / ADR-0750)

Device: RTX 4090 (CC 8.9, 128 SMs), CUDA 13.3, ncu 2026.2.0, vmaf-dev-mcp:cuda13.3.
Method: isolated worktrees, optimized libvmaf.so via LD_PRELOAD, privileged container for PMU access.
Correctness: bit-exact (max_diff = 0.0 at ADR-0214 places=4 on both workloads).

ms_ssim_decimate smem tiling — REGRESSION

The baseline kernel already achieves 95% L1 hit rate at both 576p and 1080p — it was never DRAM-bound. The smem cooperative-load breaks the hardware prefetcher, causing L1 hit rate to drop from 95% to 24%:

Workload	Baseline duration	Optimized duration	L1 hit (base)	L1 hit (opt)
WL1 576p, largest launch	5120 ns	6368 ns	95.5%	24.5%
WL2 1080p, largest launch	18656 ns	20224 ns	95.6%	23.9%

Kernel is +24% slower (WL1) and +8% slower (WL2). The estimated +68–93% was based on an incorrect DRAM-bound assumption.

adm_cm_line_kernel_8 launch_bounds — CONFIRMED POSITIVE

Workload	Baseline duration	Optimized duration	Registers (base)	Registers (opt)
WL1 576p	~14.2 µs	~14.9 µs (noise)	114	64
WL2 1080p	~68.4 µs	~62.0 µs (-9.3%)	114	64

Register reduction 114→64 (-44%) is confirmed. -9.3% kernel duration at 1080p is consistent with theoretical occupancy 33%→50%.

End-to-end (3-run median, --feature float_ms_ssim_cuda --backend cuda)

Workload	Baseline	Optimized	Delta
WL1: 576x324 48f	2469 fps	2587 fps	+4.8%
WL2: 1080p 3f	482 fps	501 fps	+3.9%

Positive end-to-end delta is driven by adm_cm alone.

Verdict: PARTIAL REVERT RECOMMENDED

• Keep: adm_cm_line_kernel_8 __launch_bounds__(128, 8) — measured -9.3% at 1080p, correctness passes.
• Revert: ms_ssim_decimate smem tiling — measured +8-24% kernel regression. The optimization path for ms_ssim_decimate at 1080p is occupancy tuning (baseline 80% occupancy) or horiz/vert fusion, not smem tiling.

Full digest at docs/research/0749-cuda-ms-ssim-decimate-adm-cm-1080p-measure.md (PR #89).

[lusoris]

Partial revert applied per PR #89 measurement: ms_ssim_decimate smem tiling reverted (L1-resident baseline made it a regression — +8 to +24% kernel duration; cooperative load broke hardware prefetcher). adm_cm launch_bounds(128, 8) kept (−9.3% at 1080p, registers 114→64 confirmed). End-to-end net: +3.9–4.8% fps from adm_cm occupancy gain alone. ADR-0744 updated to Accepted. Marking ready.

[lusoris]

4K measurement (Research-0751): adm_cm launch_bounds kernel delta at 3840x2160 is -0.3% (noise), compared to -9.3% at 1080p. The register-bound regime that the optimization targets is 8-32 waves; at 4K (32.2 waves) the scheduler is wave-saturated and the register reduction does not further improve throughput. End-to-end adm CUDA delta at 4K is +1.9% (within ±5% noise). Recommendation: ship the launch_bounds change -- it costs nothing at 4K, wins measurably at 1080p (-9.3% kernel, +3.9% end-to-end), and is bit-exact at all resolutions. The ms_ssim smem-tiling revert is confirmed correct at 4K (kernel already L1-resident, 88.1% active warps at full resolution).