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

changelog.d/added/cuda-resolution-aware-dispatch.md

@@ -0,0 +1,11 @@
+## Added: Runtime resolution-aware CUDA kernel variant dispatch (ADR-0753)
+
+A new `vmaf_cuda_workload_class(w, h)` classifier (`core/src/feature/cuda/resolution_dispatch.{h,c}`)
+maps each frame's luma pixel count to `WS_SMALL` (< 720p), `WS_MEDIUM` (720p–4K), or
+`WS_LARGE` (>= 4K) at runtime. Feature extractors use this to pick the optimal kernel
+variant without requiring separate binaries.
+
+First consumer: `integer_adm_cuda.c::adm_cm_device()` selects
+`adm_cm_line_kernel_8` (with `__launch_bounds__(128,8)`) at WS_MEDIUM and
+`adm_cm_line_kernel_8_no_bounds` at WS_SMALL / WS_LARGE, recovering the
+−9.3% kernel-time saving at 1080p without regressing 576p or 4K.

core/src/feature/cuda/AGENTS.md

@@ -540,3 +540,55 @@ See [ADR-0747](../../../../docs/adr/0747-cuda-extern-c-sweep.md).
   sub-4K throughput regardless of kernel-level changes. The primary fix is
   multi-frame SAD batching (accumulate N frames before readback synchronization).
   See [Research-0760](../../../../docs/research/0760-cuda-motion-ncu-multi-resolution-20260529.md).
+## Resolution-aware kernel variant dispatch (ADR-0753)
+
+`resolution_dispatch.h` / `resolution_dispatch.c` in this directory provide a
+lightweight `vmaf_cuda_workload_class(w, h)` classifier used to pick between
+kernel variants at runtime. The current policy table is in ADR-0753.
+
+**How to add a new resolution-aware variant:**
+
+1. In the extractor's `.cu` file, define two kernel entry points using sibling
+   macros: one WITH `__launch_bounds__` (or the other occupancy hint), one
+   WITHOUT, giving the no-hint variant a `_no_bounds` suffix.
+   Both must be inside the `extern "C" { }` block (ADR-0747).
+2. In the extractor state struct (e.g. `AdmStateCuda`), add a second
+   `CUfunction` pointer for the no-hint variant.
+3. In `init_fex_cuda`, load both pointers via `cuModuleGetFunction`.
+   Add a short comment citing the ADR-0753 policy (see existing examples).
+4. At the kernel-launch site in `submit_fex_cuda`, call
+   `vmaf_cuda_workload_class(w, h)` and branch on the result.
+   The branch structure is always a single ternary / if-else — no nested
+   policy trees. Consult the policy table in ADR-0753 for which class gets
+   the bounded variant; the pattern so far:
+   - `adm_cm`: BOUNDED at `WS_MEDIUM` only; NO_BOUNDS at `WS_SMALL` + `WS_LARGE`.
+   - `filter1d` + `ssim_vert_combine`: BOUNDED at `WS_MEDIUM` + `WS_LARGE`;
+     NO_BOUNDS at `WS_SMALL` only.
+5. Add a row to the policy table in ADR-0753 `## Decision` and to the kernel
+   list in `resolution_dispatch.h`.
+6. Note the new invariant in this file under "Rebase-sensitive invariants".
+7. Update `docs/backends/cuda/overview.md` kernel table.
+
+**Verified wirings (as of the ADR-0753 extended scope):**
+
+| Feature | BOUNDED variant (kernel name) | NO_BOUNDS variant | Policy |
+|---|---|---|---|
+| `adm_cm` | `adm_cm_line_kernel_8` | `adm_cm_line_kernel_8_no_bounds` | MEDIUM only |
+| `filter1d` | `filter1d_8_horizontal_kernel_2_17_9` | `filter1d_8_horizontal_kernel_2_17_9_no_bounds` | MEDIUM + LARGE |
+| `ssim_vert_combine` | `calculate_ssim_vert_combine` | `calculate_ssim_vert_combine_no_bounds` | MEDIUM + LARGE |
+
+- **`integer_vif_cuda.c::filter1d_8` picks `filter1d_8_horizontal_kernel_2_17_9_no_bounds`
+  at `WS_SMALL` and the bounded variant at `WS_MEDIUM`/`WS_LARGE`** (ADR-0753 extended
+  scope). `VifStateCuda` carries `func_filter1d_8_horizontal_kernel_2_17_9_no_bounds`.
+  `filter1d.cu` defines `FILTER1D_8_HORI_NO_BOUNDS(2, 17, 9)` inside `extern "C" {}`.
+  On rebase: verify both `cuModuleGetFunction` calls in `init_fex_cuda` reference
+  valid symbols. If upstream refactors the macro or adds new `fwidth` variants,
+  apply the `_NO_BOUNDS` sibling macro around the new body too.
+
+- **`integer_ssim_cuda.c::submit_fex_cuda` picks `calculate_ssim_vert_combine_no_bounds`
+  at `WS_SMALL` and the bounded variant at `WS_MEDIUM`/`WS_LARGE`** (ADR-0753 extended
+  scope). `SsimStateCuda` carries `func_vert_no_bounds`.
+  `integer_ssim/ssim_score.cu` defines both variants inside `extern "C" {}`.
+  On rebase: if upstream modifies `calculate_ssim_vert_combine`, apply the same
+  diff to `calculate_ssim_vert_combine_no_bounds` (body is identical; only the
+  `__launch_bounds__(128)` annotation differs).

