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)

.agents/roadmap_v1.md

@@ -166,6 +166,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

.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,121 @@ 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()` | `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
+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

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<int64_t>(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<int64_t> DsaTopkLaunch(Queue& q, const std::vector<float>& logits,
                                    const std::vector<int64_t>& ws,
                                    const std::vector<int64_t>& 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<int64_t> out(static_cast<size_t>(T * topk), -1);
@@ -1188,6 +1203,15 @@ std::vector<int64_t> DsaTopkLaunch(Queue& q, const std::vector<float>& logits,
       static_cast<const float*>(dl.p), static_cast<const int64_t*>(dws.p),
       static_cast<const int64_t*>(dwe.p), static_cast<int>(T), static_cast<int>(nk),
       static_cast<int>(topk), static_cast<int64_t*>(dout.p));
+  // 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");
   return out;

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<size_t>(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<int64_t> ws{-9, 0, -4, 100};
+  std::vector<int64_t> 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<size_t>(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<int64_t> ws(T, 0), we(T, nk);
+  for (const int64_t bad : {static_cast<int64_t>(0), static_cast<int64_t>(-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);