docs(cuda): ADR-0756 F3 struct-by-value audit — 20 kernels, top-5 dispatch

changelog.d/perf/cuda-f3-struct-by-value-audit.md

@@ -0,0 +1,14 @@
+<!--
+  Copyright 2026 Lusoris
+  SPDX-License-Identifier: BSD-3-Clause-Plus-Patent
+-->
+## [perf] CUDA F3 struct-by-value kernel audit (ADR-0756)
+
+Fork-wide audit of every `__global__` kernel accepting a `VmafCudaBuffer`,
+`VmafPicture`, or `AdmBufferCuda` argument by value — the "F3" pattern
+identified in PR #93. Identifies 20 kernel variants across 8 metric families
+where the struct copy hides pointer aliasing information from ptxas and
+prevents `ld.global.nc` emission. Severity-ranked by measured DRAM throughput
+from PR #77 ncu profiles. Top-5 dispatch PRs defined; `ms_ssim_vert_lcs` is
+priority-1 (structurally identical to the already-fixed `ssim_vert_combine`,
+expected -4 to -6% duration at 1080p).

core/src/feature/cuda/AGENTS.md

@@ -502,6 +502,20 @@ See [ADR-0747](../../../../docs/adr/0747-cuda-extern-c-sweep.md).
   read-many intermediate buffer must follow the same pattern.
   See [ADR-0754](../../../../docs/adr/0754-cuda-ssim-vert-combine-ldg-pinned-leak.md).
 
+## F3 pattern — VmafCudaBuffer / VmafPicture / AdmBufferCuda by-value arguments (ADR-0756)
+
+- **Passing an aggregating struct (VmafCudaBuffer, VmafPicture, AdmBufferCuda) by value to a
+  `__global__` kernel hides the embedded pointer from ptxas alias analysis.** ptxas cannot
+  emit `ld.global.nc` (the read-only L1 texture path) for any load through a pointer derived
+  from a struct field copied onto the kernel argument stack.
+- **The correct fix for inner-loop reads:** extract `const T *__restrict__` raw pointers from
+  every struct argument BEFORE the hot inner loop, then read via `__ldg(&ptr[idx])`.
+  This is the F3 fix pattern from PR #93 (Research-0754 / ADR-0754) and is the mandatory
+  approach for any new kernel that reads from `VmafCudaBuffer` or similar structs in a loop.
+- **Audit status:** ADR-0756 / Research-0756 (2026-05-29) catalogued 20 kernel variants with
+  this pattern. Priority-1 dispatch: `ms_ssim_vert_lcs` (`integer_ms_ssim/ms_ssim_score.cu`).
+  See [ADR-0756](../../../../docs/adr/0756-cuda-f3-struct-by-value-audit.md).
+
 ## Pinned-host memory free invariant after `readback_free` (ADR-0754)
 
 - **`vmaf_cuda_kernel_readback_free` NULLs `rb->host_pinned` but does NOT free it.**

docs/adr/0756-cuda-f3-struct-by-value-audit.md

