From eb2810872886af4061b48444d1bebdf6c3d3d85d Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 13 Aug 2026 00:06:52 +0000 Subject: [PATCH 1/2] fix(deepseek-v4): gate the DSA top-k window clamps and non-positive topk, attribute the launch error, and stop a guard comment overclaiming (#552) A fresh review of #505 returned PASS with 6 non-blocking findings. It could not find a single input where the two-pass threshold selection diverges from `DsaTopkSelect` -- 3,000,081 fuzzed shapes across three independent implementations, zero divergence -- and reproduced both the defect and the fix on real sm_121a. This lands the six findings. CODE 1. `DsaTopkLaunch` had no post-launch `cudaGetLastError()` while every sibling launcher here has one. That is exactly why the #505 stack overflow surfaced as `cudaStreamDestroy: an illegal memory access` -- an error raised by whatever ran next, and the text the whole #505 evidence chain ended up quoting. Now `Check(cudaGetLastError(), "dsa_topk launch")`, so the next fault in this kernel is attributable to it. 2. `topk <= 0` diverged between the arms: the host reference asserts (`deepseek_v4_dsa.cpp:76`) while the device launcher silently returned an empty vector. The launcher now throws to match. The "two independent implementations agree" gate could not see this because no case passed a non-positive topk. TESTS 3. No case exercised `win_start < 0` or `win_end > num_keys`, so mutating either clamp away left all four #505 cases green while an off-device fuzz caught both at once; on device they are out-of-bounds `logits` reads. New case drives under-run, over-run, both-at-once, and an in-range control at topk=512, and asserts no emitted key escapes the clamped window. 4. New case asserts BOTH arms refuse topk=0 and topk=-1, not merely that they agree on accepted inputs. PROSE -- these are corrections to claims that were wrong, not polish 5. The `w < topk` comment named a NaN write-past-row as the failure class it defends. That cannot occur: pass 2's predicate is satisfied by exactly `rank(th)` elements with `rank(th) <= topk`, and a NaN never satisfies it either since `better(x, NaN)` is false, so `w` cannot exceed topk for ANY input. The reviewer proved it by removing the bound and by weakening it to `w <= topk` -- both left the device suite 4/4 SUCCESS and a 3M-shape fuzz clean. The guards stay because they are free; the comment now says belt-and-braces, bounded by construction. A comment that overstates what it protects is worse than none, because the next reader trusts it. 6. "Strictly cheaper" is now qualified as a comparison to the OLD kernel only, not fitness for the real geometry: this is one thread per token row with a dependent global load per iteration, so at V4-Pro's index_topk=1024 the last row of a 4k prompt is ~4.2M serial loads in one thread. 7. Trivia: the spec cited the pre-fix span as `:624-665` (actual 624-669), and the kernel comment called `n > topk` the overflow condition -- it only selects the branch; the overflow needs `n > 512` for `chosen` or `topk > 64` for `picked`, which is why the old topk=3/nk=5 gate shape was ASan-clean. Spec section 7 also records what the review CLOSED rather than raised: five mutations left the suite green and are semantics-preserving, not coverage gaps -- each is also undetected by a 200k-shape fuzz against the oracle, which is how the reviewer separated the two. And it confirms the tie-heavy case is uniquely load-bearing: it is the only case that reddens a tie-break inversion or a value-only threshold. Recorded, not repaired: `b649a1ea2` introduced the #505 spec in the same commit as its code, where the protocol requires the spec first. Gate-reading trap worth keeping, from the reviewer's own run: `-ts='*DSA top-k*'` (suite filter) instead of `-tc=` (case filter) printed `0 passed | 0 failed | 23 skipped` beside `Status: SUCCESS!` -- a live false green from a filter matching nothing. Any scoped run of this suite must use `-tc=` and check the case count. CPU: test_cuda_deepseek_v4 25/25 (was 23, both new cases registered), test_deepseek_v4_dsa 13/13, both SUCCESS. Device arms on dgx.casa are queued behind other agents' jobs on the shared GPU lock; results land before merge. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --- .agents/roadmap_v1.md | 1 + .agents/specs/dsa-topk-bounds.md | 79 ++++++++++++++++++--- src/vt/cuda/cuda_deepseek_v4.cu | 38 +++++++--- tests/vllm/models/test_cuda_deepseek_v4.cpp | 50 +++++++++++++ 4 files changed, 150 insertions(+), 18 deletions(-) diff --git a/.agents/roadmap_v1.md b/.agents/roadmap_v1.md index e9e3a4d37..8beb9e305 100644 --- a/.agents/roadmap_v1.md +++ b/.agents/roadmap_v1.md @@ -156,6 +156,7 @@ issue is not yet placed. Keyed record: update in place, never append. | [#408](https://github.com/mudler/vllm.cpp/issues/408) | — | 12 of 54 `tests/scripts` suites are executed by nothing, and `check-test-registration.py`'s fixed `REQUIRED_TESTS` cannot see the class (found while repairing #274) | bug | | [#504](https://github.com/mudler/vllm.cpp/issues/504) | `MODEL-TEXT-deepseek-v4-deepseek-v4-for-causal-lm` | DeepSeek-V4-Pro is the same architecture as V4-Flash (zero new config keys): record the variant and gate the config descent's shape-generality (spec `specs/deepseek-v4-pro.md`) | feature | | [#505](https://github.com/mudler/vllm.cpp/issues/505) | `MODEL-TEXT-deepseek-v4-deepseek-v4-for-causal-lm` | `DsaTopkKernel` sizes `chosen[512]`/`picked[64]` by literal while `index_topk` is 512 (Flash) / 1024 (Pro); latent behind `dsa_dense` today, silent stack overflow once the real-geometry DSA residual lands (found while assessing #504) | bug | +| [#552](https://github.com/mudler/vllm.cpp/issues/552) | `MODEL-TEXT-deepseek-v4-deepseek-v4-for-causal-lm` | DSA top-k review findings: the `w < topk` guard comment overclaims what it defends, the window clamps and non-positive `topk` are ungated, and `DsaTopkLaunch` swallows its launch error (spec `specs/dsa-topk-bounds.md` §7) | bug | | [#469](https://github.com/mudler/vllm.cpp/issues/469) | — | `test_ops_glue.cpp:190`'s `CHECK_THROWS` is satisfied by the CPU kernel's second guard, not the dispatch guard it names — mutation M8 survives. Behavior is correct; test strength only | bug | ## Top-level portfolio diff --git a/.agents/specs/dsa-topk-bounds.md b/.agents/specs/dsa-topk-bounds.md index 0717a14a7..877c3df9a 100644 --- a/.agents/specs/dsa-topk-bounds.md +++ b/.agents/specs/dsa-topk-bounds.md @@ -10,7 +10,7 @@ ## 0. Scope -`DsaTopkKernel` (`src/vt/cuda/cuda_deepseek_v4.cu:624-665` pre-fix) sized two +`DsaTopkKernel` (`src/vt/cuda/cuda_deepseek_v4.cu:624-669` pre-fix) sized two thread-local arrays by literal: ```cpp @@ -68,14 +68,27 @@ asc — a total order because candidate indices are distinct): Exactly `topk` elements satisfy pass 2 under a total order, and they come out already in ascending key order, so the `O(topk^2)` emit sort disappears along with the buffers. **No per-thread scratch, no bound, no configurable limit.** Cost is -unchanged at `O(topk*n)` for pass 1 and strictly better overall. - -Two defensive additions that are not load-bearing for ordered input: pass 1 stops -if no strictly-worse element is found, and pass 2 carries a `w < topk` bound. -Both exist so a NaN row — where every float comparison is false — cannot write -past the thread's own row into the next one, which is the failure class this issue -was about. The host reference is naturally immune (it resizes to `topk`), so this -keeps the two arms equally safe rather than mirroring a weakness. +unchanged at `O(topk*n)` for pass 1 and cheaper than the revision it replaces, +which also paid the emit sort. + +That comparison is to the old kernel **only**, and is not a claim of fitness for +the real geometry (corrected per the review, #552 finding 5). Pass 1 is one thread +per token row with a dependent global load per iteration, so at V4-Pro's +`index_topk=1024` the last row of a 4k prompt is on the order of 4.2M serial +loads in a single thread. Asymptotically unchanged from the pre-fix kernel, and +out of scope here — §5 names the real-geometry path as the residual that owns it. + +Two defensive additions, neither load-bearing: pass 1 stops if no strictly-worse +element is found, and pass 2 carries a `w < topk` bound. **They are +belt-and-braces, not a defence against a known way to overrun the row** — the +emit count is bounded by construction. Pass 2's predicate is satisfied by exactly +`rank(th)` elements and `rank(th) <= topk`, and a NaN never satisfies it either +because `better(x, NaN)` is false, so `w` cannot exceed `topk` for any input. The +review proved this by removing the bound and by weakening it to `w <= topk`: both +left the device suite 4/4 SUCCESS and a 3M-shape fuzz clean. The original wording +here named a NaN write-past-row as the defended failure class, which was an +overclaim (#552 finding 1). The guards stay — they are free — but a comment that +overstates what it protects is worse than none, because the next reader trusts it. ## 4. Evidence @@ -185,6 +198,54 @@ kernel, not against our host reference. That work stays out of scope here. - Do **not** widen scope into the real-geometry DSA residual or the compressed-key-space candidate window (§5). +## 7. Review follow-up — [#552](https://github.com/mudler/vllm.cpp/issues/552) + +A fresh reviewer (never the author) returned **PASS with 6 non-blocking +findings**, having failed to find any input where the two-pass selection diverges +from `DsaTopkSelect`: **3,000,081 fuzzed shapes across three independent +implementations** — the host reference, a transcription of the kernel body, and +the reviewer's own `O(n^2)` rank-count oracle so a shared misreading could not +pass — with **zero divergence**, output slots poisoned with `-777` so an unwritten +slot could not masquerade as legitimate `-1` padding. It reproduced the ASan +overflow and the device SIGABRT independently, and ran a 12-row device mutation +table on real sm_121a. + +Landed here: + +| # | finding | resolution | +|---|---|---| +| 1 | the `w < topk` comment named a NaN write-past-row as the defended failure class, which cannot occur | reworded in the kernel and in §3: belt-and-braces, bounded by construction. Guards kept. | +| 2 | no case exercised `win_start < 0` or `win_end > num_keys`; mutating either clamp away left all four cases green | new case `OUT-OF-RANGE windows are clamped like host (#552)`, four rows driving under-run, over-run, both, and an in-range control | +| 3 | `topk <= 0`: host asserts, device silently returned an empty vector | launcher now throws to mirror `deepseek_v4_dsa.cpp:76`; new case asserts **both** arms refuse | +| 4 | `DsaTopkLaunch` had no post-launch `cudaGetLastError()`, which is why the #505 fault surfaced as `cudaStreamDestroy` | `Check(cudaGetLastError(), "dsa_topk launch")` added, matching every sibling launcher | +| 5 | "strictly cheaper" could be read as real-geometry-ready | qualified in §3 and in the kernel comment: it is a comparison to the old kernel only | +| 6 | trivia: pre-fix span cited `:624-665` (actual 624-669); the comment described `n > topk` as the overflow condition | both corrected; `n > topk` only selects the branch, the overflow needs `n > 512` or `topk > 64` | + +**Findings the review closed rather than raised.** Five mutations left the suite +green and are *semantics-preserving*, not coverage gaps — each is also undetected +by a 200k-shape fuzz against the oracle, which is how the reviewer separated the +two: removing `w < topk`, weakening it to `w <= topk`, simplifying pass 1's argmax +to `v > best_val` (equivalent, since the scan is ascending), `n <= topk` → +`n < topk` (equivalent, since at `n == topk` the full path selects all `n`), and +dropping `if (th_idx < 0) return` (reachable only via NaN). No action owed. + +**The tie-heavy case is uniquely load-bearing** — confirmed, not assumed. It is +the only case in the suite that reddens a tie-break inversion or a value-only +threshold; the widths and offset cases use distinct random logits and stay green +under both mutations. + +**Process deviation recorded, not repaired.** `b649a1ea2` introduced this spec in +the *same* commit as the code, where the protocol requires the spec to be +committed first. The substance was complete and the review supplies the +independent pass that #542 lacked, but the ordering was wrong and is noted here +rather than quietly dropped. + +**Gate-reading trap, from the reviewer's own run.** A first scoped attempt used +`-ts='*DSA top-k*'` (suite filter) instead of `-tc=` (case filter) and printed +`test cases: 0 | 0 passed | 0 failed | 23 skipped` beside `Status: SUCCESS!` — a +live false green from a filter that matched nothing. Any scoped run of this suite +must use `-tc=` and confirm the case count is non-zero. + ## Outcome **Measured.** The two literal bounds were a real device fault, not a theoretical diff --git a/src/vt/cuda/cuda_deepseek_v4.cu b/src/vt/cuda/cuda_deepseek_v4.cu index 78eaf11fb..d29a00e04 100644 --- a/src/vt/cuda/cuda_deepseek_v4.cu +++ b/src/vt/cuda/cuda_deepseek_v4.cu @@ -649,9 +649,16 @@ __global__ void DsaTopkKernel(const float* logits, const int64_t* ws, const int6 // // This replaces a `bool chosen[512]` + `int64_t picked[64]` pair of literals // that could not represent the real `index_topk` (512 on V4-Flash, 1024 on - // V4-Pro) and overflowed the thread stack on any window wider than `topk` - // (#505). Cost is unchanged at O(topk*n) for pass 1, and strictly better - // overall: the O(topk^2) emit sort is eliminated. + // V4-Pro). `n > topk` only selects THIS branch; the overflow itself needed + // `n > 512` for `chosen` or `topk > 64` for `picked`, which is why the old + // gate shape (topk=3, nk=5) was ASan-clean and the bound stayed invisible + // (#505, #552). + // + // Cost is unchanged at O(topk*n) for pass 1 and cheaper than the revision it + // replaces, which also paid an O(topk^2) emit sort. That is a comparison to the + // old kernel ONLY, not a claim of fitness for the real geometry: this is one + // thread per token row with a dependent global load per iteration, so a real + // long-context row stays expensive. See `.agents/specs/dsa-topk-bounds.md` §5. const int64_t row = static_cast(t) * nk; // `better(va, a, vb, b)` == "(va, a) outranks (vb, b)". auto better = [](float va, int64_t a, float vb, int64_t b) -> bool { @@ -675,16 +682,18 @@ __global__ void DsaTopkKernel(const float* logits, const int64_t* ws, const int6 // n > topk holds here, so a strictly worse element always exists under a // total order. `best < 0` is therefore unreachable on ordered input; it can // only arise if the row carries NaN, which makes every comparison false. Stop - // rather than reset the threshold, so pass 2 still emits a bounded prefix. + // rather than reset the threshold, so the descent cannot restart from the top. if (best < 0) break; th_val = best_val; th_idx = best; } - if (th_idx < 0) return; // pathological row: leave the -1 padding in place - // Exactly `topk` elements outrank-or-equal the threshold, so `w` lands on topk. - // The `w < topk` bound is not load-bearing for ordered input — it is here so a - // NaN row can never write past this thread's row into the next one, which is - // the failure class #505 was about. + if (th_idx < 0) return; // every candidate was NaN: leave the -1 padding + // Exactly `rank(th)` elements outrank-or-equal the threshold and `rank(th)` is + // at most `topk`, so `w` cannot exceed `topk` for ANY input — a NaN never + // satisfies this predicate either, since `better(x, NaN)` is false. The + // `w < topk` bound is therefore belt-and-braces, not a defence against a known + // way to overrun the row: the emit count is bounded by construction (#552; + // removing the bound leaves both the device suite and a 3M-shape fuzz clean). int64_t w = 0; for (int64_t s = s0; s < s1 && w < topk; ++s) { const float v = logits[row + s]; @@ -1179,6 +1188,12 @@ std::vector DsaTopkLaunch(Queue& q, const std::vector& logits, const std::vector& ws, const std::vector& we, int64_t T, int64_t nk, int64_t topk) { + // Mirror the host reference's precondition (`deepseek_v4_dsa.cpp:76` + // `VT_CHECK(topk > 0, ...)`). Without this the device arm silently returned an + // empty vector where the host throws, so the two arms disagreed on a case no + // test covered (#552). + if (topk <= 0) + throw std::runtime_error("vt cuda deepseek_v4: dsa topk: topk must be positive"); cudaStream_t s = AsStream(q); Dev dl = Upload(logits, s), dws = Upload(ws, s), dwe = Upload(we, s); std::vector out(static_cast(T * topk), -1); @@ -1188,6 +1203,11 @@ std::vector DsaTopkLaunch(Queue& q, const std::vector& logits, static_cast(dl.p), static_cast(dws.p), static_cast(dwe.p), static_cast(T), static_cast(nk), static_cast(topk), static_cast(dout.p)); + // Attribute a fault to THIS launch. Without it the #505 stack overflow + // surfaced as `cudaStreamDestroy: an illegal memory access` — an error raised + // by whatever ran next — which is what the whole #505 evidence chain ended up + // quoting. Every sibling launcher here already checks (e.g. `hc_head_ip`). + Check(cudaGetLastError(), "dsa_topk launch"); Download(out, dout.p, s); Check(cudaStreamSynchronize(s), "sync topk"); return out; diff --git a/tests/vllm/models/test_cuda_deepseek_v4.cpp b/tests/vllm/models/test_cuda_deepseek_v4.cpp index 963886530..ace1bfc7e 100644 --- a/tests/vllm/models/test_cuda_deepseek_v4.cpp +++ b/tests/vllm/models/test_cuda_deepseek_v4.cpp @@ -285,6 +285,56 @@ TEST_CASE("W7-device DSA top-k select: OFFSET window at real width matches host CHECK(got[static_cast(t * topk + j)] >= 137); } +// #552 finding 2: the kernel clamps its window with `s0 = ws[t] > 0 ? ws[t] : 0` +// and `s1 = we[t] < nk ? we[t] : nk` (cuda_deepseek_v4.cu:628-629), but no case +// passed a window that NEEDS clamping. Mutating either clamp away left all four +// #505 cases green while an off-device fuzz caught both immediately; on device +// they become out-of-bounds `logits` reads. These rows drive both clamps at a +// real width, against the same host reference. +TEST_CASE("W7-device DSA top-k select: OUT-OF-RANGE windows are clamped like host (#552)") { + if (!HasCuda()) { MESSAGE("no CUDA; skip"); return; } + vt::Backend& gpu = vt::GetBackend(vt::DeviceType::kCUDA); + QueueGuard g(gpu); + Rng r; + const int64_t topk = 512, nk = 700, T = 4; + const auto logits = Rand(r, T * nk, -3.0f, 3.0f); + + // Row 0 under-runs (ws < 0), row 1 over-runs (we > nk), row 2 does both, row 3 + // is in range as the control. Every row still has n > topk after clamping. + std::vector ws{-9, 0, -4, 100}; + std::vector we{nk, nk + 37, nk + 12, nk}; + const auto ref = dv4::DsaTopkSelect(logits, ws, we, T, nk, topk); + const auto got = dv4::DsaDevice()->topk(g.q, logits, ws, we, T, nk, topk); + REQUIRE(got.size() == ref.size()); + for (size_t i = 0; i < ref.size(); ++i) CHECK(got[i] == ref[i]); + + // No emitted key may escape the clamped window, which is what an unclamped + // read would produce. + for (int64_t t = 0; t < T; ++t) + for (int64_t j = 0; j < topk; ++j) { + const int64_t s = got[static_cast(t * topk + j)]; + CHECK(s >= 0); + CHECK(s < nk); + } +} + +// #552 finding 3: `DsaTopkSelect` asserts `topk > 0` (deepseek_v4_dsa.cpp:76) +// while the device launcher used to return an empty vector instead. The two arms +// must refuse the same inputs, not just agree on the accepted ones. +TEST_CASE("W7-device DSA top-k select: non-positive topk REFUSED on both arms (#552)") { + if (!HasCuda()) { MESSAGE("no CUDA; skip"); return; } + vt::Backend& gpu = vt::GetBackend(vt::DeviceType::kCUDA); + QueueGuard g(gpu); + Rng r; + const int64_t nk = 8, T = 2; + const auto logits = Rand(r, T * nk, -1.0f, 1.0f); + std::vector ws(T, 0), we(T, nk); + for (const int64_t bad : {static_cast(0), static_cast(-1)}) { + CHECK_THROWS(dv4::DsaTopkSelect(logits, ws, we, T, nk, bad)); + CHECK_THROWS(dv4::DsaDevice()->topk(g.q, logits, ws, we, T, nk, bad)); + } +} + TEST_CASE("W7-device attention-sink softmax + grouped output-LoRA: CUDA vs host (near-tie)") { if (!HasCuda()) { MESSAGE("no CUDA; skip"); return; } vt::Backend& gpu = vt::GetBackend(vt::DeviceType::kCUDA); From 03c3f639841d80dbfd885ea101d4c7dda8c81ddb Mon Sep 17 00:00:00 2001 From: Ettore Di Giacinto Date: Thu, 13 Aug 2026 07:48:24 +0000 Subject: [PATCH 2/2] fix(deepseek-v4): the launch-error check does NOT attribute an async fault -- device arms refuted finding 4's rationale (#552) Adds the device evidence for this change and corrects a claim I wrote in the previous commit, which the verification arm built to demonstrate it disproved. REFUTED: finding 4 held that a post-launch `cudaGetLastError()` would make the next fault in this kernel attributable to it, instead of surfacing later as `cudaStreamDestroy` the way #505's did. The `prefix_with_check` arm -- the pre-fix #505 kernel body WITH the new check in place -- still reports: what(): vt cuda: cudaStreamDestroy: an illegal memory access was encountered test_cuda_deepseek_v4.cpp:214: FATAL ERROR: test case CRASHED: SIGABRT The error text is UNCHANGED. A stack-overflow illegal access is an asynchronous execution fault; `cudaGetLastError()` right after a launch reports launch-CONFIGURATION errors (bad grid/block, shared memory over budget). Here it was not even caught by the following `cudaStreamSynchronize`, latching only at stream destruction. The check is KEPT -- it is free, it matches every sibling launcher, and it does cover the launch-configuration class. But its comment and the spec now say what it actually does. Claiming "this makes the next fault attributable" would have repeated finding 1's defect -- an overclaiming guard comment -- inside the very change that exists to correct one. DEVICE ARMS (dgx.casa, GB10 sm_121a, mandatory flags, fast path hard-verified in each run's own configure log; every arm a fresh nvcc rebuild from a pristine kernel with the binary mtime verified to advance; scoped with `-tc=`, never `-ts=`; run under `flock $HOME/gpu.lock`): baseline exit=0 6/6 SUCCESS no_topk_guard exit=1 FAILURE -> new refusal case: CHECK_THROWS did NOT throw no_ws_clamp exit=134 FAILURE -> new clamp case CRASHED, illegal memory access no_we_clamp exit=1 FAILURE -> new clamp case, 1052 failed assertions no_launch_check exit=0 SUCCESS -> nothing; unobservable by construction prefix_with_check exit=134 FAILURE -> refutes finding 4's rationale restored full suite exit=0 25/25, 90062 assertions, 0 skipped, SUCCESS Both new cases have teeth against the mutation each was written for. `no_launch_check` reddening nothing is the expected, honest outcome for a diagnostic with no observable behaviour on a passing run. The false-green shape recurred twice more: no_ws_clamp printed `assertions: 16865 | 16865 passed | 0 failed` and prefix_with_check `609 | 609 passed | 0 failed`, both beside `Status: FAILURE!`. Two infrastructure incidents, handled rather than absorbed. Round 1's clamp mutations were REFUSED by their own uniqueness assertion: the clamp pair appears in BOTH `DsaLogitsKernel` (:604-605) and `DsaTopkKernel` (:628-629), so a one-line anchor was ambiguous and the assert declined to edit a kernel it was not aiming at -- round 2 re-ran with a three-line anchor unique to DsaTopkKernel. Then the box REBOOTED mid-arm (`up 7 min`, the known GB10 unified-memory OOM-reboot class), killing the run with no marker written; the remaining arms were relaunched on the fresh box. CPU after the correction: test_cuda_deepseek_v4 25/25 SUCCESS, preflight clean. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code] --- .agents/specs/dsa-topk-bounds.md | 69 +++++++++++++++++++++++++++++++- src/vt/cuda/cuda_deepseek_v4.cu | 12 ++++-- 2 files changed, 76 insertions(+), 5 deletions(-) diff --git a/.agents/specs/dsa-topk-bounds.md b/.agents/specs/dsa-topk-bounds.md index 877c3df9a..a6ce1d8f6 100644 --- a/.agents/specs/dsa-topk-bounds.md +++ b/.agents/specs/dsa-topk-bounds.md @@ -217,10 +217,77 @@ Landed here: | 1 | the `w < topk` comment named a NaN write-past-row as the defended failure class, which cannot occur | reworded in the kernel and in §3: belt-and-braces, bounded by construction. Guards kept. | | 2 | no case exercised `win_start < 0` or `win_end > num_keys`; mutating either clamp away left all four cases green | new case `OUT-OF-RANGE windows are clamped like host (#552)`, four rows driving under-run, over-run, both, and an in-range control | | 3 | `topk <= 0`: host asserts, device silently returned an empty vector | launcher now throws to mirror `deepseek_v4_dsa.cpp:76`; new case asserts **both** arms refuse | -| 4 | `DsaTopkLaunch` had no post-launch `cudaGetLastError()`, which is why the #505 fault surfaced as `cudaStreamDestroy` | `Check(cudaGetLastError(), "dsa_topk launch")` added, matching every sibling launcher | +| 4 | `DsaTopkLaunch` had no post-launch `cudaGetLastError()` | `Check(cudaGetLastError(), "dsa_topk launch")` added, matching every sibling launcher — but the finding's stated *rationale* does not survive measurement, see below | | 5 | "strictly cheaper" could be read as real-geometry-ready | qualified in §3 and in the kernel comment: it is a comparison to the old kernel only | | 6 | trivia: pre-fix span cited `:624-665` (actual 624-669); the comment described `n > topk` as the overflow condition | both corrected; `n > topk` only selects the branch, the overflow needs `n > 512` or `topk > 64` | +### 7.1 Finding 4's rationale was REFUTED by its own verification + +The review's finding 4 held that a post-launch `cudaGetLastError()` would make the +next fault in this kernel attributable to it, instead of surfacing later as +`cudaStreamDestroy` the way #505's did. **That is wrong, and the arm built to +demonstrate it disproved it.** + +Arm `prefix_with_check` — the pre-fix #505 kernel body WITH the new +`Check(cudaGetLastError(), "dsa_topk launch")` in place: + +``` +### prefix_with_check exit=134 +[doctest] test cases: 2 | 1 passed | 1 failed | 23 skipped +[doctest] assertions: 609 | 609 passed | 0 failed | +[doctest] Status: FAILURE! + what(): vt cuda: cudaStreamDestroy: an illegal memory access was encountered +test_cuda_deepseek_v4.cpp:214: FATAL ERROR: test case CRASHED: SIGABRT +``` + +The error text is **unchanged** — still `cudaStreamDestroy`, not `dsa_topk +launch`. A stack-overflow illegal access is an *asynchronous execution* fault; +`cudaGetLastError()` immediately after a launch reports launch-*configuration* +errors (bad grid/block, shared memory over budget). Here it was not even caught by +the following `cudaStreamSynchronize`, and latched only at stream destruction. + +The check is kept: it is free, it matches every sibling launcher, and it does +cover the launch-configuration class. But its comment and this spec now say what +it actually does. Writing "this makes the next fault attributable" would have +repeated finding 1's defect — an overclaiming guard comment — in the very change +that exists to correct one. The lesson is the finding's, not the reviewer's: a +plausible rationale for a cheap, obviously-correct change is still a claim, and +this one only failed because an arm was built to test it rather than to confirm it. + +### 7.2 Device evidence for this change + +`dgx.casa`, GB10 sm_121a, mandatory flags, fast path hard-verified in the run's own +configure log (`CUTLASS found … sm120a NVFP4`, `FlashAttention-2 … [121a]`). Each +arm is a fresh nvcc rebuild from a pristine kernel with the binary mtime verified +to advance, scoped with `-tc=` (never `-ts=`), and run under `flock $HOME/gpu.lock`. + +| arm | exit | Status | reddened | +|---|---|---|---| +| baseline | 0 | `6 \| 6 passed` SUCCESS | — | +| `no_topk_guard` (drop the launcher's `topk > 0` throw) | 1 | FAILURE | the new refusal case: `CHECK_THROWS … did NOT throw at all!` | +| `no_ws_clamp` | **134** | FAILURE | the new clamp case CRASHED, SIGABRT, `illegal memory access` | +| `no_we_clamp` | 1 | FAILURE | the new clamp case, **1052** failed assertions | +| `no_launch_check` | 0 | SUCCESS | **nothing** — see §7.1; unobservable by construction | +| `prefix_with_check` | 134 | FAILURE | refutes finding 4's rationale (§7.1) | +| restored, full suite | 0 | `25 \| 25 passed \| 0 skipped`, 90062 assertions, SUCCESS | — | + +Both new cases therefore have teeth, each against the mutation it was written for. +`no_launch_check` reddening nothing is the expected and honest outcome for a +diagnostic that has no observable behaviour on a passing run. + +Note the false-green shape recurring twice more: `no_ws_clamp` printed +`assertions: 16865 | 16865 passed | 0 failed` and `prefix_with_check` printed +`609 | 609 passed | 0 failed`, both beside `Status: FAILURE!`. + +Two infrastructure incidents, both handled rather than absorbed: round 1's clamp +mutations were **refused by their own uniqueness assertion** — the clamp pair +appears in both `DsaLogitsKernel` (:604-605) and `DsaTopkKernel` (:628-629), so a +one-line anchor was ambiguous and the assert declined to edit a kernel it was not +aiming at; round 2 re-ran with a three-line anchor unique to `DsaTopkKernel`. Then +the box **rebooted mid-arm** (`up 7 min`, the known GB10 unified-memory OOM-reboot +class), killing the run with no marker written — the remaining arms were relaunched +on the fresh box. + **Findings the review closed rather than raised.** Five mutations left the suite green and are *semantics-preserving*, not coverage gaps — each is also undetected by a 200k-shape fuzz against the oracle, which is how the reviewer separated the diff --git a/src/vt/cuda/cuda_deepseek_v4.cu b/src/vt/cuda/cuda_deepseek_v4.cu index d29a00e04..d953d72fa 100644 --- a/src/vt/cuda/cuda_deepseek_v4.cu +++ b/src/vt/cuda/cuda_deepseek_v4.cu @@ -1203,10 +1203,14 @@ std::vector DsaTopkLaunch(Queue& q, const std::vector& logits, static_cast(dl.p), static_cast(dws.p), static_cast(dwe.p), static_cast(T), static_cast(nk), static_cast(topk), static_cast(dout.p)); - // Attribute a fault to THIS launch. Without it the #505 stack overflow - // surfaced as `cudaStreamDestroy: an illegal memory access` — an error raised - // by whatever ran next — which is what the whole #505 evidence chain ended up - // quoting. Every sibling launcher here already checks (e.g. `hc_head_ip`). + // Catches LAUNCH-CONFIGURATION errors (bad grid/block, shared-memory over + // budget) at the call site, and matches every sibling launcher here + // (e.g. `hc_head_ip`). It does NOT catch an asynchronous execution fault: + // MEASURED on GB10 (#552), the pre-fix #505 kernel with this very check in + // place still surfaced its illegal access later, as + // `cudaStreamDestroy: an illegal memory access`. So this does not make a + // device-side fault in the kernel attributable to this launch — a claim an + // earlier revision of this comment made and the measurement refuted. Check(cudaGetLastError(), "dsa_topk launch"); Download(out, dout.p, s); Check(cudaStreamSynchronize(s), "sync topk");