perf(cuda): size the ADM AIM CM launch by SM count instead of a fixed 8 rows (ADR-1226)

[lusoris]

Summary

adm_cm_aim_line_kernel_8 was flagged for register pressure: 255 registers, a
344-byte stack frame, 412/404 bytes of spill on sm_89, capping a 128-thread
block at two blocks per SM (16.7% occupancy).

Measuring the obvious fix first changed the fix, twice.

__launch_bounds__(128, 5) does not trade spill for occupancy — capped at 96
registers it spills less than the unconstrained build (8 B stack vs 344 B),
because at 255 registers ptxas keeps the whole unrolled theta × row working set
live instead of rematerialising it.

And it bought nothing: 0.808 ms vs 0.803 ms per call. Occupancy was never the
constraint. A 1080p frame gives buffer_h = 434, so at 8 rows per thread the
launch is ceil(434/32) × 3 = 42 blocks — against an RTX 4090's 128 SMs.
Blocks-per-SM is irrelevant when there is less than one block per SM to place.

The kernel's only parallelism is one block per BLOCKY * rows_per_thread rows ×
three orientation bands. There is no x-decomposition and there cannot cheaply be
one: each row's warp reduction must cover the whole row before the single
>> shift_inner_accum rounding, so splitting a row across blocks would round
differently from the CPU. So rows_per_thread is the knob — and lowering it is
arithmetically free, since each row is still reduced across all its columns
inside one block.

Measurements (RTX 4090, CUDA 13.3, 48 frames, warm)

Kernel time per call, host clock around the launch + stream sync:

resolution	rows=8	rows=4	rows=2	rows=1
1920x1080	0.801	0.553	0.586	0.622
640x480	0.299	0.203	0.159	—
576x324	0.300	0.178	0.141	—

The optimum tracks block count, not frame size — hence the SM-count-relative
rule (4 when that fills at least half the SMs, else 2) rather than a resolution
threshold. rows=1 regresses even at 1080p (327 blocks), past the point where
more blocks pay for the per-thread setup.

Whole-feature adm (CUDA), mean ms per frame:

resolution	before	after	delta
1920x1080	1.31	1.06	−19.1%
640x480	0.46	0.33	−28.3%
576x324	0.39	0.27	−30.8%

Numerics unchanged: vmaf_bench --validate reports identical CPU-vs-CUDA
max_diff values before and after, to the last digit, for adm2_score and all
four integer_adm_scale* outputs (5.42e-08 / 8.26e-08 / 3.15e-08 / 1.13e-07 /
4.65e-08). meson test --suite=fast: 0 failures.

__launch_bounds__ is left off deliberately — stacked on this fix it is a 3.4%
regression (0.572 vs 0.553 ms). The kernel carries a comment saying so, so it is
not helpfully re-added.

Reproducer / smoke-test command

meson setup core/build-cuda core -Denable_cuda=true -Denable_nvcc=true \
  -Denable_float=true -Db_lto=false --buildtype=release
ninja -C core/build-cuda

# Fixtures: ref_/dis_<W>x<H>.yuv (48 frames) in one directory
VMAF_TEST_DATA=/path/to/fixtures core/build-cuda/tools/vmaf_bench \
  --resolution 1920x1080 --frames 48 --gpu-only        # adm (CUDA) row
VMAF_TEST_DATA=/path/to/fixtures core/build-cuda/tools/vmaf_bench \
  --validate --resolution 1920x1080 --frames 8         # ADM/CU max_diff rows

# Register/spill counts behind the table:
nvcc --fatbin -gencode=arch=compute_89,code=sm_89 -Xptxas -v \
  core/src/feature/cuda/integer_adm/adm_cm.cu -o /dev/null <the build's -I set>

Discard the first three benchmark runs — the 4090 clocks up over roughly that
many.

Deep-dive deliverables (ADR-0108)

• Research digest — no digest needed: the investigation is the ADR's ## Measurements section, which carries the full register/spill sweep, the per-kernel timings and the whole-feature deltas.
• Decision matrix — ADR-1226 ## Alternatives considered (five options, including the x-decomposition that would break CPU bit-exactness).
• AGENTS.md invariant note — no rebase-sensitive invariants beyond the rebase note below, which carries all four.
• Reproducer / smoke-test command — above.
• Changelog fragment — changelog.d/changed/1226-cuda-adm-cm-aim-grid-occupancy.md, CHANGELOG.md regenerated.
• Rebase note — docs/rebase-notes.md § ADR-1226: the _8 entry point is gone, __launch_bounds__ is absent on purpose, rows_per_thread must not go back up, and sm_count == 0 is a supported state.

Docs (rule 10)

No user-discoverable surface changes: no CLI flag, no public header, no build
option, no output-schema change. The performance delta is in the changelog.

State (rule 13)

no state delta: performance change, no bug opened or closed.

🤖 Generated with Claude Code