W2 (KERNEL-SSM-MAMBA): the Mamba2 SSD CUDA arm — 9/9 mutations caught against the corrected bound (#496) - #675
Merged
Merged
Conversation
…ps (#496) W2 of .agents/specs/mamba2-ssd.md. These are the failing tests, committed before the kernels they gate, per the implementer contract. The three suites gain a `#ifdef VLLM_CPP_CUDA` section that runs the SAME inputs through the device arm. They fail for the intended reason: no NATIVE kernel is registered for kMamba2ChunkScan / kMamba2StateUpdate / kRmsNormGatedGroup on DeviceType::kCUDA. That reason is NOT "GetOp throws", and the difference is the whole point of one assertion in these suites. GB10 is `integrated && pageable_memory_access` (cuda_backend.cu Registrar), so `Backend::UnifiedMemory()` is TRUE and `ReferenceTierEligible(kCUDA)` with it. On a GetOp miss the provider seam does not throw: it installs the CPU HOST kernel as a `kReferenceProviderName` provider and runs THAT over the device pointers (op_provider.h, "portable reference tier"). Every numeric assertion in a device arm would then pass while nothing ran on the GPU -- the device arm gated by running the host arm twice. So every CUDA case calls `RequireNativeCudaProvider`, which reads `GetOpProviderStats(op, kCUDA).last_selected` and refuses `vt-cpu-ref`. These are EAGER dispatches, not a captured graph, so the counter is genuinely populated ([[graph-replay-does-no-host-dispatch-counters-read-zero]]). The declared equivalence contract is written down here BEFORE the kernel, in the head comment of the SSD suite's CUDA section: * the CUDA arm keeps f32 accumulation throughout and does NOT mirror the tile downcasts in upstream's Triton dots (ssd_chunk_state.py:283-285, ssd_chunk_scan.py:266-269, :359-363) -- those are the input-precision requirement of `tl.dot`, i.e. of a tensor-core MMA, and every one of those tiles is loaded `.to(tl.float32)` and computed in f32 right up to the MMA. The memory format is unchanged, so this is not a "too wide" dtype; * G1, the primary gate, is the device output against the SAME independent double-precision sequential reference at the SAME upstream-ported tolerances the host arm is held to; * G2, device-vs-host, is a DERIVED bound: `rtol(K) = 4*(K + 2)*2^-24` over a recurrence of length K. CUDA's `expf` is documented to <= 2 ulp and glibc's to <= 0.5, so a product of K decay factors carries <= 2.5*K*u of libm disagreement, and the length-K f32 summation adds the standard (K-1)*u -- 3.5*K*u, rounded up to integers. Everything else is held identical by construction: each device output element is accumulated in ONE thread over the host arm's index range in the host arm's direction, so summation order is not a second source. A BYTE COMPARE IS NOT REACHABLE, and the libm difference is exactly why. The slack actually used is REPORTED on every comparison, so a bar that stopped doing work would be visible rather than silently absorbing a defect. Also lands the mutation-proof §8.2 records as owed: the decode kernel's `CheckMamba2ANegative` at cpu_ops.cpp:1877 was pinned by NO test -- deleting it left test_ops_mamba2_state_update fully green while the same deletion on its chunk-scan twin reds. The new "A must be negative" SUBCASE mirrors test_ops_mamba2_ssd.cpp:900 and additionally pins that the guard is a SIGN test, not an accidental magnitude floor (A = -1e-30 is accepted). FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…he gate host after a worktree loss (#496) FOLLOWING_AGENTS_PROTOCOL NOT AUTHORED BY THE COMMITTER. This is an operator recovery of a fresh implementer's work after the fourth external deletion of an isolation worktree this session. The implementer had built and run this green on both boxes and staged it; the worktree was removed before the commit. The bytes survived on the gate host and are restored here unchanged: md5 cbb1f928f4 for cuda_mamba2_ssd.cuh and 37a0404433 for cuda_gdn.cu, matching what the implementer reported before the loss. The declared equivalence contract, which the implementer decided BEFORE writing the kernel and recorded in the kernel header and all three test headers: The CUDA arm keeps f32 accumulation throughout and deliberately does NOT mirror upstream's tile downcasts. Those casts -- b.to(x_ptr.dtype.element_ty) at ssd_chunk_state.py:283-285, cb.to(...)/prev_states.to(...) at ssd_chunk_scan.py:266-269,359-363 -- are the input-precision requirement of tl.dot, a tensor-core MMA. Every one of those tiles is loaded .to(tl.float32) and computed in f32 right up to the MMA. These are scalar-FMA kernels with no MMA, so mirroring the downcast would copy a constraint we do not have. The inter-chunk `passed` buffer is allocated at state_dtype, NOT the host arm's f32 working width that spec 8.2 F9 warned W2 must not inherit. A byte compare against the host arm is NOT reachable, and the downcasts are not why: the two arms call different libms (CUDA expf <= 2 ulp, glibc <= 0.5). Everything else is identical by construction. So the primary gate is the device output against the same double-precision sequential reference at the same upstream-ported tolerances the host arm uses, on the same inputs -- which separates "device defect" from "wrong threshold". The derived device-vs-host bar is rtol(K) = 4*(K+2)*2^-24, derived from 2.5 ulp of libm disagreement per decay factor through a product of at most K plus (K-1)*u summation error. No number was tuned and no tolerance was widened; each comparison logs the fraction of budget actually used through MESSAGE rather than INFO, because doctest prints INFO only on failure and an unaudited bar would have been a false claim. Evidence already captured on the gate host: Release build for 121a with CUTLASS 4.5.0, fa2 ENABLED and Marlin NVFP4 enabled, 0 warnings; RED run SIGSEGV on all three binaries; GREEN run ssd 11 cases / 2069 assertions, state_update 10 / 5965, gated_norm 12 / 3723, all Status SUCCESS, exit 0, with zero reference-tier lines. That RED SIGSEGV is a real shared-seam defect, filed as #547 and deliberately not fixed in flow: GB10 reports Backend::UnifiedMemory() == true, so ReferenceTierEligible(kCUDA) is true, and with no native kernel GetOp installs the CPU host kernel as a vt-cpu-ref provider and runs it over cudaMalloc pointers. include/vt/backend.h already says a cudaMalloc pointer is not host-dereferenceable on GB10; op_provider.cpp:515-526 gates on UnifiedMemory() where it needs DeviceMemoryIsHostAddressable(). Every CUDA case here now calls RequireNativeCudaProvider, so a device arm can never be gated by running the host arm twice. STILL OWED, and this branch is NOT landable until a fresh implementer finishes it: the 8 scripted CUDA mutations, compute-sanitizer, the Debug arm, a full ctest on the gate host, the spec section 8.3 that records the contract above (its only copy was a staged blob and is presumed lost), and an origin/main re-merge. A fresh review follows that, not this commit. Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…e $8.2 decode SUBCASE (#496) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…aught, sanitizer clean, Debug arm green (#496) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…tributed failure (#496) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
… -- a second job locks /tmp/gpu.lock (#496) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…e attribution re-run is REMOTE_UNVERIFIED (#496) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…, and the memory-safety claim outran the kernels (#496) Tightening pass on PR #566 (`row/KERNEL-SSM-MAMBA-SSD-W2-FINISH` @ `1e819144e`), whose fresh review returned PASS with five findings. The review's own work -- the equivalence contract's MMA reasoning, the reformulated mutations on CPU twins, the grep, the recovered-byte md5s, the CI baseline subtraction -- is not re-done here. F1 (MEDIUM) -- the derivation, not the number, was wrong. `cuda_mamba2_ssd.cuh` and the test headers both claimed the elementary functions were the ONLY admitted source of device-vs-host divergence. They are not. Host C++ is pinned `-ffp-contract=off` (CMakeLists.txt:55) so `a*b + c` keeps two roundings; nothing passes `--fmad=false` to nvcc, so the header 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 MEASURED that idiom already (.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, which the old constant did not survive: libm 2.5*K*u (<= 2.5 ulp per decay factor, CUDA expf <= 2, glibc <= 0.5) 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 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, the worst of 55 comparisons goes 7.66% -> 6.13% of budget, the driver shapes 0.32%/0.18% -> 0.26%/0.14%, and mutant M3 962173% -> 769738%, still caught by four orders of magnitude. The bound moved because the DERIVATION gained a term the build actually emits, and the header comment, all three test comments and spec §8.3 now say the same true thing. F2 (MEDIUM) -- both halves taken: the free 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 of them. Three are memory-unsafe, and the stated reason for dropping them -- a D2H plus a stream sync -- does not apply, because the values are already in device registers and the decode kernel has always clamped its `state_indices` slot for free on that basis. * `M2StatePassKernel`: `chunk_end` clamped 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` additionally makes `si == si_prev` at c == 0 and indexes `passed` at chunk -1 -- a hole the finding did not name and this pass 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. The claim is narrowed at the same time, because 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 -- and the header and §8.3 now say so instead of folding it into 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 and the mutant is memory-unsafe, failing partly for that rather than 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 sentence 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 in one spec gave two causes for 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, and `Release()` frees all five before reporting rather than leaking the remainder on a mid-sequence free failure. EVIDENCE. `df -h /` 85% used / 67G free before and after every result. test_ops_mamba2_ssd 8/8 1175/1175 SUCCESS! test_ops_mamba2_state_update 6/6 2469/2469 SUCCESS! test_ops_mamba2_gated_norm 9/9 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. OMITTED_GATES -- the CUDA arm was neither built nor run. `dgx.casa` has been unreachable since 06:50 CEST and this box has no nvcc and no GPU. Two substitutes were run and neither is offered as the device gate: 1. The `.cuh` 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 arity and type check on all 7 launches. EXIT 0. Proved ARMED by deleting the `S` argument on a scratch copy: `too few arguments to function M2ChunkScanKernel`, exit 1. 2. A CPU twin of the two clamped index computations, 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. Still owed on device: the three CUDA arms, `compute-sanitizer memcheck` on the new case (which is what actually proves memory safety -- an out-of-bounds write into a pool allocation commonly does not fault), a mutation re-sweep against the moved bound, and §8.4's `~/w2ssd/refail.log`, still REMOTE_UNVERIFIED. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…g pass (#496) Clean Release rebuild of every target -- 817/817 ninja edges, 0 warnings, 0 errors at `-Wall -Wextra -Werror` -- then full `ctest -j 4`: 100% tests passed, 0 failed out of 403, CTEST_EXIT=0, 22.35 s, 2 skipped. `df -h /` 87% used / 60G free. Recorded with the caveat that matters: this is the CPU-ONLY lane, so it is a much smaller and faster gate than §8.4's 431-test GPU-host run and is not a substitute for it. None of §8.4's ten failures is reachable from a build with no CUDA. The device arm remains `omitted_gates` while `dgx.casa` is unreachable. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…D bound (#496) §8.4 reports "the worst one used 7.66% of rtol(K) = 4*(K+2)*2^-24". That is still what the run produced, but §8.5 moved the bound to 5*(K+2)*2^-24, so read without a pointer it now looks like a statement about the current bar. Cross-referenced rather than restated: the captured numbers stay as captured, with the conversion (6.13%, 0.26%/0.14%, 769738%) named next to them and the arithmetic in §8.3 point 6. Rewriting a measurement to match a later derivation is how a record stops being evidence. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…ed the gap it does not cover (#496) Two inaccuracies in the comment written one commit ago, found re-reading it: It said "two of the dropped ones are memory-UNSAFE", which reads as "and the other four are not". The header enumerates six, of which the `cu_chunk_seqlens` tiling and per-chunk length checks are ALSO memory-unsafe and are NOT clamped. The two this case covers are the two that are memory-unsafe AND bounded by a register-local clamp. Narrowing the count in the very comment that exists to stop a claim outrunning the kernels would have re-introduced F2 at a smaller scale, so the comment now points at the header's full list and names the uncovered gap explicitly. It also cited §8.4 for the owed `compute-sanitizer memcheck`; that is recorded in §8.5. Comment-only. Rebuilt clean (392/392, 0 warnings) and re-ran: 8/8, 1175/1175, `Status: SUCCESS!`. The `#ifdef VLLM_CPP_CUDA` region was re-checked with a `-DVLLM_CPP_CUDA -fsyntax-only` compile at -Wall -Wextra -Werror, exit 0. `df -h /` 83% used / 73G free. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
… reproduced standalone on an idle box (#496) FOLLOWING_AGENTS_PROTOCOL The last unattributed ctest failure is closed, and the contention hypothesis this spec had recorded as "plausible, not proven" is REFUTED rather than quietly kept: test_minimax_h3 reproduces STANDALONE, serially, under the lock, on a box that had just rebooted and had no CUDA process resident. It is #486 (open) with #516 as its root cause -- which is what row/pool-device-key is repairing, and therefore why that baseline branch passed a test main-based branches fail. The difference was the baseline carrying a FIX, not this branch carrying a defect. The re-run reproduces #486's recorded signature number for number: 38 | 36 passed | 2 failed | 41 skipped, 42724 assertions, the same two case names, the same `cudaFree: invalid argument`, the same SIGSEGV. All ten ctest failures are now pre-existing and tracked. None is W2's. Also fixed a live protocol defect found while doing it: the re-run had been relaunched as `flock -w 3600 $HOME/gpu.lock ./w2refail.sh`, an OUTER flock wrapping a script that takes the SAME lock on its own fd. flock locks an open file DESCRIPTION, so the inner acquisition blocked against its own parent -- a self-deadlock HOLDING the shared GPU mutex with three other agents' jobs queued behind it. Killed the stack, the lock passed straight to a waiter, and relaunched with a single acquisition. Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…- deterministic, not starved (#496) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…vidence; 32G build tree reclaimed (#496) FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…erator-run on dgx, not CI (#496) Two corrections to §8.5, both narrowing what this branch claims. 1. THE CUDA COMPILE AXIS IS CLOSED, AND IT WAS NOT MY RESULT. `dgx.casa` returned at 06:57 UTC after a reboot, which supersedes §8.4's REMOTE_UNVERIFIED note on the host being away. The operator transferred this branch by `git archive` -- never rsync, which has previously overwritten goldens into false passes -- 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 the project's -Werror flags, disk unchanged at 64G either side so this is not the stale-binary false green, and the three fast-path features READ OUT OF the configure log rather than assumed -- an absent CUTLASS exits 0 too, so "the build succeeded" alone proves nothing. That retires the risk my shim arity-check could only approximate: the new `S` kernel parameter and the `M2Scratch` guard compile through `cuda_gdn.cu` at the arch this ships on. Recorded as OPERATOR-RUN. I did not run it, this box has no nvcc, and `cuda-fat-build` has still never completed on this branch -- so no CI job has compiled this code either. The shim check and the F2 clamp CPU twin are kept as what an implementer could establish unaided, not restated as the gate. 2. THE REMAINING GATES ARE BLOCKED BY THE GPU LOCK, NOT BY THE HOST. The old wording said "omitted_gates until dgx.casa returns". It has returned, and they are still owed -- `$HOME/gpu.lock` is held by other coordinators' jobs with 8h timeouts. A reachable host is not an available GPU, and running these against a contended one reproduces exactly the undetected-contention defect §8.4 already records. Owed: the three CUDA arms (EXECUTION, which a compile does not supply); `compute-sanitizer memcheck` on the new clamp case, which is what actually proves F2 because an out-of-bounds write into a cudaMallocAsync pool commonly does not fault; the 9-mutation re-sweep against the MOVED bound, since a widened bound is precisely the change that could stop a mutation reddening; and §8.4's refail.log. Item 3 is the one a reader is most likely to wave through. The re-scaled margins in §8.3 point 6 make it very likely to hold, and very likely is not a result. 3. THE CANCELLED-vs-FAILED TRAP, WHICH TWO OF US HIT FROM OPPOSITE ENDS. Every Actions run for #592 -- 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` dying mid-build after passing two steps. That is an Actions-side event, not a verdict on any 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. A per-check listing cannot separate "this branch is red" from "the pool was killed" -- only the run-level `conclusion` can. Carried in the spec because anyone subtracting a CI baseline on this repo will otherwise attribute an infrastructure event to a diff. A re-run was triggered at 07:47 and was still queued at 07:50. Whatever it reports is the CI result; this commit does not claim one. Record-only; no code change. `df -h /` 84% used / 73G free. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
FOLLOWING_AGENTS_PROTOCOL Neither W2 branch contained the other. W2-FINISH carries the three evidence commits that resolve the merge precondition -- test_minimax_h3 ATTRIBUTED to #486/#516, all ten ctest failures reproducing STANDALONE, and the evidence-log location. W2-FIX carries the F1/F2 repair and the operator-run nvcc compile. Both are needed; this consolidates them so one squash lands the row. Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
FOLLOWING_AGENTS_PROTOCOL Routine re-merge before landing. `git merge --no-edit` writes git's default message with no trailers, which check-commit-trailers rejects -- it exempts nothing, merges included. Amended rather than left, because the same trailer-less merge commit on a sibling branch is exactly what reddens the push lane when a branch is landed with a local --no-ff instead of a squash. Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…ement (#496) FOLLOWING_AGENTS_PROTOCOL Spec 8.6. The four items 8.5 listed as owed are closed -- three by measurement on the gate host, one by attribution. The CUDA arms and compute-sanitizer ran under one lock acquisition: ssd 12/2095, state_update 10/5965, gated_norm 12/3723, all SUCCESS, memcheck ERROR SUMMARY 0 errors on both. The assertion counts are quoted because a suite that skipped every CUDA case would also exit 0; a non-zero count is what separates "the device arms ran" from "the binary started". The re-sweep ran against the MOVED bound, 5*(K+2)*u, which is the whole point: 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. Controls first at the same counts as the arms, so an empty run is excluded. 9 of 9 CAUGHT. Every mutant built in a phase that takes no lock, run under one acquisition, source restored and md5-verified after each. Every form keeps every variable read, because a mutation that fails to compile under -Werror=all-warnings is not a caught mutation, it is a suite that never ran -- two of the original eight were exactly that. M6 is the one worth reading: it aborts while printing "assertions: 2577 | 2577 passed | 0 failed", a clean assertions line on a FAILING run, and is caught only because the harness reads the exit code. That is this repo's recurring trap in its sharpest form, and it is why item 3 could not have been discharged by inspecting summaries. test_minimax_h3 is ATTRIBUTED, not waived: reproduced standalone on an idle box under the lock, TEST_EXIT=139, and it is #486 with root cause #516, signature-for-signature. The independent baseline that PASSED was row/pool-device-key, the branch that FIXES #516 -- so the baseline carried a fix, this branch does not carry a defect. The contention hypothesis 8.4 recorded as "plausible, not proven" is REFUTED rather than quietly retained. Records only, no product code touched. Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
FOLLOWING_AGENTS_PROTOCOL main moved twice more during the gate run. Merged with an explicit message because `git merge --no-edit` writes git's default, which check-commit-trailers rejects -- it exempts nothing, merges included. Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
This was referenced Aug 13, 2026
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Completes W2 of
.agents/specs/mamba2-ssd.md(#496) — the CUDA arm of the three Mamba2 SSD ops. Supersedes #566 and #592, which had diverged:W2-FINISHcarried the evidence commits that resolve the merge precondition,W2-FIXcarried the F1/F2 repair. Neither contained the other; this consolidates both so one squash lands the row.What the fresh reviews found, and what the repair did
F1 — the derived bound omitted nvcc FMA contraction. Host C++ is pinned
-ffp-contract=off; nothing passes--fmad=falseto nvcc, soacc += xv * bvcontracts on device and not on host. The stated model was3.5·K·u, contraction adds~1.0·K·u, and the bound was4·(K+2)·u— provable only forK ≤ 18, while the driver shapes run at K = 200. Repaired by carrying the term:5·(K+2)·u, arithmetic shown in the header, all three test comments and §8.3.-fmad=falsewas rejected deliberately — it is a per-TU flag on a header included by the hot GDN TU. The repo-wide gap is filed as #591.F2 — the memory-safety claim outran the kernels. The validator checks metadata shape/dtype only; every value check lived in the CPU kernel. Two dropped checks were memory-unsafe: an out-of-bounds write past a
cudaMallocAsyncallocation and an out-of-bounds read ofinitial_states. Repaired with both the free register-local clamps and a narrowed claim enumerating all six dropped checks. A third hole the review did not name was found and closed:seq_idx[0] < 0indexespassedat chunk −1.Gates — all four owed items discharged, operator-run
sm_121aCONFIGURE_EXIT=0 BUILD_EXIT=0 WARNINGS=0 ENOSPC=0, CUTLASS/FA2/Marlin confirmed enabled in the configure logSUCCESS!compute-sanitizer memcheckERROR SUMMARY: 0 errors, both suites5·(K+2)·uAssertion counts are quoted because a suite that skipped every CUDA case would also exit 0.
M6 is why the re-sweep could not be waived. It aborts (
exit=134) while printingassertions: 2577 | 2577 passed | 0 failed— a clean assertions line on a failing run — and is caught only because the harness reads the exit code. The original sweep was scored against the old bound, and a widened tolerance is precisely what stops a mutation reddening; "very likely to hold" is not a result.Every mutant keeps every variable read, because a mutation that fails to compile under
-Werror=all-warningsis not a caught mutation but a suite that never ran — two of the original eight were exactly that.test_minimax_h3— attributed, not waivedReproduced standalone on an idle box under the lock (
TEST_EXIT=139,ENOSPC=0). It is #486, root cause #516, signature-for-signature. The independent baseline that passed wasrow/pool-device-key— the branch that fixes #516. The baseline carried a fix; this branch does not carry a defect. The contention hypothesis is refuted, not quietly retained.Not claimed
CI is
REMOTE_UNVERIFIED: every run on the predecessor branches endedcancelled, including a repo-wide mass cancellation of 20 runs across 7 branches.cuda-fat-buildnever completed, which is why the compile was run directly on the gate host instead. Windows reds are themainbaseline (#514, #584).🤖 Generated with Claude Code