refactor(cuda): bring the CUDA kernels and their headers to the lint and HISS standard (ADR-1142)

CHANGELOG.md

@@ -2796,6 +2796,17 @@ make `core/AGENTS.md` a generated index over `AGENTS.d/` topic pages ([ADR-1454]
   `isnumeric()` no longer reads a `string_view` as a C string. No score or API change.
 
 
+- **The CUDA kernels, their headers and two CUDA tests meet the clang-tidy
+  standard.** The 62 files of the `cuda` lane that were not shared C headers
+  owned by the `cpu` lane are cleaned: device pointers are rebuilt from a
+  kernel argument's `CUdeviceptr` with one macro (`cuda_device_ptr.cuh`),
+  helpers live in anonymous namespaces, the long VIF statistic and the
+  SpEED / CAMBI / ADM / SSIM / PSNR-HVS kernels are split into helpers, and
+  the shared device headers carry the ADR-1138 brackets for the C includers.
+  No score moves: the CUDA twins return the same bits on the Netflix pair, both
+  1080p checkerboard pairs and a 10-bit pair, and every CUDA device test passes.
+
+
 - The HIP kernels of integer VIF, float VIF, PSNR-HVS and SSIMULACRA 2, and the
   shared GPU headers `ordered_sum.h`, `adm_angle_flag.h`, `ff_math.h`,
   `ciede_ff_math.h`, `adm_cm_accumulator.h`, `integer_adm.h`,
@@ -3265,6 +3276,14 @@ make `core/AGENTS.md` a generated index over `AGENTS.d/` topic pages ([ADR-1454]
   unchanged.
 
 
+- `vmaf` prints `problem scoring picture N: libvmaf returned E` when libvmaf
+  fails to score a frame it has read (a feature extractor refused the frame or
+  its options), where it printed `problem reading pictures`, which read like
+  the input read failure `problem while reading pictures` (exit 102). The exit
+  status is unchanged. `docs/usage/cli.md` lists the case in the exit-code
+  table.
+
+
 - **CodeQL include-non-header alert #1309 resolved with internal test accessors and CI guard.**
   `core/test/test_feature_backend_twin.c` linked directly against `libvmaf` instead
   of unity-including `core/src/libvmaf.c`. Narrow internal accessors
@@ -4583,6 +4602,14 @@ make `core/AGENTS.md` a generated index over `AGENTS.d/` topic pages ([ADR-1454]
   bisect cap tests stub the scoring step the way their docstring says.
 
 
+- `list_backends`, `probe_backend` and `vmaf_version` of both MCP servers (Python
+  and Go) now take a GPU backend as available only when `vmaf --list-backends`
+  reports it usable. They read `vmaf --help` before, which names every backend
+  on every build, so a CPU-only `vmaf` was reported with CUDA, SYCL, HIP and
+  Metal and the backend allowlist admitted them. A `vmaf` that cannot print the
+  report is treated as CPU-only.
+
+
 - **The Go MCP server is checked against the Python server's own tool
   list.** Its parity tests compared it with 15 hand-copied tool names, so the
   four sidecar tools and every property type went unchecked. The Python server
@@ -5961,6 +5988,14 @@ make `core/AGENTS.md` a generated index over `AGENTS.d/` topic pages ([ADR-1454]
   change by up to 3.6e-7.
 
 
+- **`vif_sycl` with `vif_fused=true` returns the CPU's scores.** From 1920x1080
+  up, scales 1 to 3 differed from the CPU `vif` on every frame (by up to 4.9e-4
+  at 3840x2160): one fused launch read a scale from the downsampled buffers it
+  was writing the next scale into. The fused scales now alternate between two
+  buffers, which costs 4 MB of device memory at 3840x2160; the default separate
+  passes were not affected.
+
+
 - **SYCL: no kernel uses scratch memory on Xe2 (Arc B580, Arc Pro B60).** On an Arc B580
   the term kernel of `float_adm_sycl` spilled 128 bytes, and
   `test_sycl_kernel_scratch` failed; its scores were still exact there, but a
@@ -6261,6 +6296,12 @@ make `core/AGENTS.md` a generated index over `AGENTS.d/` topic pages ([ADR-1454]
   working directory, so the suite no longer leaves a `-version` file behind.
 
 
+- The vmaf-tune Python tests no longer start the host's `vmaf` through the
+  backend probe (a suite-wide fixture keeps the default probe off `PATH`), and
+  `go test ./pkg/fast/` runs the vmaf CLI of the build under test
+  (`VMAF_BIN` or `core/build-cpu`) instead of `vmaf` on `PATH`.
+
+
 - **`vmaf-tune corpus --two-pass` encodes with `libx265`.** x265 refuses
   `-crf` in the second pass (exit 183), so every libx265 two-pass cell failed.
   A cell at a CRF now runs pass 1 at that CRF, measures the bitstream with

changelog.d/changed/tidy-zero-cuda-lane.md

@@ -0,0 +1,9 @@
+- **The CUDA kernels, their headers and two CUDA tests meet the clang-tidy
+  standard.** The 62 files of the `cuda` lane that were not shared C headers
+  owned by the `cpu` lane are cleaned: device pointers are rebuilt from a
+  kernel argument's `CUdeviceptr` with one macro (`cuda_device_ptr.cuh`),
+  helpers live in anonymous namespaces, the long VIF statistic and the
+  SpEED / CAMBI / ADM / SSIM / PSNR-HVS kernels are split into helpers, and
+  the shared device headers carry the ADR-1138 brackets for the C includers.
+  No score moves: the CUDA twins return the same bits on the Netflix pair, both
+  1080p checkerboard pairs and a 10-bit pair, and every CUDA device test passes.

changelog.d/fixed/cli-scoring-error-message.md

@@ -0,0 +1,6 @@
+- `vmaf` prints `problem scoring picture N: libvmaf returned E` when libvmaf
+  fails to score a frame it has read (a feature extractor refused the frame or
+  its options), where it printed `problem reading pictures`, which read like
+  the input read failure `problem while reading pictures` (exit 102). The exit
+  status is unchanged. `docs/usage/cli.md` lists the case in the exit-code
+  table.

changelog.d/fixed/mcp-backends-from-list-backends.md

@@ -0,0 +1,6 @@
+- `list_backends`, `probe_backend` and `vmaf_version` of both MCP servers (Python
+  and Go) now take a GPU backend as available only when `vmaf --list-backends`
+  reports it usable. They read `vmaf --help` before, which names every backend
+  on every build, so a CPU-only `vmaf` was reported with CUDA, SYCL, HIP and
+  Metal and the backend allowlist admitted them. A `vmaf` that cannot print the
+  report is treated as CPU-only.

changelog.d/fixed/sycl-vif-fused-rd-race.md

@@ -0,0 +1,6 @@
+- **`vif_sycl` with `vif_fused=true` returns the CPU's scores.** From 1920x1080
+  up, scales 1 to 3 differed from the CPU `vif` on every frame (by up to 4.9e-4
+  at 3840x2160): one fused launch read a scale from the downsampled buffers it
+  was writing the next scale into. The fused scales now alternate between two
+  buffers, which costs 4 MB of device memory at 3840x2160; the default separate
+  passes were not affected.

changelog.d/fixed/vmaf-tune-tests-no-path-vmaf.md

@@ -0,0 +1,4 @@
+- The vmaf-tune Python tests no longer start the host's `vmaf` through the
+  backend probe (a suite-wide fixture keeps the default probe off `PATH`), and
+  `go test ./pkg/fast/` runs the vmaf CLI of the build under test
+  (`VMAF_BIN` or `core/build-cpu`) instead of `vmaf` on `PATH`.

cmd/vmafx-mcp/impl.go

@@ -29,6 +29,7 @@ import (
 
 	"github.com/VMAFx/vmafx/pkg/libvmaf"
 	"github.com/VMAFx/vmafx/pkg/modeleval"
+	"github.com/VMAFx/vmafx/pkg/scorebackend"
 )
 
 // ---------------------------------------------------------------------------
@@ -52,8 +53,11 @@ var (
 	probeCache = map[string]map[string]bool{}
 )
 
-// probeBackends returns the set of backends the vmaf binary advertises
-// (via --help flags). "cpu" is always included.
+// probeBackends returns the set of backends the vmaf binary can use on this
+// host. It reads `vmaf --list-backends` (ADR-1874) through pkg/scorebackend:
+// the help text names every backend on every build, so it says nothing about
+// the binary. "cpu" is always included; a binary that cannot print the report
+// is treated as CPU-only (scorebackend.Detect writes the reason to stderr).
 func probeBackends(vmafBin string) map[string]bool {
 	probeMu.Lock()
 	defer probeMu.Unlock()
@@ -63,22 +67,20 @@ func probeBackends(vmafBin string) map[string]bool {
 	}
 
 	advertised := map[string]bool{"cpu": true}
-	ctx, cancel := context.WithTimeout(context.Background(), 5*time.Second)
+	ctx, cancel := context.WithTimeout(context.Background(), backendProbeTimeout)
 	defer cancel()
 
-	out, err := exec.CommandContext(ctx, vmafBin, "--help").CombinedOutput()
-	if err == nil {
-		blob := string(out)
-		for _, name := range []string{"cuda", "sycl", "hip", "metal"} {
-			if strings.Contains(blob, "--no_"+name) {
-				advertised[name] = true
-			}
-		}
+	for _, name := range scorebackend.Detect(ctx, scorebackend.Options{VMAFBin: vmafBin}) {
+		advertised[name] = true
 	}
 	probeCache[vmafBin] = advertised
 	return advertised
 }
 
+// backendProbeTimeout bounds one `vmaf --list-backends` run, which initialises
+// every compiled GPU backend once (scorebackend allows 60 s for the same run).
+const backendProbeTimeout = 60 * time.Second
+
 // ---------------------------------------------------------------------------
 // String arg helpers
 // ---------------------------------------------------------------------------

cmd/vmafx-mcp/impl_handlers_test.go

@@ -680,35 +680,76 @@ func TestAddRawTool_HandlerErrorBecomesIsError(t *testing.T) {
 }
 
 // ---------------------------------------------------------------------------
-// probeBackends — advertised backends from --help output.
-// We write a tiny fake vmaf shell script that emits --no_cuda to verify
-// parsing.
-// ---------------------------------------------------------------------------
-
-func TestProbeBackends_ParsesHelpOutput(t *testing.T) {
-	tmp := t.TempDir()
-	script := filepath.Join(tmp, "vmaf")
-	// Script prints a --help output that advertises cuda.
-	if err := os.WriteFile(script, []byte("#!/bin/sh\necho '--no_cuda --no_sycl'\n"), 0o700); err != nil {
+// probeBackends — usable backends from `vmaf --list-backends` (ADR-1874).
+// A tiny fake vmaf shell script prints the report; the --help text of the
+// real CLI names every backend on every build, so it must not decide.
+// ---------------------------------------------------------------------------
+
+// fakeVmafReport writes a fake vmaf that prints report for --list-backends and
+// a help text naming every --no_<backend> flag for anything else.
+func fakeVmafReport(t *testing.T, report string) string {
+	t.Helper()
+	script := filepath.Join(t.TempDir(), "vmaf")
+	body := "#!/bin/sh\n" +
+		"if [ \"$1\" = \"--list-backends\" ]; then\ncat <<'EOF'\n" + report + "\nEOF\nexit 0\nfi\n" +
+		"echo '--no_cuda --no_sycl --no_hip --no_metal'\n"
+	if err := os.WriteFile(script, []byte(body), 0o700); err != nil {
 		t.Fatal(err)
 	}
-
-	// Ensure this binary isn't in the cache from a previous test run.
 	probeMu.Lock()
 	delete(probeCache, script)
 	probeMu.Unlock()
+	return script
+}
+
+func TestProbeBackends_ReadsListBackendsReport(t *testing.T) {
+	script := fakeVmafReport(t, `{"backends": [
+  {"name": "cpu", "compiled": true, "usable": true},
+  {"name": "cuda", "compiled": true, "usable": true},
+  {"name": "sycl", "compiled": false, "usable": false},
+  {"name": "hip", "compiled": true, "usable": false, "init_status": -19},
+  {"name": "metal", "compiled": false, "usable": false}]}`)
 
 	advertised := probeBackends(script)
 	if !advertised["cpu"] {
 		t.Error("cpu always true")
 	}
 	if !advertised["cuda"] {
-		t.Error("cuda should be true: script emitted '--no_cuda'")
+		t.Error("cuda should be true: the report lists it usable")
 	}
-	if !advertised["sycl"] {
-		t.Error("sycl should be true: script emitted '--no_sycl'")
+	for _, name := range []string{"sycl", "hip", "metal"} {
+		if advertised[name] {
+			t.Errorf("%s should be false: the report does not list it usable "+
+				"(the help text names it on every build)", name)
+		}
 	}
-	if advertised["hip"] {
-		t.Error("hip should be false: script did not emit '--no_hip'")
+}
+
+func TestProbeBackends_CPUOnlyBuildReportsNoGPU(t *testing.T) {
+	script := fakeVmafReport(t, `{"backends": [
+  {"name": "cpu", "compiled": true, "usable": true},
+  {"name": "cuda", "compiled": false, "usable": false},
+  {"name": "sycl", "compiled": false, "usable": false},
+  {"name": "hip", "compiled": false, "usable": false},
+  {"name": "metal", "compiled": false, "usable": false}]}`)
+
+	advertised := probeBackends(script)
+	if len(advertised) != 1 || !advertised["cpu"] {
+		t.Errorf("a CPU-only build advertises only cpu, got %v", advertised)
+	}
+}
+
+func TestProbeBackends_NoReportMeansCPUOnly(t *testing.T) {
+	script := filepath.Join(t.TempDir(), "vmaf")
+	if err := os.WriteFile(script, []byte("#!/bin/sh\nexit 2\n"), 0o700); err != nil {
+		t.Fatal(err)
+	}
+	probeMu.Lock()
+	delete(probeCache, script)
+	probeMu.Unlock()
+
+	advertised := probeBackends(script)
+	if len(advertised) != 1 || !advertised["cpu"] {
+		t.Errorf("a binary without --list-backends is CPU-only, got %v", advertised)
 	}
 }

core/include/libvmaf/libvmaf_cuda.h

@@ -41,7 +41,9 @@ extern "C" {
  * `VmafMetalState`. One state pins one device; callers that want multi-GPU
  * fan-out create one state per device and one VmafContext per state.
  */
+/* NOLINTBEGIN(modernize-use-using): C header included by C and C++ translation units; C has no `using`. ADR-1138. */
 typedef struct VmafCudaState VmafCudaState;
+/* NOLINTEND(modernize-use-using) */
 
 /**
  * @struct VmafCudaConfiguration
@@ -55,9 +57,11 @@ typedef struct VmafCudaState VmafCudaState;
  * VmafContext until `vmaf_close()` returns exactly 0. A nonzero close retains
  * that dependency for retry.
  */
+/* NOLINTBEGIN(modernize-use-using): C header included by C and C++ translation units; C has no `using`. ADR-1138. */
 typedef struct VmafCudaConfiguration {
     void *cu_ctx; /**< Optional CUcontext (cast from `CUcontext`); NULL → create one. */
 } VmafCudaConfiguration;
+/* NOLINTEND(modernize-use-using) */
 
 /**
  * Initialize VmafCudaState.
@@ -152,12 +156,14 @@ VMAF_EXPORT int vmaf_cuda_import_state(VmafContext *vmaf, VmafCudaState *cu_stat
  *
  * Stable enumerator values — append-only across libvmaf releases.
  */
+/* NOLINTBEGIN(performance-enum-size): C header included by C and C++ translation units; C has no fixed enum underlying type across the required toolchains (ADR-1470). ADR-1138. */
 enum VmafCudaPicturePreallocationMethod {
     VMAF_CUDA_PICTURE_PREALLOCATION_METHOD_NONE = 0,
     VMAF_CUDA_PICTURE_PREALLOCATION_METHOD_DEVICE,
     VMAF_CUDA_PICTURE_PREALLOCATION_METHOD_HOST,
     VMAF_CUDA_PICTURE_PREALLOCATION_METHOD_HOST_PINNED,
 };
+/* NOLINTEND(performance-enum-size) */
 
 /**
  * @struct VmafCudaPictureConfiguration
@@ -172,6 +178,7 @@ enum VmafCudaPicturePreallocationMethod {
  * Storage tier is selected by `pic_prealloc_method` (see
  * `VmafCudaPicturePreallocationMethod`).
  */
+/* NOLINTBEGIN(modernize-use-using): C header included by C and C++ translation units; C has no `using`. ADR-1138. */
 typedef struct VmafCudaPictureConfiguration {
     struct {
         unsigned w;                   /**< Per-plane width in samples. */
@@ -181,6 +188,7 @@ typedef struct VmafCudaPictureConfiguration {
     } pic_params;                     /**< Per-picture shape (width/height/bpc/pixel-format). */
     enum VmafCudaPicturePreallocationMethod pic_prealloc_method; /**< Storage tier selector. */
 } VmafCudaPictureConfiguration;
+/* NOLINTEND(modernize-use-using) */
 
 /**
  * Config and preallocate VmafPictures for use during CUDA feature extraction.

core/src/cuda/AGENTS.md

@@ -416,3 +416,21 @@ whether context pop still has to run; `pinned_alloc_unwind` takes `PINNED_UNWIND
 because it replaced `free_priv` -> `free_data` -> `fail_no_data` cascade. Adding resource to
 either path means adding stage to helper, not second exit. `CHECK_CUDA_GOTO` labels stay
 — they are macro's jump targets and now call helper.
+
+## Device pointers and device-code lint (ADR-1142, RC3 tidy-zero)
+
+- Kernel arg structs carry `CUdeviceptr` as `uint64_t`. Kernel turns one back into pointer with
+  `VMAF_CUDA_DPTR(T, address)` (`cuda_device_ptr.cuh`), never `reinterpret_cast` on integer. One
+  macro, one NOLINT(performance-no-int-to-ptr).
+- Macro, not function: function wrapper (`reinterpret_cast` or `__builtin_bit_cast`) turned 134 of
+  `speed_score.cu`'s 378 `LDG.E.CONSTANT` into plain loads (sm_89, nvcc 13.4). Do not turn it into
+  an inline function.
+- No designated initializers (`{.a = 1}`) in `.cu` / `.cuh`: nvcc's MSVC host frontend rejects them
+  (`preflight.sh --stage msvcism`). Fill aggregate field by field (`make_*()` helper). No
+  constructor: it trips `misc-non-private-member-variables-in-classes`.
+- Helpers and device structs of a `.cu` live in `namespace {}`, no `static`. `__global__` kernels
+  and symbols the host looks up by name stay `extern "C"` / external.
+- Header included by C host (`typedef`, plain `enum`): `NOLINTBEGIN(modernize-use-using ...)`
+  with ADR-1138 / ADR-1470 wording, like `core/include/libvmaf/model.h`.
+- Refactor proof: sm_89 SASS before / after (`cuobjdump -sass`) identical, or the diff explained
+  (commutative operand swap, register names).

core/src/cuda/cuda_device_ptr.cuh

@@ -0,0 +1,35 @@
+/**
+ *
+ *  Copyright 2026 Lusoris
+ *  SPDX-License-Identifier: EUPL-1.2
+ *
+ *  Device pointer from the 64-bit address a kernel argument struct carries.
+ *
+ *  The host fills the argument structs of the CUDA feature kernels with
+ *  `CUdeviceptr` values (unsigned 64-bit integers), because the structs are
+ *  plain C. A kernel turns each address back into a typed pointer once, at
+ *  its top. clang-tidy's performance-no-int-to-ptr rejects that conversion
+ *  (an integer-to-pointer cast hides the pointer's provenance from the
+ *  optimiser); there is no integer-free way to spell it, so the conversion
+ *  lives in this one macro, with its one lint suppression.
+ *
+ *  It is a macro and not a function on purpose. Measured on sm_89 with nvcc
+ *  13.4 on core/src/feature/cuda/speed/speed_score.cu: the same conversion
+ *  inside a `__forceinline__` function template (reinterpret_cast or
+ *  __builtin_bit_cast alike) turns 134 of the file's 378 LDG.E.CONSTANT
+ *  (ld.global.nc) loads into plain loads, which gives up the compiler's
+ *  freedom to treat those loads as invariant. The macro keeps the SASS of the
+ *  direct cast.
+ */
+
+#ifndef VMAF_SRC_CUDA_CUDA_DEVICE_PTR_CUH_
+#define VMAF_SRC_CUDA_CUDA_DEVICE_PTR_CUH_
+
+#include <stdint.h>
+
+/* `T` may be const-qualified; `address` is an unsigned 64-bit device address. */
+/* NOLINTBEGIN(performance-no-int-to-ptr): a CUdeviceptr is an integer by the C ABI of the kernel argument structs and a function wrapper costs ld.global.nc loads (measured above); ADR-1142 allows a cited NOLINT for a load-bearing invariant. */
+#define VMAF_CUDA_DPTR(T, address) (reinterpret_cast<T *>(address))
+/* NOLINTEND(performance-no-int-to-ptr) */
+
+#endif /* VMAF_SRC_CUDA_CUDA_DEVICE_PTR_CUH_ */

core/src/cuda/cuda_helper.cuh

@@ -122,7 +122,7 @@ static inline int vmaf_cuda_result_to_errno(int cu_err_code)
 #ifdef DEVICE_CODE
 namespace
 {
-typedef unsigned long long int uint64_cu;
+using uint64_cu = unsigned long long int;
 
 /* Warp sum of non-negative 64-bit terms whose total may pass INT64_MAX. Every
  * lane of the warp must call it: the shuffles take the full mask. */

core/src/feature/cuda/float_adm/float_adm_device.h

@@ -38,9 +38,11 @@
 #include "feature/float_adm_gpu_common.h"
 
 /* The argument blocks under the names the CUDA sources use. */
+/* NOLINTBEGIN(modernize-use-using): C header included by C and C++ translation units; C has no `using`. ADR-1138. */
 typedef FloatAdmGpuBands FloatAdmCudaBands;
 typedef FloatAdmGpuDecoupleArgs FloatAdmCudaDecoupleArgs;
 typedef FloatAdmGpuTermArgs FloatAdmCudaTermArgs;
 typedef FloatAdmGpuRowArgs FloatAdmCudaRowArgs;
+/* NOLINTEND(modernize-use-using) */
 
 #endif /* VMAF_SRC_FEATURE_CUDA_FLOAT_ADM_FLOAT_ADM_DEVICE_H_ */