Research-0734 — CUDA 13.3 Fix-List Deep Audit: Per-Issue Exposure Mapping¶

Date: 2026-05-28 Author: Lusoris (Claude agent) Branch: docs/cuda-13.3-fix-list-deep-audit-20260528 Supersedes / extends: PR #64 (thread-reconvergence + __mul24 + C++23 unblock)

Executive Summary¶

This digest audits every "Fixed an issue where…" / "Resolved an issue where…" entry in the CUDA 13.3 (and 13.2 / 13.1 / 13.0) release notes and maps each to our codebase.

No new CRITICAL findings beyond the thread-reconvergence issue already tracked by PR #64. However, two items merit action:

1. CRITICAL (already tracked — still unresolved for our build env): The thread- reconvergence compiler bug [6156910] is present in NVCC since 12.8. Our Containerfile currently pins cuda-toolkit-13-2. Every .cu kernel compiled with that toolchain is potentially generating wrong register values after nested-divergent branches. This affects all 25+ .cu kernels in core/src/feature/cuda/, most critically float_adm_score.cu, adm_decouple.cu, integer_vif/filter1d.cu, and integer_adm/adm_dwt2.cu, which all contain multiple-level nested if divergence (42–86 branch points per file). The fix is in NVCC 13.3 only.

2. HIGH — cuFFT floating-point exceptions [5923044]: If multi-GPU cuFFT plans are ever used (not currently in tree, but cuFFT header is available), cufftXtQueryPlan can generate FP exceptions. Not exposed today; annotated below as scope-guarded.

3. MEDIUM — ptxas WGMMA data race: Affects Hopper/Blackwell wgmma.mma_async instruction sequences only. Our kernels do not use wgmma instructions; not exposed.

4. MEDIUM — NVRTC + lld/mold initialisation failure [5020829]: Our build does not use NVRTC (runtime compilation). Not exposed.

Severity Rubric¶

Full Per-Issue Classification Table¶

CUDA 13.3¶

CUDA 13.2 / 13.2 Update 1¶

CUDA 13.1 / 13.1 Update 1¶

CUDA 13.0 / 13.0 Updates¶

Summary Counts¶

Action Items¶

1. Containerfile CUDA pin bump (CRITICAL): Upgrade cuda-toolkit-13-2 → cuda-toolkit-13-3 in dev/Containerfile (line 186). Reconvergence bug [6156910] is live in 13.2. ADR not required (pin-version change, not architectural). Run /cross-backend-diff after upgrade.
2. Dockerfile main image pin (CRITICAL): Dockerfile FROM line references nvidia/cuda:13.2.1-devel-ubuntu24.04. This should be bumped to 13.3.x when that image is published.
3. Annotate __syncwarp call in speed/speed_score.cu: The single __syncwarp call (line 279) is in a warp-uniform code path but sits near a conditional. After the 13.3 pin bump, run the GPU parity gate to confirm no delta.

Reproducer / Audit Commands¶

# 1. Fetch the release notes
# https://docs.nvidia.com/cuda/cuda-toolkit-release-notes/index.html

# 2. Enumerate our CUDA sources
find core/src/feature/cuda core/src/cuda -name '*.cu' -o -name '*.cuh' | sort

# 3. Grep for __mul24 (confirmed zero)
grep -rn '__mul24\|__umul24' core/src/

# 4. Grep for NVRTC (confirmed zero)
grep -rn 'nvrtc\|NVRTC' core/src/ mcp-server/

# 5. Grep for WGMMA (confirmed zero)
grep -rn 'wgmma\|mma\.sync\|wmma' core/src/

# 6. Count nested divergence (worst offenders)
grep -c 'if\b' core/src/feature/cuda/integer_vif/filter1d.cu        # 86
grep -c 'if\b' core/src/feature/cuda/float_adm/float_adm_score.cu   # 42
grep -c 'if\b' core/src/feature/cuda/integer_adm/adm_decouple.cu    # 18

# 7. Verify Containerfile CUDA version
grep 'cuda-toolkit' dev/Containerfile

References¶

• CUDA Toolkit 13.3 Release Notes: https://docs.nvidia.com/cuda/cuda-toolkit-release-notes/index.html
• PR #64: thread-reconvergence + __mul24 + C++23 unblock
• PR #65: __mul24 audit (zero exposure confirmed)
• Research-0734: this document