DSpark: proposal depth is not clamped to the draft's block_size — OOB read, LOSSLESS=0 at depth == block_size

[nickmopen]

Summary

kProposalDepth is clamped to 15 but never against the draft checkpoint's block_size, while
the buffers it indexes are sized by block_size. With the released DSpark draft
(RadixArk/Qwen3.8-27B-DSpark, block_size = 7), a depth of 7 or more reads past the end of two
allocations. At depth == 7 this silently breaks losslessness; at depth > 7 speculative decode dies.

The >= kDeepMinSeq branch selects exactly 7 by default, so the default long-context path
requests the first unsafe value.

Mechanism

• runtime/src/models/dflash_draft.cpp:435 — h_out is allocated B ints (B = block_size).
• runtime/src/models/dflash_draft.cpp:1367 — readout is for (int t = 0; t <= kProposalDepth; t++) out_argmax[t] = s.h_out[t];, i.e. kProposalDepth + 1 entries.
• runtime/src/models/qwen35.cpp:3010 — set_dflash_capture(true, dc.target_layer_ids, B) sizes the capture buffer for B rows (dflash_max_rows, used at :2749).
• runtime/src/models/qwen35.cpp:3320 — the compact verify requests vn = kProposalDepth + 1 rows; dflash_warm_verify(kProposalDepth + 1, ...) at :3193 does the same.
• runtime/src/models/qwen35.cpp:3085-3112 — depth is clamped only as v > 15 ? 15 : v, with no reference to block_size. BW in dflash_draft.cpp is clamped to c.block_size, so the draft computes at most B rows while the caller reads kProposalDepth + 1.

Evidence

compute-sanitizer (--tool memcheck, ctx=128, COMPACT_VERIFY=1 PROPOSALS=7):

========= Invalid __global__ read of size 2 bytes
=========     Address 0x74d67d3f2880 is out of bounds
=========         Host Frame: sparkinfer::DFlashDraftModel::forward_block(...) in libsparkinfer_runtime.so
=========         Host Frame: sparkinfer::Qwen35Model::dflash_generate(...) in libsparkinfer_runtime.so

Losslessness, dspark_tau_check, ctx=128, SPARKINFER_DFLASH_COMPACT_VERIFY=1, 2 reps each:

SPARKINFER_DFLASH_PROPOSALS	vn = depth+1 vs B=7	LOSSLESS
5	6 — fits	1, 1
6	7 — fits exactly	1, 1
7	8 — overruns	0, 0
8	9 — overruns	run aborts: [dflash] verify failed, [FAIL] dflash_generate produced nothing

The boundary is exactly depth == block_size, deterministic across reps.

Suggested fix

--- a/runtime/src/models/qwen35.cpp
+++ b/runtime/src/models/qwen35.cpp
@@ -3108,8 +3108,15 @@ std::vector<int> Qwen35Model::dflash_generate(...)
-    const int kProposalDepth = kProposalDepthEnv > 0 ? kProposalDepthEnv
-                             : ((n + max_new) >= kDeepMinSeq ? 7 : 3);
+    const int kProposalDepthWanted = kProposalDepthEnv > 0 ? kProposalDepthEnv
+                                  : ((n + max_new) >= kDeepMinSeq ? 7 : 3);
+    // Clamp to the draft's own block. The verify runs kProposalDepth+1 rows, the capture buffer is
+    // sized for B rows (set_dflash_capture(..., B) above) and the draft's h_out is allocated B ints
+    // (dflash_draft.cpp:435), so a depth >= B reads past both. B == 7 for the released DSpark draft
+    // and the >= kDeepMinSeq branch asks for exactly 7, so the default long-context path selects
+    // the unsafe value. The draft cannot propose more than B-1 tokens after the seed in any case,
+    // so this only removes rows that were never backed by a real proposal.
+    const int kProposalDepth = (B > 1 && kProposalDepthWanted > B - 1) ? B - 1 : kProposalDepthWanted;

With this applied (same box, same build):

case	before	after
ctx=128 default (the scored config)	LOSSLESS=1, tau 1.0079, 86.05 tok/s	LOSSLESS=1, tau 1.0079, 86.72 tok/s
ctx=128, compact, depth 7	LOSSLESS=0	LOSSLESS=1, tau 1.2075
ctx=128, compact, depth 8	aborts	LOSSLESS=1, tau 1.2075

The scored ctx=128 path is unaffected (it selects depth 3, well inside B-1).

Scope / what I did not establish

I also saw [FAIL] dflash_generate produced nothing at ctx≈12400 on stock defaults, but a
control at PROPOSALS=3 fails there too, so that is a separate problem and this patch does not
fix it. I have not root-caused it and it may be specific to my synthetic prompt (bench_prompt.txt
repeated to 12400 ids) or to dspark_tau_check's env pins at long context — flagging it only so it
is not conflated with the overrun above.

Environment

• origin/main @ 9837dbd, RTX 5090 (sm_120), CUDA 12.8, -DCMAKE_CUDA_ARCHITECTURES=120
• target gittensor-model-hub/Qwen3.8-27B-NVFP4-RTX5090, draft RadixArk/Qwen3.8-27B-DSpark
• repro: build/runtime/dspark_tau_check <MODEL_DIR> <DRAFT_DIR> 128 $(ids) with
  SPARKINFER_DFLASH_COMPACT_VERIFY=1 SPARKINFER_DFLASH_PROPOSALS=7

Happy to open a PR with the patch if useful — I held off because the DSpark bot auto-closes any PR
that doesn't move decode@128, and this is a correctness fix that deliberately leaves that metric flat.