Skip to content

fix(cuda)!: wrap all host-looked-up __global__ kernels in extern "C" (P0 silent-corruption sweep) - #80

Merged
lusoris merged 1 commit into
masterfrom
audit/cuda-extern-c-name-mangling-sweep-20260528
May 28, 2026
Merged

fix(cuda)!: wrap all host-looked-up __global__ kernels in extern "C" (P0 silent-corruption sweep)#80
lusoris merged 1 commit into
masterfrom
audit/cuda-extern-c-name-mangling-sweep-20260528

Conversation

@lusoris

@lusoris lusoris commented May 28, 2026

Copy link
Copy Markdown
Contributor

Summary

  • Full sweep of all 24 .cu kernel files in core/src/feature/cuda/ and core/src/cuda/ against every cuModuleGetFunction host-side lookup.
  • One broken file found: integer_ssim/integer_ssim_score.cu — three __global__ entry points not wrapped in extern "C", causing --feature ssim --backend cuda to silently return -EINVAL from init_fex_cuda since the file was introduced.
  • Fix: three-line insertion of extern "C" { } around entry points. All 23 other kernel files confirmed safe.
  • CI guard added: scripts/dev/check-cuda-extern-c.sh fails if any kernel referenced by cuModuleGetFunction is found outside extern "C". Mandatory invariant documented in core/src/feature/cuda/AGENTS.md.

This is a companion to PR #77 (which fixed the same pattern in ssim_score.cu). The ADR-0747 sweep confirms the audit is now complete across the full codebase.

Root cause

nvcc compiles .cu files as C++. Without extern "C", the symbols receive C++ name-mangling (_Z31integer_ssim_horiz_8bpc...). cuModuleGetFunction passes the plain C name and receives CUDA_ERROR_NOT_FOUND. The feature appears to init successfully up to that point but then silently produces no output.

Test plan

  • bash scripts/dev/check-cuda-extern-c.sh exits 0 (all 46 looked-up kernels confirmed wrapped).
  • Container smoke: docker run --rm --gpus all -v $(pwd):/workspace -w /workspace vmaf-dev-mcp:cuda13.3 bash -c 'ninja -C build-audit && build-audit/tools/vmaf --reference python/test/resource/yuv/src01_hrc00_576x324.yuv --distorted python/test/resource/yuv/src01_hrc01_576x324.yuv --width 576 --height 324 --pixel_format 420 --bitdepth 8 --feature ssim --backend cuda -o /tmp/scores-ssim.json' — exits 0 + emits scores.
  • Cross-backend correctness vs --backend cpu at places=4 (ADR-0214).

Deliverables checklist (ADR-0108)

  • (a) Research digest: docs/research/research-0747-cuda-extern-c-sweep.md
  • (b) Decision matrix: docs/adr/0747-cuda-extern-c-sweep.md §Alternatives considered
  • (c) AGENTS.md invariant: core/src/feature/cuda/AGENTS.md + scripts/dev/check-cuda-extern-c.sh
  • (d) Reproducer: bash scripts/dev/check-cuda-extern-c.sh
  • (e) Changelog: changelog.d/fixed/cuda-extern-c-sweep.md
  • (f) Rebase notes: docs/rebase-notes.md

State / bug tracking (CLAUDE §12 r13)

  • docs/state.md row added (T-CUDA-EXTERN-C-SWEEP-0747-2026-05-28 → Recently closed).

FFmpeg patches (CLAUDE §12 r14)

  • no rebase impact: fix is internal to .cu device code; no public C-API surfaces or headers changed.

🤖 Generated with Claude Code

@lusoris
lusoris enabled auto-merge (squash) May 28, 2026 20:59
@lusoris lusoris mentioned this pull request May 28, 2026
Closed
9 tasks done
@lusoris @claude
fix(cuda)!: wrap all host-looked-up __global__ kernels in extern "C" …
be587bb 
…(P0 silent-corruption sweep)

Full audit of all 24 .cu kernel files in core/src/feature/cuda/ and
core/src/cuda/ against every cuModuleGetFunction host-side lookup.