core/src/feature/cuda/integer_adm/adm_cm.cu

@@ -361,8 +361,42 @@ adm_cm_line_kernel(AdmBufferCuda buf, int h, int w, int top, int bottom, int lef
  * Scale 0 already used the fused pattern (adm_cm_line_kernel_8); scales 1-3
  * were migrated in PR perf/adm-cm-cuda-warp-reduce-fusion. */
 
-#define ADM_CM_LINE(rows_per_thread)                                                               \
-    __global__ void adm_cm_line_kernel_##rows_per_thread(                                          \
+/* ADR-0753: Two variants of adm_cm_line_kernel_8 for resolution-aware dispatch.
+ *
+ *   adm_cm_line_kernel_8            — WITH __launch_bounds__(128, 8)
+ *     Reduces register count on sm_89 from 114 to the bounds-guided budget,
+ *     improving theoretical occupancy in the 8–32 wave regime (WS_MEDIUM: 1080p).
+ *     Measured −9.3% kernel time at 1080p (Research-0749). Neutral at 576p / 4K.
+ *
+ *   adm_cm_line_kernel_8_no_bounds  — WITHOUT __launch_bounds__
+ *     Compiler allocates the full register budget (114 regs, sm_89).
+ *     Used at WS_SMALL (< 720p) and WS_LARGE (>= 4K) where the bounds hint
+ *     shows no measurable gain and the extra register spilling from the hint
+ *     could hurt future kernels on lower-SM GPUs.
+ *
+ * The dispatch site in integer_adm_cuda.c::adm_cm_device() picks the variant
+ * via vmaf_cuda_workload_class(w, h) — a single integer compare per frame.
+ * See resolution_dispatch.h and ADR-0753.
+ */
+
+/* Macro for the bounds-guided variant (WS_MEDIUM). */
+#define ADM_CM_LINE_BOUNDED(rows_per_thread)                                                       \
+    __launch_bounds__(128, 8) __global__ void adm_cm_line_kernel_##rows_per_thread(                \
+        AdmBufferCuda buf, int h, int w, int top, int bottom, int left, int right, int start_row,  \
+        int end_row, int start_col, int end_col, int src_stride, int csf_a_stride, int buffer_h,   \
+        int buffer_stride, int32_t *accum_per_block, AdmFixedParametersCuda params, int scale,     \
+        int64_t *accum_global, WarpShift ws, const uint32_t shift_inner_accum,                     \
+        const uint32_t add_shift_inner_accum)                                                      \
+    {                                                                                              \
+        adm_cm_line_kernel<rows_per_thread>(                                                       \
+            buf, h, w, top, bottom, left, right, start_row, end_row, start_col, end_col,           \
+            src_stride, csf_a_stride, buffer_h, buffer_stride, accum_per_block, params, scale,     \
+            accum_global, ws, shift_inner_accum, add_shift_inner_accum);                           \
+    }
+
+/* Macro for the no-bounds variant (WS_SMALL, WS_LARGE). */
+#define ADM_CM_LINE_NO_BOUNDS(rows_per_thread)                                                     \
+    __global__ void adm_cm_line_kernel_##rows_per_thread##_no_bounds(                              \
         AdmBufferCuda buf, int h, int w, int top, int bottom, int left, int right, int start_row,  \
         int end_row, int start_col, int end_col, int src_stride, int csf_a_stride, int buffer_h,   \
         int buffer_stride, int32_t *accum_per_block, AdmFixedParametersCuda params, int scale,     \