@@ -0,0 +1,64 @@
+<!--
+  Copyright 2026 Lusoris
+  SPDX-License-Identifier: BSD-3-Clause-Plus-Patent
+-->
+# ADR-0756: CUDA F3 struct-by-value kernel audit (scope + dispatch order)
+
+- **Status**: Accepted
+- **Date**: 2026-05-29
+- **Deciders**: lusoris
+- **Tags**: `cuda`, `perf`, `research`
+
+## Context
+
+PR #93 identified "F3" — a pattern where a `__global__` kernel accepts a
+`VmafCudaBuffer` (or `VmafPicture`, `AdmBufferCuda`) by value. Because
+`VmafCudaBuffer.data` is a `CUdeviceptr` (opaque integer) nested inside a
+struct copy, ptxas cannot infer that the underlying pointer is non-aliased and
+cannot emit `ld.global.nc` (the read-only L1 texture path). PR #93 applied
+the in-kernel `__ldg()` extraction fix to `calculate_ssim_vert_combine` and
+measured -4.2% kernel duration at 1080p (Research-0754).
+
+The scope of remaining instances across the CUDA kernel suite was not
+enumerated before PR #93 merged. This ADR records the fork-wide audit result
+and the chosen dispatch order for follow-on PRs.
+
+## Decision
+
+We will address the F3 pattern using the in-kernel `__ldg()` extraction
+strategy (extract raw `const T *` from each struct arg before the hot inner
+loop; read via `__ldg()`). We will apply this in severity order per
+Research-0756: `ms_ssim_vert_lcs` first (PR-1), then `ms_ssim_horiz`,
+`ciede_8bpc/16bpc`, `adm_decouple`. Kernels whose inner loop is already
+smem-resident after the tile-load phase (motion, adm_cm) are explicitly
+out of scope for F3 treatment — the DRAM-bound portion is the tile load,
+not a repeated inner-loop global read.
+
+## Alternatives considered
+
+| Option | Pros | Cons | Why not chosen |
+|---|---|---|---|
+| Host-side extraction (pass raw `CUdeviceptr` to `cuLaunchKernel`) | Full compiler visibility; enables ptxas to use `ld.global.nc` without `__ldg()` | Requires modifying every call site in `_cuda.c` dispatch files; higher risk, larger diff | Deferred; revisit if post-F3 ncu shows remaining DRAM headroom |
+| AoS → SoA buffer restructure (F1) | Best long-term coalescing | Structural change to `VmafCudaBuffer`; API impact | Deferred per Research-0754 decision section |
+| Skip F3 for all kernels that are launch-starved at 576p | Correct that at 576p F3 rarely matters | Ignores 1080p+ production workloads | Not chosen; fork targets 1080p/4K production |
+
+## Consequences
+
+- **Positive**: Each dispatched PR carries at most 30–40 LOC diff, is
+  bit-identical (zero score drift), and is independently deployable. The
+  in-kernel pattern was proven by PR #93 with a live ncu A/B.
+- **Negative**: Multiple separate PRs rather than one large refactor.
+  Call-site code in `_cuda.c` files still passes structs; the compiler
+  sees the `__ldg()` hint but not full `__restrict__` aliasing. Full
+  benefit requires the eventual host-side extraction pass.
+
+## References
+
+- req: "Per PR #93: this affects every kernel that takes its inputs as
+  VmafCudaBuffer struct copies" (per user direction, 2026-05-29)
+- Research-0754 (`calculate_ssim_vert_combine` ncu A/B)
+- Research-0756 (this audit)
+- Research-0736 (SSIM ncu hotpath)
+- Research-0737 (MS-SSIM ncu hotpath)
+- Research-0734 (ADM ncu hotpath)
+- ADR-0108 (deep-dive deliverables)

docs/adr/README.md

@@ -767,3 +767,4 @@ ADRs may exist there for local session continuity, but the tracked
 | [ADR-0717](0717-vmafx-node-ffmpeg-latest.md) | vmafx-node ffmpeg version policy: pin to latest stable tag (n8.2); multi-stage Dockerfile with cpu/cuda/rocm/sycl variants | Accepted | node, ffmpeg, docker, phase4b, fork-local |
 | [ADR-0752](0752-perf-bench-multi-resolution.md) | Accepted | Multi-resolution performance benchmark baseline — `scripts/perf/bench-multi-resolution.sh` + versioned JSON in `testdata/perf_multi_resolution.json` |
 | [ADR-0754](0754-cuda-ssim-vert-combine-ldg-pinned-leak.md) | CUDA SSIM vert_combine: __ldg() read-only cache routing for 5x11 intermediate loads + __launch_bounds__(128) register budget hint + fix pinned-host memory leak in close_fex_cuda. | Accepted | 2026-05-29 | cuda, performance, correctness, ssim, fork-local |
+| [ADR-0756](0756-cuda-f3-struct-by-value-audit.md) | Fork-wide audit of CUDA kernels accepting VmafCudaBuffer/VmafPicture/AdmBufferCuda by value (F3 pattern); severity ranking by measured DRAM throughput; dispatch order for in-kernel __ldg() extraction follow-on PRs. | Accepted | 2026-05-29 | cuda, performance, research, fork-local |

docs/rebase-notes.md

@@ -40206,3 +40206,17 @@ All files modified are fork-local:
 
 No source files were modified (audit-only).  No Netflix upstream commit
 will collide with these additions on `sync-upstream`.
+## research/cuda-f3-struct-by-value-audit-20260529 (2026-05-29)
+
+No rebase impact: this PR adds documentation-only files (research digest, ADR,
+changelog fragment, state.md row). No CUDA source files are modified.
+`VmafCudaBuffer`, `VmafPicture`, and `AdmBufferCuda` definitions are
+unchanged; no upstream Netflix/vmaf commit will collide with this PR's diff.
+
+Fork-local files added/modified:
+`docs/research/research-0756-cuda-f3-struct-by-value-audit.md` (new),
+`docs/adr/0756-cuda-f3-struct-by-value-audit.md` (new),
+`docs/adr/README.md` (new row),
+`changelog.d/perf/cuda-f3-struct-by-value-audit.md` (new),
+`docs/state.md` (new row),
+`docs/rebase-notes.md` (this entry).