fix(KERNEL-SSM-MAMBA): the derived bound omitted nvcc FMA contraction, and the memory-safety claim outran the kernels (#496)

[localai-bot]

Tightening pass on the fresh review of #566 (row/KERNEL-SSM-MAMBA-SSD-W2-FINISH @ 1e819144e, W2 of .agents/specs/mamba2-ssd.md). That review's verdict was PASS; these are its two MEDIUM and three LOW findings. Branched from 1e819144e with origin/main merged, so it is standalone and does not stack on #566.

Issue: #496.

The review verified the equivalence contract's central MMA reasoning, reconstructed the reformulated mutations on CPU twins and proved each reds, and confirmed the grep, the recovered-byte md5s and the CI baseline subtraction. None of that is re-done here.

F1 (MEDIUM) — the derived bound's derivation omitted nvcc FMA contraction

The header and both test headers claimed the elementary functions were the only admitted source of device-vs-host divergence. False as written. Host C++ is pinned -ffp-contract=off (CMakeLists.txt:55) so a*b + c keeps two roundings; nothing passes --fmad=false to nvcc, so cuda_mamba2_ssd.cuh compiles at the default --fmad=true and every acc += a*b in it is a single-rounding fma whose host twin is not. This repo has already measured that idiom (.agents/benchmark-record.md:532). CMakeLists.txt:41-56 carves CUDA out of the contraction policy because "GPU parity tests compare GPU-vs-GPU" — G2 is exactly the case that carve-out does not cover.

The arithmetic:

term	bound	source
libm	2.5·K·u	≤ 2.5 ulp per decay factor (CUDA expf ≤ 2, glibc ≤ 0.5), through a product of at most K
summation	(K-1)·u	length-K f32 sum; this is what amplifies the libm term
contraction	K·u	the K product roundings the host keeps and the device's fma does not
total	4.5·K·u − u

• old 4·(K+2)·u: 4.5K − 1 ≤ 4K + 8 ⟺ K ≤ 18. Not provable at the driver shapes, which run at K = T = 200.
• new 5·(K+2)·u: 4.5K − 1 ≤ 5K + 10 ⟺ 0.5K + 11 ≥ 0. All K ≥ 0.

-fmad=false was weighed and REJECTED, not overlooked: nvcc takes it per translation unit and this is a header included by cuda_gdn.cu:48, so applying it means de-contracting every GDN decode kernel in that TU (a measured hot path) or splitting a new src/vt/ TU, which §8.3 already records as blocked on #515. Slowing a shipped kernel to make a bound's prose true is the wrong trade.

Nothing was hidden numerically. Re-scaling §8.4's audit by 4/5: worst of 55 comparisons 7.66% → 6.13% of budget, driver shapes 0.32%/0.18% → 0.26%/0.14%, mutant M3 962173% → 769738% — still caught by four orders of magnitude. The constant moved because the derivation gained a term the build emits. The header comment, all three test comments and spec §8.3 now say the same true thing.

F2 (MEDIUM) — the memory-safety claim outran the kernels

Both halves of the finding are taken: the clamps AND the narrowed claim.

The shared validator checks metadata shape/dtype/device only (CheckI32Meta, ops.cpp:1717-1723); every value check lives in the CPU kernel (cpu_ops.cpp:1622-1648), so the device arm silently drops six. The stated reason for dropping them — a D2H plus a stream sync — does not apply to the memory-unsafe ones: those values are already in device registers, and the decode kernel has always clamped its state_indices slot for free on that basis.

Clamped:

• M2StatePassKernel — chunk_end to nchunks, chunk_start to 0. Unclamped, lci[b] >= nchunks makes M2Store(passed, …) an out-of-bounds write past the cudaMallocAsync allocation (W1 finding F7's device half).
• M2ChunkScanKernel — si_ok = si >= 0 && si < S, and !si_ok opens the chunk from a zero previous state. Unclamped, an out-of-range seq_idx[c] reads initial_states out of bounds, and a seq_idx[0] < 0 makes si == si_prev at c == 0 and indexes passed at chunk −1 — a hole the finding did not name, found while writing the clamp. The kernel gained an S parameter for it.

In-contract behaviour is bit-identical: in contract si is always in range and lci always below nchunks, so neither clamp can fire.

Narrowed: clamping two does not make the arm memory-safe. The cu_chunk_seqlens tiling and per-chunk length checks are NOT clamped and a violation IS memory-unsafe (a garbage ccs indexes x/B/C/z/out out of bounds in every stage). The header and §8.3 now say that, instead of a blanket "the device kernels remain MEMORY SAFE".

Pinned, not asserted in prose: a device-only case runs both violations against in-contract reference runs whose result each clamp is defined to reproduce, so the assertions are exact rather than tolerances.

F3 / F4 (LOW, record accuracy)

• M7's §8.4 label overstated the device mutant: the launcher's nblocks = rows * args.n_groups is not mutated while the kernel's n_groups is forced to 1, so blocks blk >= rows run past the tensor. The mutant is memory-unsafe and fails partly for that, not purely on whole-row variance. The guarantee is pinned by the reviewer's clean CPU twin; only the label was wrong.
• §8.2's residual still gave the downcasts as the reason W2 cannot byte-compare, written when W2 was expected to mirror them. §8.3 supersedes it: both arms stay f32 and the reasons are libm and contraction. Two sentences, two causes, one fact — reconciled.

F5 (LOW) — taken

The five per-call scratch buffers are held by an M2Scratch scope guard, so a throw on the Nth cudaMallocAsync no longer leaks the N−1 before it; Release() also frees all five before reporting rather than leaking the remainder on a mid-sequence free failure.

Evidence (implementer, CPU box)

df -h / logged beside every result: 83–87% used / 60–73G free throughout.

suite	test cases	assertions	status
test_ops_mamba2_ssd	8 / 8 passed	1175 / 1175	SUCCESS!
test_ops_mamba2_state_update	6 / 6 passed	2469 / 2469	SUCCESS!
test_ops_mamba2_gated_norm	9 / 9 passed	2107 / 2107	SUCCESS!

Identical to the pre-change counts, as expected — every code change is inside #ifdef VLLM_CPP_CUDA or in the .cuh. Status: was read, not assertions: alone. Clean Release rebuild: 817/817 edges, 0 warnings, 0 errors under -Wall -Wextra -Werror. Full ctest: 403/403 passed, 0 failed. scripts/agent-preflight.sh --staged: RC=0.

The CUDA arm COMPILES under real nvcc — operator-run, not CI

dgx.casa returned at 06:57 UTC after a reboot. The operator (not me — this box has no nvcc and no GPU) transferred this branch by git archive (never rsync) and built it with real nvcc:

NVCC: cuda_13.0.r13.0
CUTLASS found at ~/cutlass-4.5.0; enabling sm120a NVFP4 cutlass GEMM
Marlin NVFP4 W4A16 MoE GEMM enabled (vendored) for [121a]
FlashAttention-2 prefill/decode: ENABLED for arch(es) [121a]
CONFIGURE_EXIT=0   BUILD_EXIT=0   WARNINGS=0   ENOSPC=0

Zero errors, zero warnings under -Werror, disk unchanged at 64G either side (so not the stale-binary false green), and the three fast-path features read out of the configure log rather than assumed — an absent CUTLASS also exits 0, so "the build succeeded" alone would prove nothing.

That retires the risk the shim check below could only approximate: the new S kernel parameter and the M2Scratch guard compile through cuda_gdn.cu at the shipping arch. Attributed to the operator, not to CI — cuda-fat-build has still never completed on this branch.

⚠️ The red checks on this PR are CANCELLED, not FAILED

Every Actions run here — four SHAs, both workflows — ended conclusion: cancelled. Three were my own follow-up pushes. The fourth was killed at 07:44:45–07:44:55 together with every run in the repository (20 of 20 across 7 branches); windows-msvc-cpu died mid-build after passing two steps. That is an Actions-side event, not a verdict on this diff.

gh pr checks renders a cancelled job as fail. My watcher reported 16 failures and the operator's reported 20; the true count both times was zero. Only the run-level conclusion distinguishes the two — gh run view <id> --json status,conclusion. Whatever the current run reports is the CI result; nothing here claims one.

omitted_gates — blocked by the GPU LOCK, not the host

Two substitutes were run by the implementer before the operator's compile existed; neither is offered as the device gate:

1. .cuh compile + arity check. Compiled at -std=c++20 -Wall -Wextra -Werror against CUDA shims, with each Kernel<<<cfg>>>(args) rewritten to M2Sink(cfg), Kernel(args) — dropping the launch config while preserving the argument-count and type check on all 7 launches. EXIT 0. Proved armed by deleting the S argument the F2 clamp added, on a scratch copy: too few arguments to function M2ChunkScanKernel, exit 1. Superseded by the operator's real nvcc build above.
2. F2 clamp CPU twin. The two index computations walked over the device case's own shape. Unclamped, every claimed hole reproduced: index 2432 and −4096 into a 2048-element passed, 463232 into a 2048-element initial_states, −512 for the c == 0 hole. Clamped, all land in [0, 1920], the lci clamp reproduces the in-contract index range exactly, and both seq_idx violations read no previous state at all. Not superseded by any compile.

Still owed — and the blocker is now $HOME/gpu.lock, not the machine. The host is back; the lock is held by other coordinators' jobs with 8-hour timeouts, and running against a contended GPU would reproduce the undetected-contention defect §8.4 already records.

1. the three suites' CUDA arms (Release and Debug) — execution, which a compile does not supply;
2. compute-sanitizer memcheck on the new clamp case. This is what actually proves F2: an out-of-bounds write into a cudaMallocAsync pool commonly does not fault, so a green run is necessary and not sufficient;
3. the 9-mutation re-sweep against the moved bound 5·(K+2)·u — §8.4's sweep was scored against 4·(K+2)·u, and a widened bound is precisely the change that could stop a mutation reddening. The re-scaled margins make it very likely to hold; very likely is not a result;
4. §8.4's ~/w2ssd/refail.log, still REMOTE_UNVERIFIED.

The two Windows CI jobs are the main baseline established in §8.4 (#512 fixed by #583, #514, #584) — subtracted, not inherited.

🤖 Generated with Claude Code

[localai-bot]

Superseded by #675, which consolidates this branch with the evidence commits from #566 and lands the row.