@@ -376,9 +410,10 @@ adm_cm_line_kernel(AdmBufferCuda buf, int h, int w, int top, int bottom, int lef
     }
 
 extern "C" {
-// 128 = warps_per_thread * val_per_thread = 32 * 4 -- assuming 32 threads per warp, this might change in the future
+// 128 = warps_per_thread * val_per_thread = 32 * 4 -- assuming 32 threads per warp
 /* adm_cm_reduce_line_kernel_4 removed: fused into i4_adm_cm_line_kernel_fused (scales 1-3). */
-ADM_CM_LINE(8); // adm_cm_line_kernel_8
+ADM_CM_LINE_BOUNDED(8);   /* adm_cm_line_kernel_8          — WS_MEDIUM (1080p)    */
+ADM_CM_LINE_NO_BOUNDS(8); /* adm_cm_line_kernel_8_no_bounds — WS_SMALL / WS_LARGE */
 }
 
 /* ============================================================================

core/src/feature/cuda/integer_adm_cuda.c

@@ -34,6 +34,7 @@
  * feature/integer_adm.h. No separate adm_options.h include is needed here. */
 #include "drain_batch.h"
 #include "picture_cuda.h"
+#include "feature/cuda/resolution_dispatch.h"
 
 #include <assert.h>
 
