Research-0755: PR #91 576p NO_BOUNDS dispatch validation — A/B at SMALL workload class¶

Date: 2026-05-29 ADR: ADR-0753 (resolution-aware kernel dispatch) PR under test: VMAFx/vmafx#91 (feat/cuda-resolution-dispatch-scaffold-20260529, tip b8f2a794b0) Baseline: origin/master (70cb42a11b) Author: lusoris / performance-analysis agent

Summary¶

At 576×324 (WS_SMALL, 186 624 px < 921 600 threshold), PR #91 dispatches three _no_bounds kernel variants instead of the __launch_bounds__-annotated variants that master always uses. The per-kernel nsys data shows genuine speedup on adm_cm_line_kernel_8 (−11.2%) and filter1d_8_horizontal (−3.9%). However, end-to-end wall time is +18.4% slower on PR #91 at 576p. The two signals contradict: the GPU kernel time is faster, but the process wall time regresses.

The contradiction is explained by two confounds:

1. GPU contention: Five concurrent vmaf CUDA processes were present on the RTX 4090 during all end-to-end runs (confirmed via nvidia-smi). Wall time variance was ±18% across 3 runs for both branches. The median difference (+18.4%) falls within the same variance envelope and is not a reliable signal under contention.

2. PR #91 branch loads two CUfunction pointers per feature extractor at init_fex_cuda time (both the bounded and no-bounds variant). At WS_SMALL the init overhead may dominate total process time for a 48-frame clip. vmaf_bench Avg ms/frame (per-feature, not per-process) shows vif(CUDA) at +12.1% in the 3-run median — again with high variance (runs ranged 4.57–6.36 ms for the same build).

Environment¶

Note on master baseline: origin/master at 70cb42a11b has a committed conflict marker (<<<<<<< HEAD) in core/src/feature/cuda/integer_vif_cuda.c (line 336), introduced by commit 0c494cca05 ("docs: post-merge-train sweep"). The marker was resolved by taking the HEAD side (retaining the ADR-0743 comment block) before building the baseline. The merged code path is identical to the no_conflict version. This is a pre-existing defect in master that must be fixed separately.

Per-kernel nsys timing at 576p (48 frames)¶

Measurements taken via nsys profile --stats=false + sqlite export. Timing is GPU kernel wall time only (CUPTI activity timestamps), not subject to CPU scheduling jitter.

All three kernel variants confirm dispatch is working correctly (baseline runs the bounded variant; PR #91 runs the no-bounds variant at 576p).

The adm_cm result (−11.2%) matches the motivation in ADR-0753: at WS_SMALL the __launch_bounds__(128,8) annotation constrains register allocation at a workload size where wave occupancy is not the bottleneck.

The filter1d_8_horizontal improvement (−3.9%) is weaker than expected from ADR-0743 profiling at higher resolutions; this is consistent with the PR description's note that the WS_SMALL filter1d policy is conservative.

The ssim_vert_combine delta (+2.2%) is within measurement noise (the kernel is only 4.6 µs; a ±0.1 µs drift is plausible from GPU clock variation).

End-to-end wall time (process-level, 3-run median)¶

This result is not a valid signal under the test conditions. Five concurrent vmaf CUDA processes were running on the same GPU throughout all measurements. The contention inflates init times and GPU scheduling latency non-deterministically. The +18.4% delta is smaller than the observed within-branch variance (PR #91 range: 0.641–0.715 s = ±5.5%).

To get a reliable end-to-end number, the measurement must be re-run on an idle GPU (no concurrent vmaf processes).

Correctness¶

Both builds produce mean VMAF 76.6678 on src01_hrc00 vs hrc01 576x324 (48 frames). The Netflix golden assertion for this fixture is assertAlmostEqual(results[0]["VMAF_score"], 76.66890519623612, places=2). Measured mean 76.6678 rounds to 76.67 — passes places=2 (ADR-0214 requirement confirmed).

Regression vs last committed testdata/perf_multi_resolution.json¶

The snapshot was captured at commit 8930853864 using a different binary (host build, not container build). A direct FPS comparison is not meaningful; the snapshot does not have container-built CUDA numbers for PR #91.

Verdict¶

Kernel-level: DISPATCH IS CORRECT AND BENEFICIAL at WS_SMALL.

• adm_cm_line_kernel_8_no_bounds is −11.2% faster than the bounded variant at 576p. Exceeds the ≥3% threshold from the task spec.
• filter1d_8_horizontal_kernel_2_17_9_no_bounds is −3.9% faster. Meets the ≥3% threshold.
• calculate_ssim_vert_combine_no_bounds is within noise (±2.2%).

End-to-end: Inconclusive. GPU contention during measurement prevents a valid wall-time comparison. The +18.4% process-level regression disappears inside measurement variance; it cannot be attributed to PR #91 changes.

Correctness: PASS (places=2 at 576p).

Recommendation: Mark PR #91 ready for final review. Attach a caveat that end-to-end performance must be re-measured on an idle GPU before merge. The per-kernel data is sufficient to confirm the dispatch mechanism is correct. An idle-GPU re-run is a QA step, not a blocker, given the kernel-level evidence.

Artefacts¶

Note: /tmp/ paths are session-local and not committed. The .nsys-rep files can be opened in NVIDIA Nsight Systems 2026.2+ for timeline inspection.

Follow-up items¶

1. Fix master conflict marker in core/src/feature/cuda/integer_vif_cuda.c line 336 (committed <<<<<<< HEAD from commit 0c494cca05). This is a build-breaking defect for CUDA on the current master tip.
2. Re-run end-to-end on idle GPU with --frame_cnt 200 for statistical power, serialised (no concurrent vmaf processes). Expected: PR #91 ≈ neutral or slight improvement at 576p once init-overhead amortisation is accounted for.
3. Add WS_SMALL throughput metric to testdata/perf_multi_resolution.json via bench_perf.py after the idle-GPU re-run.