lusoris · May 28, 2026 · May 28, 2026
@@ -0,0 +1,11 @@
+- **`--feature ssim --backend cuda` silently broken since introduction.**
+  `integer_ssim/integer_ssim_score.cu` defined three `__global__`
+  kernels (`integer_ssim_horiz_8bpc`, `integer_ssim_horiz_16bpc`,
+  `integer_ssim_vert_combine`) without an `extern "C"` block. nvcc
+  C++ name-mangling caused `cuModuleGetFunction` to return
+  `CUDA_ERROR_NOT_FOUND`, so `init_fex_cuda` returned `-EINVAL` and
+  the `ssim` feature was never computed. Fixed by wrapping all three
+  entry points in `extern "C" { }`. A CI audit script
+  (`scripts/dev/check-cuda-extern-c.sh`) and an `AGENTS.md` invariant
+  note prevent recurrence. (ADR-0747; Research-0747; sweep of all 24
+  `.cu` kernel files confirmed no other instances.)
@@ -416,3 +416,29 @@ reason (e.g. a specific `dp4a` / `imad.lo.u32` targeting pattern), they must:
    constraint (CLAUDE.md §15).
 4. Document the constraint in `docs/backends/cuda/overview.md` under
    "Known gaps / CUDA version notes".
+## extern "C" invariant — mandatory for every new CUDA kernel TU (ADR-0747)
+
+Every `__global__` kernel that the host looks up by name via
+`cuModuleGetFunction` **must** be defined inside an `extern "C" { }`
+block in its `.cu` file.
+
+Rationale: nvcc compiles `.cu` files as C++ by default. Without
+`extern "C"`, the kernel symbol receives C++ name-mangling (e.g.,
+`_Z31integer_ssim_horiz_8bpc...`). `cuModuleGetFunction` uses the
+plain C name and receives `CUDA_ERROR_NOT_FOUND`, silently disabling
+the feature. This was found by a full sweep (Research-0747) that
+identified `integer_ssim/integer_ssim_score.cu` as broken; the
+pattern caused `--feature ssim --backend cuda` to fail silently from
+the file's introduction until this fix.
+
+Rules:
+- Wrap only `__global__` entry points. `__device__` helpers,
+  `__constant__` arrays, and `#define` macros do not need wrapping.
+- For macro-expanded kernel instantiations (e.g., the `FILTER1D_*`
+  and `ADM_CSF_KERNEL` patterns), place the macro invocations inside
+  the `extern "C" { }` block, not the macro definition.
+- CI gate: `scripts/dev/check-cuda-extern-c.sh` fails if any kernel
+  referenced by `cuModuleGetFunction` is found outside `extern "C"`.
+  Run it locally before pushing CUDA kernel changes.
+
+See [ADR-0747](../../../../docs/adr/0747-cuda-extern-c-sweep.md).
@@ -73,6 +73,8 @@ __device__ static const int32_t ISSIM_KERNEL[ISSIM_K_SZ] = {2, 9, 28, 55, 68, 55
  * Writes six arrays of width*height int64_t:
  *   d_mux, d_muy, d_x2, d_xy, d_y2, d_w
  */
+extern "C" {
+
 __global__ void integer_ssim_horiz_8bpc(const uint8_t *__restrict__ ref, ptrdiff_t ref_stride,
                                         const uint8_t *__restrict__ cmp, ptrdiff_t cmp_stride,
                                         int64_t *__restrict__ d_mux, int64_t *__restrict__ d_muy,
@@ -271,3 +273,4 @@ integer_ssim_vert_combine(const int64_t *__restrict__ d_mux_h, const int64_t *__
         partial_weights[blk] = block_wgt;
     }
 }
+} /* extern "C" */
@@ -0,0 +1,68 @@
+# ADR-0747 — CUDA `extern "C"` invariant for host-looked-up kernels
+
+**Status:** Accepted
+**Date:** 2026-05-28
+**Deciders:** lusoris
+**Related:** [Research-0747](../research/research-0747-cuda-extern-c-sweep.md)
+
+## Context
+
+nvcc compiles `.cu` files as C++. Any `__global__` kernel that is not
+declared inside an `extern "C"` block receives a C++ mangled symbol in
+the compiled PTX/cubin. When the host uses `cuModuleGetFunction` to
+resolve kernels by their plain C names, a mangled symbol causes
+`CUDA_ERROR_NOT_FOUND`, silently disabling the feature.
+
+PR #77 fixed one instance of this pattern in
+`integer_ssim/ssim_score.cu`. A full audit (Research-0747) identified
+one additional broken file introduced after PR #77:
+`integer_ssim/integer_ssim_score.cu`, which provides the `"ssim"`
+feature on the CUDA backend. All three of its entry points
+(`integer_ssim_horiz_8bpc`, `integer_ssim_horiz_16bpc`,
+`integer_ssim_vert_combine`) were missing `extern "C"` wrapping,
+silently breaking `--feature ssim --backend cuda` since the file was
+introduced.
+
+## Decision
+
+1. Wrap all three `__global__` entry points in
+   `integer_ssim/integer_ssim_score.cu` inside `extern "C" { }`.
+2. Add a CI script (`scripts/dev/check-cuda-extern-c.sh`) that fails
+   if any `__global__` function referenced by `cuModuleGetFunction` is
+   not covered by an `extern "C"` block.
+3. Add a mandatory invariant to `core/src/feature/cuda/AGENTS.md` so
+   that all future contributors are informed of the requirement before
+   writing new CUDA kernel TUs.
+
+## Alternatives considered
+
+- **Rename all kernels to a `C`-friendly symbol scheme using NVRTC
+  or embedding a `__attribute__((visibility("default")))` compiler
+  annotation**: rejected. These approaches are more complex and do not
+  address the root cause (C++ name-mangling). `extern "C"` is the
+  standard CUDA idiom used by the NVIDIA SDK examples and by every
+  other kernel in this codebase.
+- **Switch from `cuModuleGetFunction` (driver API) to the CUDA runtime
+  API**: out of scope for this PR. The driver API is intentional per
+  ADR-0001 (deferred CUDA context creation). Changing the dispatch
+  pattern would require pervasive refactoring across all 15 host
+  extractor files.
+- **No-op (accept the breakage)**: rejected. `--feature ssim --backend
+  cuda` silently produces no output in the broken state, which violates
+  the correctness-first principle (see CLAUDE.md §6 and
+  `feedback_correctness_first.md` memory entry).
+
+## Consequences
+
+- `--feature ssim --backend cuda` is restored to working order.
+- The CI script prevents future regressions of this class.
+- No numerical changes: `extern "C"` affects only symbol naming, not
+  kernel code generation.
+- The fix is a one-file, three-line insertion; rebase risk is minimal.
+
+## References
+
+- Research-0747 (`docs/research/research-0747-cuda-extern-c-sweep.md`)
+- `core/src/feature/cuda/ssim_cuda.c:122–127` (host lookup sites)
+- `core/src/meson.build:784` (PTX generation confirming which `.cu`
+  maps to `integer_ssim_score_ptx`)
@@ -550,3 +550,4 @@ ADRs may exist there for local session continuity, but the tracked
 | [ADR-0712](0712-ide-config-multilang-refresh.md) | IDE config audit and refresh for multi-language post-rebrand VMAFX: clangd, gopls, rust-analyzer, and Python LSP wired to the post-rebrand directory layout. | Accepted | 2026-05-28 | ide, clangd, gopls, rust-analyzer, build, fork-local |
 | [ADR-0714](0714-vmafx-operator-skeleton.md) | vmafx-operator kubebuilder skeleton + CRDs: `VmafxJob`, `VmafxNode`, `VmafxModelTraining` in API group `vmafx.dev/v1`; Stage 1 stub reconcilers; Helm chart integration; envtest suite. | Accepted | go, k8s, operator, crd, controller-runtime, phase4b, fork-local |
 | [ADR-0717](0717-vmafx-node-ffmpeg-latest.md) | vmafx-node ffmpeg version policy: pin to latest stable tagged release (n8.2); multi-stage Dockerfile with shared ffmpeg-builder-cpu stage; four variants (cpu/cuda/rocm/sycl); startup encoder-inventory probe. | Accepted | 2026-05-28 | node, ffmpeg, docker, phase4b, fork-local |
+| [ADR-0747](0747-cuda-extern-c-sweep.md) | CUDA `extern "C"` invariant for host-looked-up kernels: full sweep of 24 `.cu` files; one broken file found and fixed (`integer_ssim/integer_ssim_score.cu`); CI audit script added. | Accepted | 2026-05-28 | cuda, correctness, build, fork-local |
@@ -282,6 +282,20 @@ per-extractor coverage matrix.
   `core/src/feature/cuda/AGENTS.md`.
 
 
+- **Integer SSIM `extern "C"` sweep (fixed, ADR-0747)** — A full audit
+  of all 24 `.cu` kernel files confirmed that `integer_ssim/integer_ssim_score.cu`
+  was the only file with `__global__` kernels referenced by
+  `cuModuleGetFunction` but not wrapped in `extern "C"`.  This caused
+  `--feature ssim --backend cuda` to silently return `-EINVAL` from
+  `init_fex_cuda` (the driver returned `CUDA_ERROR_NOT_FOUND` for all
+  three kernel names) since the file was introduced.  Fixed in this PR by
+  wrapping the three entry points in `extern "C" { }`.  A CI script
+  (`scripts/dev/check-cuda-extern-c.sh`) prevents recurrence.
+  The analogous bug in `ssim_score.cu` was fixed earlier in PR #77.
+
+See [metrics/features.md](../../metrics/features.md) for the
+per-extractor coverage matrix.
+
 ## References
 
 - [CUDA C++ Best Practices Guide](https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/)

@@ -40034,3 +40034,21 @@ Critical issues that must be fixed before merge:
 - PR #58 `ref.cpp`: `make_unique` / C-caller `free()` allocator mismatch
 
 No rebase impact from the review itself; all findings are fixes required in those PRs.
+## `core/src/feature/cuda/integer_ssim/` — `extern "C"` on new kernels (ADR-0747)
+
+Any upstream or fork PR that adds a new `__global__` kernel to a `.cu`
+file under `core/src/feature/cuda/` or `core/src/cuda/` must wrap the
+entry point in `extern "C" { }` if it is also referenced by
+`cuModuleGetFunction` in the host `.c` glue.
+
+The invariant is enforced by `scripts/dev/check-cuda-extern-c.sh`.
+Run it locally before pushing. On upstream sync, if Netflix adds new
+CUDA kernels to their `libvmaf/src/feature/cuda/` tree, check whether
+those kernels use `extern "C"` in the upstream source and mirror the
+pattern here.
+
+This invariant was formalised after the audit that found
+`integer_ssim/integer_ssim_score.cu` missing `extern "C"`, silently
+breaking `--feature ssim --backend cuda` since introduction (PR #77
+fixed the analogous break in `ssim_score.cu`; ADR-0747 fixes
+`integer_ssim_score.cu`).
@@ -0,0 +1,97 @@
+# Research-0747 — CUDA `extern "C"` name-mangling sweep
+
+**Date:** 2026-05-28
+**Branch:** `audit/cuda-extern-c-name-mangling-sweep-20260528`
+**Status:** Complete — 1 broken file found and fixed.
+
+## Motivation
+
+PR #77 discovered that `integer_ssim/ssim_score.cu` had three
+`__global__` kernels not wrapped in `extern "C"`, causing
+`cuModuleGetFunction` to fail silently because the C++ name-mangler
+decorated the symbols while the driver API lookup used the plain C
+names. The fix for that file was merged, but no systematic sweep was
+performed for the rest of the codebase.
+
+This research covers every `__global__` CUDA kernel in
+`core/src/feature/cuda/` and `core/src/cuda/` that is also looked
+up by name via `cuModuleGetFunction` on the host side.
+
+## Inventory
+
+24 `.cu` files contain `__global__` definitions. 46 distinct kernel
+names are looked up by `cuModuleGetFunction` across 15 host `.c`
+files.
+
+### Classification
+
+| File | `extern "C"` present | All looked-up kernels covered | Status |
+|------|----------------------|-------------------------------|--------|
+| `integer_ssim/ssim_score.cu` | Yes (line 37) | Yes (`calculate_ssim_*`) | SAFE |
+| `integer_ssim/integer_ssim_score.cu` | **No** | N/A — 3 kernels exposed bare | **BROKEN** |
+| `integer_adm/adm_csf.cu` | Yes (line 205) | Yes (`adm_csf_kernel_1_4`, `i4_adm_csf_kernel_1_4`) | SAFE |
+| `integer_adm/adm_csf_den.cu` | Yes (end of file) | Yes (`adm_csf_den_scale_line_kernel_8_128`, `adm_csf_den_s123_line_kernel_8_128`) | SAFE |
+| `integer_adm/adm_dwt2.cu` | Yes (end of file) | Yes (all DWT + DWT2 kernels) | SAFE |
+| `integer_adm/adm_cm.cu` | Yes (two blocks, lines 143 + 378) | Yes (`adm_cm_line_kernel_8`, `i4_adm_cm_line_kernel_fused`) | SAFE |
+| `integer_adm/adm_decouple.cu` | Yes (line 41) | No lookup — `adm_decouple_kernel` / `adm_decouple_s123_kernel` are not looked up by `cuModuleGetFunction`; they are device-side | SAFE |
+| `float_vif/float_vif_score.cu` | Yes (line 116) | Yes (`float_vif_compute`, `float_vif_decimate`) | SAFE |
+| `integer_ciede/ciede_score.cu` | Yes (line 31) | Yes | SAFE |
+| `integer_psnr_hvs/psnr_hvs_score.cu` | Yes (line 27) | Yes (`psnr_hvs`) | SAFE |
+| `integer_ms_ssim/ms_ssim_score.cu` | Yes (line 42) | Yes | SAFE |
+| `ssimulacra2/ssimulacra2_blur.cu` | Yes (line 65) | Yes (5 kernels including `ssimulacra2_transpose`) | SAFE |
+| `ssimulacra2/ssimulacra2_mul.cu` | Yes (line 21) | Yes (`ssimulacra2_mul3`) | SAFE |
+| `integer_cambi/cambi_score.cu` | Yes (line 99) | Yes | SAFE |
+| `integer_motion_v2/motion_v2_score.cu` | Yes (line 55) | Yes | SAFE |
+| `speed/speed_score.cu` | Yes (line 71) | Yes (5 kernels) | SAFE |
+| `float_motion/float_motion_score.cu` | Yes (line 49) | Yes | SAFE |
+| `float_psnr/float_psnr_score.cu` | Yes (line 27) | Yes | SAFE |
+| `float_adm/float_adm_score.cu` | Yes (line 101) | Yes (6 kernels) | SAFE |
+| `integer_motion/motion_score.cu` | Yes (line 62) | Yes | SAFE |
+| `integer_moment/moment_score.cu` | Yes (line 36) | Yes | SAFE |
+| `integer_vif/filter1d.cu` | Yes (line 832) | Yes (10 macro-expanded kernels) | SAFE |
+| `integer_psnr/psnr_score.cu` | Yes (line 34) | Yes | SAFE |
+| `integer_adm_cuda.c` (note: `.c` file with grep hit on comment) | n/a — not a `.cu` file | n/a | n/a |
+
+### The one broken file
+
+**`core/src/feature/cuda/integer_ssim/integer_ssim_score.cu`**
+
+- Compiled as `integer_ssim_score_ptx` (confirmed in
+  `core/src/meson.build` line 784).
+- Loaded by `core/src/feature/cuda/ssim_cuda.c` via
+  `cuModuleLoadData(&module, integer_ssim_score_ptx)`.
+- Three `__global__` kernel entry points:
+  - `integer_ssim_horiz_8bpc`
+  - `integer_ssim_horiz_16bpc`
+  - `integer_ssim_vert_combine`
+- All three are looked up by name in `ssim_cuda.c` lines 122–127.
+- The file had **zero** `extern "C"` declarations. nvcc compiles
+  `.cu` files as C++ by default; without `extern "C"`, the symbols
+  receive C++ name-mangling (`_Z31integer_ssim_horiz_8bpc...`).
+  `cuModuleGetFunction` passes the plain C names and receives
+  `CUDA_ERROR_NOT_FOUND`, which the host `CHECK_CUDA_GOTO` macro
+  logs and returns `-EINVAL` from `init_fex_cuda`. The feature
+  `"ssim"` via `--backend cuda` is therefore silently non-functional
+  from the time this file was introduced.
+
+## Fix
+
+Added `extern "C" {` immediately before the first `__global__`
+definition (line 76) and `} /* extern "C" */` at end of file.
+Device-only helpers (`__device__ static`) and `#define` constants are
+not affected.
+
+## Audit script
+
+`scripts/dev/check-cuda-extern-c.sh` codifies this check for CI.
+It fails with exit code 1 if any `__global__` function in a `.cu`
+file is referenced by `cuModuleGetFunction` but is not inside an
+`extern "C"` block.
+
+## Before / after smoke
+
+No container build is available in the worktree environment (no
+GPU), but the name-mangling evidence is definitive without a binary:
+the function names looked up by the host are plain C identifiers; if
+the compiled PTX exports only mangled names, the lookup always fails.
+The fix eliminates the mismatch.
@@ -207,6 +207,7 @@ landed fix yet._
 | T-CUDA-MUL24-AUDIT-2026-05-28 | CUDA `__mul24` silent-corruption sweep (Research-0734). Audited all 78 files under `core/src/feature/cuda/` and `core/src/cuda/` for `__mul24` / `__umul24` / `__mul24hi` usage. Zero instances found; no scores affected by the CUDA 11.1–13.3 silent-corruption bug. Prohibition invariant added to `core/src/feature/cuda/AGENTS.md`; known-issue note added to `docs/backends/cuda/overview.md`. | [Research-0734](research/0734-cuda-mul24-corruption-audit.md) | audit/cuda-mul24-corruption-20260528 | `grep -rn '__mul24\|__umul24\|__mul24hi' core/src/feature/cuda/ core/src/cuda/` — no output (exit 1). | (2026-05-28) |
 | T-POST-RENAME-DRIFT-SWEEP-2026-05-28 | Post-ADR-0700 path drift sweep: fixed 9 stale `libvmaf/` and `python/vmaf/` directory references across `Makefile`, `Dockerfile`, `.github/codeql-config.yml`, `.vscode/c_cpp_properties.json`, `.zed/settings.json`, `.claude/skills/add-gpu-backend/scaffold.sh`, `scripts/dev/project_modernization_audit.py`, `README.md`, `AGENTS.md`, and `.claude/skills/add-model/SKILL.md`. No ADR (maintenance fix). | no ADR | chore/post-rename-drift-sweep-20260528 | `grep -rn 'libvmaf/[a-z]' . --include='*.sh' --include='*.py' --include='Makefile' ... \| grep -vE '(libvmaf\.so\|core/include/libvmaf\|...)` returns zero actionable hits after this PR. | (2026-05-28) |
 | T-CUDA-VIF-FILTER1D-NCU-HOTPATH-2026-05-28 | CUDA VIF `filter1d.cu` profiled with ncu 2026.2.0.0 on RTX 4090 (sm_89). Primary hotspot: `filter1d_8_horizontal_kernel_2_17_9` (35 % of VIF filter time, 20.8 µs avg). Diagnosis: launch-width-limited (0.84 waves), register-pressure ceiling (56 regs → 75 % theoretical occupancy), L2 hot-slice imbalance (+46 %). Next: evaluate `val_per_thread` 2→4 for scale-0 horizontal pass. | [Research-0734](research/0734-cuda-vif-filter1d-ncu-hotpath-20260528.md) | research/cuda-vif-ncu-hotpath-20260528 | `docker run ... ncu ... /workspace/core/build-ncu/tools/vmaf ...` reproducer in Research-0734. | (2026-05-28) |
+| T-CUDA-EXTERN-C-SWEEP-0747-2026-05-28 | Full sweep of all 24 `.cu` kernel files across `core/src/feature/cuda/` and `core/src/cuda/` found one broken file: `integer_ssim/integer_ssim_score.cu` (three `__global__` kernels not wrapped in `extern "C"`, causing `--feature ssim --backend cuda` to silently fail since introduction). Fixed by adding `extern "C" { }` wrapping. CI audit script added at `scripts/dev/check-cuda-extern-c.sh`; invariant documented in `core/src/feature/cuda/AGENTS.md`. All other 23 kernel files confirmed safe. | [ADR-0747](adr/0747-cuda-extern-c-sweep.md) / [Research-0747](research/research-0747-cuda-extern-c-sweep.md) | audit/cuda-extern-c-name-mangling-sweep-20260528 | `bash scripts/dev/check-cuda-extern-c.sh` exits 0 | (2026-05-28) |
 | T-VMAFX-NODE-IMPL-2026-05-28 | vmafx-node Go worker binary (Phase 4b.2) shipped: cgo libvmaf scoring, GPU vendor detection (`pkg/gpu/`), ffmpeg hardware encoder support, ONNX inference registry (`pkg/ai/`), gRPC VmafxController client (`gen/go/controller/`), Prometheus metrics, SIGTERM graceful drain, multi-variant Docker images, Helm node Deployment. | [ADR-0713](adr/0713-vmafx-node-impl.md) | feat/vmafx-node-go-worker-binary | `go test ./pkg/gpu/ ./pkg/ai/ ./cmd/vmafx-node/ -v` — all pass; `go vet ./...` clean. | (2026-05-28) |
 | T-VMAFX-SIDECAR-TRAINING-RESEARCH-0733-2026-05-28 | Research digest for Phase 4b.7 sidecar online training architecture complete. Evaluates three architecture options; recommends Option A (Python sidecar container per node). Specifies gRPC triple transport, EMA-SGD training loop with replay buffer, atomic checkpoint swap, `VmafxModelTraining` CRD, and operator reconcile loop. | [ADR-0709](adr/0709-vmafx-phase4b-distributed-platform.md) / [Research-0733](research/0733-vmafx-sidecar-training-architecture.md) | docs/research-0733-vmafx-sidecar-training-arch | doc-only PR; no test runner applicable — research digest filed. | (2026-05-28) |
 | T-VMAFX-RUST-PILOT-TAD-2026-05-28 | TAD (Temporal Absolute Difference) feature extractor implemented in Rust and wired into libvmaf.so via cbindgen + Meson custom_target. Proves the Phase 4 Rust-in-libvmaf integration story end-to-end. New `--feature tad` signal available; does not affect existing VMAF scores. Workspace `Cargo.toml` established at repo root. | [ADR-0707](adr/0707-vmafx-rust-pilot-feature.md) | feat/tad-rust-pilot | `cargo test --manifest-path core/src/feature/rust/tad/Cargo.toml` — 5/5 Rust unit tests pass; `meson test -C build-tad-pilot test_tad_rust` — 4/4 C smoke tests pass | (2026-05-28) |