@@ -83,7 +84,9 @@ typedef struct AdmStateCuda {
         func_adm_csf_den_scale_line_kernel, func_adm_csf_den_s123_line_kernel,
         // adm_cm kernel
         /* func_adm_cm_reduce_line_kernel_4 removed: fused into i4_adm_cm_line_kernel_fused. */
-        func_adm_cm_line_kernel_8, func_i4_adm_cm_line_kernel_fused,
+        func_adm_cm_line_kernel_8,           /* with __launch_bounds__(128,8): WS_MEDIUM   */
+        func_adm_cm_line_kernel_8_no_bounds, /* without bounds hint: WS_SMALL / WS_LARGE   */
+        func_i4_adm_cm_line_kernel_fused,
         /* AIM CM kernels (ADR-0746): */
         func_adm_cm_aim_line_kernel_8, func_i4_adm_cm_aim_line_kernel_fused;
 
@@ -495,8 +498,13 @@ int adm_cm_device(AdmStateCuda *s, AdmBufferCuda *buf, int w, int h, int src_str
                         &shift_inner_accum,
                         &add_shift_inner_accum};
 
-        CHECK_CUDA_RETURN(cu_f, cuLaunchKernel(s->func_adm_cm_line_kernel_8,
-                                               DIV_ROUND_UP(buffer_stride, BLOCKX),
+        /* ADR-0753: pick __launch_bounds__ variant only at WS_MEDIUM (1080p range).
+         * At WS_SMALL the kernel is launch-overhead-bound; at WS_LARGE the register
+         * pressure regime changes and the bounds hint shows zero gain (Research-0751). */
+        const enum WorkloadSize adm_cm_ws = vmaf_cuda_workload_class(w, h);
+        CUfunction adm_cm_fn = (adm_cm_ws == WS_MEDIUM) ? s->func_adm_cm_line_kernel_8 :
+                                                          s->func_adm_cm_line_kernel_8_no_bounds;
+        CHECK_CUDA_RETURN(cu_f, cuLaunchKernel(adm_cm_fn, DIV_ROUND_UP(buffer_stride, BLOCKX),
                                                DIV_ROUND_UP(buffer_h, BLOCKY * rows_per_thread), 3,
                                                BLOCKX, BLOCKY, 1, 0, c_stream, args, NULL));
     }
@@ -1373,10 +1381,19 @@ static int init_fex_cuda(VmafFeatureExtractor *fex, enum VmafPixelFormat pix_fmt
                     fail);
 
     /* adm_cm_reduce_line_kernel_4 removed: fused into i4_adm_cm_line_kernel_fused. */
+    /* ADR-0753: load both resolution variants of adm_cm_line_kernel_8.
+     * The bounded variant (WS_MEDIUM) carries __launch_bounds__(128,8) which
+     * saves ~9.3% kernel time at 1080p by reducing register pressure.
+     * The no-bounds variant (WS_SMALL / WS_LARGE) lets the compiler allocate
+     * the full register budget where the occupancy win does not materialise. */
     CHECK_CUDA_GOTO(
         cu_f,
         cuModuleGetFunction(&s->func_adm_cm_line_kernel_8, adm_cm_module, "adm_cm_line_kernel_8"),
         fail);
+    CHECK_CUDA_GOTO(cu_f,
+                    cuModuleGetFunction(&s->func_adm_cm_line_kernel_8_no_bounds, adm_cm_module,
+                                        "adm_cm_line_kernel_8_no_bounds"),
+                    fail);
     CHECK_CUDA_GOTO(cu_f,
                     cuModuleGetFunction(&s->func_i4_adm_cm_line_kernel_fused, adm_cm_module,
                                         "i4_adm_cm_line_kernel_fused"),

core/src/feature/cuda/integer_ssim/ssim_score.cu

@@ -213,4 +213,66 @@ __launch_bounds__(128) __global__
     }
 }
 
+/* ADR-0753: _no_bounds sibling of calculate_ssim_vert_combine.
+ * Omits __launch_bounds__(128) for WS_SMALL (<720p) where occupancy pressure
+ * is absent (the workload is wave-limited, not register-limited) and at
+ * WS_LARGE (>=4K) where the kernel is fully saturated and the bounds hint
+ * adds no measurable gain.  The body is identical to calculate_ssim_vert_combine;
+ * only the launch-bounds annotation is stripped.
+ * __ldg() loads are retained: they are correct at all resolutions (the 5 horiz
+ * intermediates are always write-once / read-many). */
+__global__ void calculate_ssim_vert_combine_no_bounds(
+    VmafCudaBuffer h_ref_mu_buf, VmafCudaBuffer h_cmp_mu_buf, VmafCudaBuffer h_ref_sq_buf,
+    VmafCudaBuffer h_cmp_sq_buf, VmafCudaBuffer h_refcmp_buf, VmafCudaBuffer partials,
+    unsigned w_horiz, unsigned w_final, unsigned h_final, float c1, float c2)
+{
+    const unsigned x = blockIdx.x * blockDim.x + threadIdx.x;
+    const unsigned y = blockIdx.y * blockDim.y + threadIdx.y;
+    const float *__restrict__ h_ref_mu = reinterpret_cast<const float *>(h_ref_mu_buf.data);
+    const float *__restrict__ h_cmp_mu = reinterpret_cast<const float *>(h_cmp_mu_buf.data);
+    const float *__restrict__ h_ref_sq = reinterpret_cast<const float *>(h_ref_sq_buf.data);
+    const float *__restrict__ h_cmp_sq = reinterpret_cast<const float *>(h_cmp_sq_buf.data);
+    const float *__restrict__ h_refcmp = reinterpret_cast<const float *>(h_refcmp_buf.data);
+
+    float my_ssim = 0.0f;
+    if (x < w_final && y < h_final) {
+        float ref_mu = 0.0f, cmp_mu = 0.0f, ref_sq = 0.0f, cmp_sq = 0.0f, refcmp = 0.0f;
+        for (int v = 0; v < K; v++) {
+            const unsigned src_y = y + (unsigned)v;
+            const unsigned src_idx = src_y * w_horiz + x;
+            const float w = G[v];
+            ref_mu += w * __ldg(&h_ref_mu[src_idx]);
+            cmp_mu += w * __ldg(&h_cmp_mu[src_idx]);
+            ref_sq += w * __ldg(&h_ref_sq[src_idx]);
+            cmp_sq += w * __ldg(&h_cmp_sq[src_idx]);
+            refcmp += w * __ldg(&h_refcmp[src_idx]);
+        }
+        const float ref_var = ref_sq - ref_mu * ref_mu;
+        const float cmp_var = cmp_sq - cmp_mu * cmp_mu;
+        const float covar = refcmp - ref_mu * cmp_mu;
+        const float mu_xy = ref_mu * cmp_mu;
+        const float num = (2.0f * mu_xy + c1) * (2.0f * covar + c2);
+        const float den = (ref_mu * ref_mu + cmp_mu * cmp_mu + c1) * (ref_var + cmp_var + c2);
+        my_ssim = num / den;
+    }
+
+    __shared__ float s_warp_sums[BLOCK_SIZE / 32];
+    float warp_sum = my_ssim;
+    for (int off = 16; off > 0; off >>= 1)
+        warp_sum += __shfl_down_sync(0xffffffff, warp_sum, off);
+    const int tid = threadIdx.y * blockDim.x + threadIdx.x;
+    const int lane = tid % 32;
+    const int warp_id = tid / 32;
+    if (lane == 0)
+        s_warp_sums[warp_id] = warp_sum;
+    __syncthreads();
+    if (tid == 0) {
+        float block_sum = 0.0f;
+        for (int i = 0; i < BLOCK_SIZE / 32; i++)
+            block_sum += s_warp_sums[i];
+        const unsigned block_idx = blockIdx.y * gridDim.x + blockIdx.x;
+        reinterpret_cast<float *>(partials.data)[block_idx] = block_sum;
+    }
+}
+
 } /* extern "C" */

