Research-0756: CUDA F3 Struct-by-Value Kernel Audit¶

• Status: Active
• Workstream: perf/cuda-f3-struct-by-value-audit (no ADR — research only)
• Date: 2026-05-29
• References: PR #93 (F3 finding deferral), Research-0754, Research-0734 through 0738

Background¶

PR #93 (perf/cuda-ssim-vert-combine-ldg-launch-bounds-leak-20260529) identified "F3" as the pattern where a __global__ kernel accepts a VmafCudaBuffer (or similar aggregating struct) by value. This hides the underlying CUdeviceptr pointer from the compiler's alias analysis and prevents ptxas from classifying those loads as non-aliased reads eligible for:

• the read-only L1 texture cache path (ld.global.nc), and
• __restrict__-based alias-free optimisation within the inner loop.

The fix used in PR #93 for calculate_ssim_vert_combine was to extract const float *__restrict__ raw pointers from each VmafCudaBuffer argument before the inner loop and then read via __ldg(). That approach is measured at -4.2% kernel duration at 1080p (Research-0754).

This digest performs the fork-wide inventory and severity-ranks the affected kernels by their measured DRAM throughput from PR #77's ncu profiles.

Struct definitions¶

Full kernel inventory¶

AFFECTED — takes VmafCudaBuffer or similar struct by value¶

Partial F3 (VmafCudaBuffer is output-only or single read — lower severity):

SAFE-PTR — takes raw const T * __restrict__ pointers¶

Severity ranking (DRAM throughput from PR #77 ncu measurements)¶

The critical question for F3 severity is: does the kernel do significant global-memory reads in a hot inner loop, and are those reads currently missing the read-only L1 cache path?

Below top-5:

• calculate_psnr_kernel_*: VmafCudaBuffer sse is write-only; the VmafPicture inputs read at most 1 pixel/thread; launch-count and integer arithmetic dominate. F3 severity: LOW.
• calculate_moment_kernel_*: Same pattern as psnr — 2 loads per pixel, arith-light kernel, warp-reduce dominated. F3 severity: LOW.
• adm_cm_line_kernel_8 / i4_adm_cm_line_kernel_fused: Already receiving __launch_bounds__(128, 8) attention (PR #97). The AdmBufferCuda struct access pattern is indirect (buf.ref_dwt2.band_h[i*stride+j]); the inner loop is already register-limited, not DRAM-limited. Adding __ldg() here requires extracting ~6 per-band sub-pointers which is more invasive without proportionate DRAM gain.
• calculate_motion_score_kernel_8bpc: Input is VmafPicture.data[0] but is loaded once per tile into shared memory. Post-smem-load inner loop does not re-read global memory. F3 severity: NEGLIGIBLE.
• motion_v2_kernel_8bpc: VmafCudaBuffer sad is write-only (single atomic). Input args already const uint8_t *__restrict__. F3 severity: NONE (only output struct).

Fix strategies¶

In-kernel __ldg() extraction (low-risk, proven in PR #93)¶

// Before inner loop, once per kernel invocation:
const float *h_ref_mu = __ldg_ptr(reinterpret_cast<const float *>(h_ref_mu_buf.data));
// ... then use h_ref_mu[src_idx] throughout the loop

This is the approach PR #93 applied to calculate_ssim_vert_combine. It requires no call-site changes, is bit-identical (measured: 0.00e+00 max diff on all frames in Research-0754), and yields measurable improvement where the kernel is DRAM-bound at usable wave counts (1080p: -4.2% duration).

Host-side extraction (higher impact, more invasive)¶

Extract raw CUdeviceptr values from each VmafCudaBuffer on the host before cuLaunchKernel, then pass as raw kernel arguments. This removes the struct entirely from the kernel signature and allows ptxas to see const float * __restrict__ directly, enabling full ld.global.nc without __ldg(). Requires changing both the .cu kernel signature and every cuLaunchKernel call site in the corresponding _cuda.c dispatcher. Higher effort, not bit-risk, but more invasive. Recommended only where the kernel is also a candidate for AoS→SoA restructuring (F1, deferred in PR #93).

Top-5 PRs to dispatch from this audit¶

Kernels confirmed already fixed¶

• calculate_ssim_vert_combine (ssim_score.cu:153) — fixed in PR #93 / Research-0754. Measured -4.2% kernel duration at 1080p.

ncu profile commands for unvalidated kernels¶

# ms_ssim_vert_lcs (needs separate launch-count to capture past decimate)
ncu -k "regex:ms_ssim_vert_lcs" --set basic --launch-count 10 \
  core/build-ncu/tools/vmaf --feature float_ms_ssim_cuda --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_1_0.yuv \
  --width 1920 --height 1080 --pixel_format 420 --bitdepth 8 -o /dev/null
# ncu section: LaunchStats + MemoryWorkloadAnalysis

# ciede (1080p, both bpc variants)
ncu -k "regex:calculate_ciede" --set basic --launch-count 6 \
  core/build-ncu/tools/vmaf --feature ciede_cuda --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_1_0.yuv \
  --width 1920 --height 1080 --pixel_format 420 --bitdepth 8 -o /dev/null

Summary¶

Total AFFECTED kernels (F3 pattern present): 20 kernel variants across 8 families.

Of those, high severity (inner hot loop, not smem-tiled, no existing pointer extraction): 5 (ms_ssim_vert_lcs, ms_ssim_horiz, ciede 8/16bpc, adm_decouple).

Already fixed: calculate_ssim_vert_combine (PR #93).

Dispatch priority: PR-1 (ms_ssim_vert_lcs) first — identical pattern to PR #93, highest confidence of a measurable win at 1080p.

Related¶

• PR #93 / Research-0754 (F3 first fix — ssim_vert_combine)
• Research-0736 (SSIM ncu hotpath)
• Research-0737 (MS-SSIM ncu hotpath)
• Research-0734 (ADM ncu hotpath)
• Research-0738 (cross-metric summary)
• ADR-0756 (this audit)