fix(cuda): reassemble the int64 warp reduction, correct the ADM doc, add nvcc --threads

[lusoris]

Summary

Three findings from the CUDA Tile C++ adoption audit (ADR-1224),
which concluded Tile should not be adopted. The incidental findings turned
out to be worth more than the headline.

1. Dropped carry in an int64 warp reduction

integer_ssim_score.cu split the int64_t per-pixel weight into lo/hi
32-bit halves and summed each independently across the warp, recombining
only afterwards:

int32_t lo = (int32_t)(warp_wgt & 0xffffffffLL);
int32_t hi = (int32_t)((warp_wgt >> 32) & 0xffffffffLL);
for (int off = 16; off > 0; off >>= 1) {
    lo += __shfl_down_sync(0xffffffffu, lo, off);   /* carry out of lo is lost */
    hi += __shfl_down_sync(0xffffffffu, hi, off);
}

A carry out of the low half is dropped, and lo itself can overflow a signed
int32 (undefined behaviour, C17 6.3.1.4p1). warp_reduce(int64_t) at
cuda_helper.cuh:119 already reassembles the shuffled halves into an int64
before adding — this was a second, subtly different hand-rolled copy.

Not reachable today: the 9-tap per-pixel weight sum stays far below 2³¹, so
no score is currently wrong. It would surface at large frame sizes or under a
future weight scaling.

2. docs/backends/cuda/overview.md contradicted itself

One bullet said integer_adm_cuda lacks adm_csf_mode: 2 and that libvmaf
falls back to the CPU reference extractor. A section dated the same day said
the whole ADM family runs on the device for the default model. Three independent
checks agree with the later section:

• integer_adm_cuda.c:840 declares adm_csf_mode
• T-GPU-ADM-CSF-MODE-NOT-PORTED-2026-09-05 is closed
• a live CUDA run of the default model emits no computing it on the CPU notice

The stale bullet is corrected.

3. nvcc --threads

New nvcc_threads meson option (default 4). nvcc parallelises across the six
gencode architectures, so this shortens the critical-path kernel:
adm_cm.cu 8.5 s → 2.1 s, serial wall 41.9 s → 14.2 s.

Verified byte-identical here, not assumed — all 22/22 fatbins compare
equal between --threads 4 and --threads 1, so it cannot move a score. ninja
already builds the 22 fatbin targets concurrently, so the effective thread count
is this value × the ninja job count; hence an option rather than a hardcoded
number.

Reproducer / smoke test

meson setup build-cuda core -Denable_cuda=true -Denable_sycl=false -Db_lto=false
ninja -C build-cuda && meson test -C build-cuda --suite=fast     # 164/164

# byte-identity of the --threads change
ninja -C build-cuda $(cd build-cuda && ninja -t targets | grep -oE '^src/[a-z0-9_]+\.fatbin')
mkdir -p /tmp/t4 && cp build-cuda/src/*.fatbin /tmp/t4/
meson configure build-cuda -Dnvcc_threads=1
rm -f build-cuda/src/*.fatbin
ninja -C build-cuda $(cd build-cuda && ninja -t targets | grep -oE '^src/[a-z0-9_]+\.fatbin')
for f in /tmp/t4/*.fatbin; do cmp "$f" "build-cuda/src/$(basename $f)" || echo "DIFFERS: $f"; done

Measured on this workstation (RTX 4090, nvcc 13.3.73): byte-identical: 22 / 22.

Deliberately not adopted

--split-compile — three identical invocations produced three different
fatbin hashes. That breaks build reproducibility, which the keyless
Sigstore/SLSA release story depends on. Recorded in the ADR so it is not tried
again.

Why Tile itself is out

Summarised in the ADR. The short version: integer ct::mma() is 8-bit-operand
only and ~⅔ of the tree is int16/32/64; the longest contraction anywhere is a
17-tap constant filter; adm_csf_den.cu:101 and adm_cm.cu:266 apply a
rounding shift inside the accumulation, which is not a sum of products at
any dtype; tile extents must be powers of two (cuda_tile.h:749) and our rows
are 576 and 1920 wide; ct::sum() has no association parameter and no published
reduction-order guarantee; and --fatbin/--tilefatbin are not composable, with
--tilefatbin emitting no PTX at all.

The ADR also records what would change the answer, and explicitly retires
three arguments that did not survive verification — the architecture floor,
the CI pin, and a "not user-controllable" quote that does not exist in NVIDIA's
documentation.

Deep-dive deliverables (ADR-0108)

• Research digest — no digest needed: the analysis lives in ADR-1224 in full, including the measured numbers, the refuted claims and the re-evaluation triggers.
• Decision matrix — docs/adr/1224-cuda-tile-not-adopted.md ## Alternatives considered (four options; the runner-up "adopt for SpEED's covariance" fails on four independent grounds).
• AGENTS.md invariant note — no rebase-sensitive invariants: the int64 fix removes a local copy in favour of an existing shared helper, and the build option is fork-local.
• Reproducer / smoke-test command — see above.
• CHANGELOG fragment — changelog.d/fixed/1224-cuda-int64-warp-carry.md, changelog.d/fixed/1224-cuda-adm-doc-contradiction.md, changelog.d/changed/1224-nvcc-threads.md.
• Rebase note in docs/rebase-notes.md — entry ADR-1224 — CUDA Tile not adopted; audit findings banked (2026-09-07).

Docs (rule 10)

docs/backends/cuda/overview.md — the stale
ADM bullet is corrected. nvcc_threads is documented in
core/meson_options.txt with the oversubscription caveat.

Bug status (rule 13 / ADR-0165)

• docs/state.md — closes T-CUDA-SSIM-INT64-WARP-CARRY-2026-09-07, opens T-CUDA-ADM-CM-REGISTER-PRESSURE-2026-09-07 (REG:255 with 336–344 B spill on the default model's hot path; its own measured PR).

🤖 Generated with Claude Code