Finding: integer_ssim/integer_ssim_score.cu is the only file with
__global__ kernels referenced by cuModuleGetFunction but not wrapped
in extern "C". nvcc compiles .cu as C++ by default; the three entry
points (integer_ssim_horiz_8bpc, integer_ssim_horiz_16bpc,
integer_ssim_vert_combine) received C++ name-mangling, causing the
driver to return CUDA_ERROR_NOT_FOUND for all three lookups.
init_fex_cuda returned -EINVAL, silently disabling --feature ssim
--backend cuda since the file was introduced (PR #77 fixed the
analogous break in ssim_score.cu).

Fix: add extern "C" { } around all three __global__ entry points.
Device-only helpers (__device__ static, __constant__) are not affected.
All 23 other kernel files are confirmed safe.

Deliverables (ADR-0108):
  (a) Research digest: docs/research/research-0747-cuda-extern-c-sweep.md
  (b) ADR-0747: docs/adr/0747-cuda-extern-c-sweep.md
  (c) AGENTS.md invariant + CI script: scripts/dev/check-cuda-extern-c.sh
  (d) Reproducer: bash scripts/dev/check-cuda-extern-c.sh
  (e) Changelog: changelog.d/fixed/cuda-extern-c-sweep.md
  (f) Rebase notes: docs/rebase-notes.md

Affected features: ssim (--backend cuda).

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
@lusoris
lusoris force-pushed the audit/cuda-extern-c-name-mangling-sweep-20260528 branch from 6c3eadb to be587bb Compare May 28, 2026 21:36
@lusoris
lusoris merged commit 396e8ac into master May 28, 2026
41 of 63 checks passed
@lusoris
lusoris deleted the audit/cuda-extern-c-name-mangling-sweep-20260528 branch May 28, 2026 21:36
This was referenced May 29, 2026
Merged
Closed
lusoris added a commit that referenced this pull request May 29, 2026
@lusoris @claude
chore(release): set release-please draft mode for manual version review
a0949ff 
The next release-please tag is projected as 4.0.0-lusoris.0 (MAJOR) because
five commits carry '!' breaking markers since 3.0.0-lusoris.0. Three of those
markers are incorrect:

  PR #52  feat!: sunset legacy native build modes
          CI-matrix pruning; no public API, CLI flag, or header removed.
          '!' unwarranted.

  PR #80  fix(cuda)!: wrap __global__ kernels in extern "C"
          Internal CUDA kernel mangling fix (P0 silent-corruption bug).
          No public API change. '!' unwarranted.

  PR #108 fix(cuda)!: remove committed conflict marker
          Three-line literal-marker deletion. Not a breaking change.
          '!' unwarranted.

Two '!' commits ARE correctly marked:

  PR #47  feat(core)!: drop Vulkan backend
          Removed libvmaf_vulkan.h (public header), CLI flags
          --backend vulkan / --vulkan_device / --vulkan-require-fp64,
          and public enum values. Genuine public API removal.

  PR #87  feat!: sunset VmafLegacyQualityRunner
          Removed importable Python class VmafLegacyQualityRunner.
          Genuine public surface removal.

Net assessment: 2/5 breaking markers warrant a MAJOR bump; the other 3
are bug-fix or CI-maintenance commits mislabelled with '!'. The fork
tracks Netflix upstream v3.x; jumping to 4.0.0 prematurely would
misrepresent the version relative to upstream and surprise downstream
users.

Safest mitigation without rewriting history: set "draft": true in
release-please-config.json so the next release PR opens as DRAFT. The
maintainer reviews the proposed version, adjusts if needed, then
un-drafts to merge. This adds one manual gate without masking future
real majors.

