feat(cuda): runtime resolution-aware kernel variant dispatch (ADR-0753)

[lusoris]

Summary

• Introduces vmaf_cuda_workload_class(w, h) — a pure-C classifier mapping luma pixel count to WS_SMALL (< 720p) / WS_MEDIUM (720p–4K) / WS_LARGE (>= 4K).
• Wires the first concrete consumer: adm_cm_device() in integer_adm_cuda.c selects adm_cm_line_kernel_8 (with __launch_bounds__(128,8)) at WS_MEDIUM and adm_cm_line_kernel_8_no_bounds otherwise — recovering the −9.3% 1080p gain without regressions at 576p or 4K.
• Policy and design rationale documented in ADR-0753 + Research-0753.

Why DRAFT

Architecture review needed before retrofitting other kernels. The dispatch infrastructure is intentionally minimal; future consumers (filter1d WS_LARGE split, motion CPU fallback at WS_SMALL) follow the same pattern.

Policy table (ADR-0753)

Optimisation	WS_SMALL	WS_MEDIUM	WS_LARGE
adm_cm __launch_bounds__(128,8)	SKIP	APPLY	SKIP
filter1d __ldg + __launch_bounds	SKIP	APPLY	APPLY
ms_ssim_decimate smem tiling	SKIP	SKIP	SKIP

Reproducer / smoke test

# Build with CUDA
meson setup build -Denable_cuda=true && ninja -C build

# 576p — WS_SMALL path: should pick adm_cm_line_kernel_8_no_bounds
./build/tools/vmaf --feature adm --backend cuda \
  --reference python/test/resource/yuv/src01_hrc00_576x324.yuv \
  --distorted python/test/resource/yuv/src01_hrc01_576x324.yuv \
  --width 576 --height 324 --pixel_format yuv420p --bitdepth 8

# 1080p — WS_MEDIUM path: should pick adm_cm_line_kernel_8
./build/tools/vmaf --feature adm --backend cuda \
  --reference python/test/resource/yuv/checkerboard_1920_1080_10_3_0_0.yuv \
  --distorted python/test/resource/yuv/checkerboard_1920_1080_10_3_0_1_0.yuv \
  --width 1920 --height 1080 --pixel_format yuv420p --bitdepth 10

# Parity gate (must pass at places=4)
make test-netflix-golden

Six deliverables (ADR-0108)

• Research digest: docs/research/0753-cuda-resolution-aware-dispatch-design.md
• Decision matrix: ADR-0753 ## Alternatives considered (compile-time, auto-tune, single-variant, per-SM-count)
• AGENTS.md invariant: core/src/feature/cuda/AGENTS.md — "How to add a new resolution-aware variant"
• Reproducer: build commands above
• Changelog: changelog.d/added/cuda-resolution-aware-dispatch.md
• Rebase notes: docs/rebase-notes.md entry for ADR-0753

Files touched

• core/src/feature/cuda/resolution_dispatch.{h,c} — new, pure C
• core/src/feature/cuda/integer_adm/adm_cm.cu — bounded + no-bounds macro variants
• core/src/feature/cuda/integer_adm_cuda.c — include, struct field, init, dispatch
• core/src/meson.build — register resolution_dispatch.c

🤖 Generated with Claude Code

[lusoris]

Extended to wire filter1d + ssim_vert_combine per the same policy. Now 3 kernels use the dispatch (adm_cm, filter1d, ssim_vert_combine).

[lusoris]

Validation at 576p (WS_SMALL) — Research-0755 (PR #103):

Per-kernel nsys (CUPTI timestamps, 48 frames, src01 576×324):

Kernel	Baseline avg	NO_BOUNDS avg	Delta
adm_cm_line_kernel_8	24.66 µs	21.90 µs	−11.2%
filter1d_8_horizontal	18.69 µs	17.96 µs	−3.9%
ssim_vert_combine	4.61 µs	4.71 µs	+2.2% (noise)

Correctness: VMAF mean 76.6678 (both builds) — passes ADR-0214 places=2.

End-to-end wall time: Inconclusive. Five concurrent vmaf CUDA processes on the GPU during measurement produced ±18% run-to-run variance, masking any signal. An idle-GPU re-run is needed before quoting a wall-time number.

Verdict: NO_BOUNDS dispatch confirmed beneficial at SMALL for adm_cm (−11.2%) and filter1d_8_horizontal (−3.9%). Dispatch mechanism is correct. Recommend marking ready after idle-GPU end-to-end re-run.

Side note — master defect found: origin/master (70cb42a) has a committed merge conflict marker at core/src/feature/cuda/integer_vif_cuda.c line 336 (from commit 0c494cca05, post-merge-train sweep). This makes master CUDA-build-broken without a manual patch. Should be fixed as a hotfix before PR #91 merges.

Full digest: #103