core/src/feature/cuda/integer_ssim_cuda.c

@@ -42,6 +42,7 @@
 #include "picture.h"
 #include "picture_cuda.h"
 #include "cuda_helper.cuh"
+#include "feature/cuda/resolution_dispatch.h"
 
 #define SSIM_BLOCK_X 16
 #define SSIM_BLOCK_Y 8
@@ -57,7 +58,8 @@ typedef struct SsimStateCuda {
 
     CUfunction func_horiz_8;
     CUfunction func_horiz_16;
-    CUfunction func_vert;
+    CUfunction func_vert;           /* with __launch_bounds__(128): WS_MEDIUM + WS_LARGE */
+    CUfunction func_vert_no_bounds; /* without bounds hint: WS_SMALL                    */
     int scale_override;
     /* `enable_chroma` option: when false, only luma is dispatched.
      * Default false mirrors CPU integer_ssim.c PR #939. */
@@ -175,6 +177,11 @@ static int init_fex_cuda(VmafFeatureExtractor *fex, enum VmafPixelFormat pix_fmt
         cu_f, cuModuleGetFunction(&s->func_horiz_16, module, "calculate_ssim_horiz_16bpc"), fail);
     CHECK_CUDA_GOTO(cu_f, cuModuleGetFunction(&s->func_vert, module, "calculate_ssim_vert_combine"),
                     fail);
+    /* ADR-0753: no-bounds sibling for WS_SMALL (<720p) where occupancy pressure is absent. */
+    CHECK_CUDA_GOTO(cu_f,
+                    cuModuleGetFunction(&s->func_vert_no_bounds, module,
+                                        "calculate_ssim_vert_combine_no_bounds"),
+                    fail);
 
     CHECK_CUDA_GOTO(cu_f, cuCtxPopCurrent(NULL), fail_after_pop);
 
@@ -303,7 +310,12 @@ static int submit_fex_cuda(VmafFeatureExtractor *fex, VmafPicture *ref_pic, Vmaf
     /* Pass 2 — vertical + SSIM combine. Grid sized over
      * (W-10) × (H-10). The horiz pass writes happen-before
      * the vert pass reads on the same stream — implicit
-     * stream ordering, no extra event needed. */
+     * stream ordering, no extra event needed.
+     *
+     * ADR-0753: resolution-aware dispatch. BOUNDED variant (__launch_bounds__(128))
+     * at WS_MEDIUM + WS_LARGE (>=720p); NO_BOUNDS at WS_SMALL (<720p).
+     * __ldg() loads are present in BOTH variants (correct and beneficial at all
+     * resolutions per ADR-0754). */
     void *params2[] = {
         (void *)s->h_ref_mu,
         (void *)s->h_cmp_mu,
@@ -317,8 +329,14 @@ static int submit_fex_cuda(VmafFeatureExtractor *fex, VmafPicture *ref_pic, Vmaf
         &s->c1,
         &s->c2,
     };
-    CHECK_CUDA_RETURN(cu_f, cuLaunchKernel(s->func_vert, grid_x, grid_y, 1, SSIM_BLOCK_X,
-                                           SSIM_BLOCK_Y, 1, 0, stream, params2, NULL));
+    {
+        const enum WorkloadSize ssim_vert_ws =
+            vmaf_cuda_workload_class((int)s->width, (int)s->height);
+        CUfunction ssim_vert_fn =
+            (ssim_vert_ws == WS_SMALL) ? s->func_vert_no_bounds : s->func_vert;
+        CHECK_CUDA_RETURN(cu_f, cuLaunchKernel(ssim_vert_fn, grid_x, grid_y, 1, SSIM_BLOCK_X,
+                                               SSIM_BLOCK_Y, 1, 0, stream, params2, NULL));
+    }
 
     /* DtoH copy of the partials on our private stream. */
     CHECK_CUDA_RETURN(cu_f, cuEventRecord(s->lc.submit, stream));