The version chosen on review should be 3.1.0-lusoris.0: the two genuine
breaking changes (Vulkan + LegacyRunner) are fork-local extensions with
no upstream counterpart, and the fork has not bumped its
upstream-tracking MAJOR.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
lusoris added a commit that referenced this pull request May 29, 2026
@lusoris @claude
chore(release): set release-please draft mode for manual version review
8a4b316 
Set `"draft": true` in `release-please-config.json` root package so the
next release PR opens as a draft, requiring manual review before merge.
Prevents an unintended `4.0.0` major bump caused by three incorrectly
marked breaking commits (PR #52, #80, #108).

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
lusoris added a commit that referenced this pull request May 29, 2026
@lusoris @claude
chore(release): set release-please draft mode for manual version review
6ea0651 
Set `"draft": true` in `release-please-config.json` root package so the
next release PR opens as a draft, requiring manual review before merge.
Prevents an unintended `4.0.0` major bump caused by three incorrectly
marked breaking commits (PR #52, #80, #108).

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
lusoris added a commit that referenced this pull request Jun 2, 2026
@lusoris @claude
chore(release): set release-please draft mode for manual version review
2b9997f 
Set `"draft": true` in `release-please-config.json` root package so the
next release PR opens as a draft, requiring manual review before merge.
Prevents an unintended `4.0.0` major bump caused by three incorrectly
marked breaking commits (PR #52, #80, #108).

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
lusoris added a commit that referenced this pull request Jun 2, 2026
@lusoris @claude
chore(ci): CI workflow hygiene bundle (#141 TSan/Python + #165 ubuntu…
fd63d7f 
…-24.04 pin + #173 Renovate grouping + #183 release-please draft) (#520)

* chore(ci): nightly workflow audit — remove redundant TSan, fix retention + Python ver (ADR-0793)

Three targeted fixes surfaced by a periodic nightly-CI audit:

1. Remove the `tsan` job from nightly.yml. sanitizers.yml already fires
   TSan on every push to master (ADR-0710); the daily cron duplicate burned
   ~45 runner-minutes/night for zero additional signal.

2. Add explicit `retention-days` to both nightly artifacts: 14 d for the
   clang-tidy-full-report (diagnostic value expires quickly) and 30 d for
   nightly-benchmark-results (one month of period-over-period comparisons).
   Both previously defaulted to GitHub's 90-day retention.

3. Fix `python-version: "3.14.5"` → `"3.12"` in nightly-bisect.yml.
   Python 3.14 is a pre-release alpha series with no stable release; the
   step name ("Set up Python 3.12") was correct and the version string was
   wrong. This would cause the job to fail trying to download a non-existent
   release.

No functional changes to test coverage. All nightly signal is preserved.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>

* chore(ci): pin ubuntu-latest to ubuntu-24.04 across all non-Docker runners (ADR-0802)

Replace floating `ubuntu-latest` runner alias with `ubuntu-24.04` across 15
workflow files to prevent silent toolchain drift when GitHub promotes the alias
to Ubuntu 26.04 (expected H2 2026). Add ADR-0802 documenting the pin policy.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>

* chore(deps): Renovate — Go/Cargo grouping, off-hours schedule, PR-cap (ADR-0812)

- Global schedule: "at any time" → "before 6am on weekdays" (Europe/Vienna);
  vulnerability alerts retain their existing "at any time" override.
- Add gomod packageRule: minor+patch grouped, auto-merged Monday mornings;
  major individual, manual review.
- Add cargo packageRule: same group-and-automerge pattern; major manual.
- prConcurrentLimit: 12 → 10 (Go grouping reduces PR count per cycle).
- ADR-0812 documents the decision and alternatives.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>

* chore(release): set release-please draft mode for manual version review

Set `"draft": true` in `release-please-config.json` root package so the
next release PR opens as a draft, requiring manual review before merge.
Prevents an unintended `4.0.0` major bump caused by three incorrectly
marked breaking commits (PR #52, #80, #108).

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>

* chore(ci): add bundle CHANGELOG fragment for CI workflow hygiene PRs #141 #165 #173 #183

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>

---------

Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>
Co-authored-by: Lusoris <lusoris@pm.me>
@lusoris lusoris added this to the 1.0.0 — First release milestone Sep 4, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

1.0.0 — First release

Development

Successfully merging this pull request may close these issues.

1 participant

@lusoris