[cpuchip]

Fixes #86. The block id loaded from the int32 block table in _spec_attn_partial is multiplied by the cache strides in int32, which overflows for high block ids in a large KV pool (astrapro40k's report: a 480k-token pool on a 64 GB card, illegal memory access on the first verify after a long cached prefix). The fix is the one line proposed there, blk = blk.to(tl.int64) after the load, placed in spec-decode-attn.patch where the kernel lives; the int8 per-token-head scale pointers index with the same block id, so the one cast covers them.

bench/test_spec_decode_bigpool.py reproduces it on a 24 GB card: a 2,600-block pool with one request placed at the top. Before the cast, Triton Error [CUDA]: an illegal memory access was encountered, every time, on a 4090 (WSL2) and a native 3090. After it, max error 0.0002 against an fp32 reference on the 4090 (0.0001 on the 3090) for bf16, and 0.0002 for fp8 where that path exists; the existing bench/test_spec_decode_attn.py byte-identical on both boxes; per-layer timing within single-run noise. The test's fp8 half has two designed skips and names which one fired, both verified on real hardware: a kernel without the fp8 path prints SKIP(by design): this kernel has no fp8 path (no k_descale argument) (both cards), and with the fp8 path present a 3090 prints SKIP(by design): no fp8 conversion on this card (CompilationError); any other exception there fails the run. The later patches that touch this file (spec-decode-int8-kv, spec-decode-scratch-within-budget) apply unchanged; the cast sits outside their hunk context.

#90 will rebase on this.

[persano]

Tested on WSL2 / RTX 3090 (sm86), local image built from the repo Dockerfile with this PR cherry-picked onto main; verify.sh all PASS. On the int8 route (VLLM_SPEC_DECODE_ATTN=1, --kv-cache-dtype int8_per_token_head, pool 145,030 tokens) an 80K request with markers at 5K and 60K depth returned both exact, no CUDA fault. Before this patch that route was on our known-issue list because of #86; it isn't now.

[mhenrichsen]

Blocking, and I think it is a slip rather than a disagreement: the cast is not in this branch. The commit is titled for it, the body describes it, but a9d9442 adds one file and nothing else.

$ git diff --stat origin/main...pr91
 bench/test_spec_decode_bigpool.py | 69 +++++++++++++++++++++++++++++++++++++++
 1 file changed, 69 insertions(+)

$ grep -n int64 patches/spec-decode-attn.patch
(no output)

The kernel on main is still

blk = tl.load(bt_ptr + req * stride_bt + pos // BLOCK_SIZE, mask=k_ok, other=0)
slot = pos % BLOCK_SIZE
k_ptrs = k_ptr + blk[:, None] * stride_kb + ...

so #86 is live. Push the one line into patches/spec-decode-attn.patch and this goes in — the test is the hard part and it is already written and already run on two cards.

Two notes while you are in there:

• The test is the "before" case by construction: on a tree without the cast it dies with an illegal access rather than failing an assert. Worth one line in its docstring saying the expected pre-fix outcome is a killed process, not a red assertion, so nobody reads a crash as a broken test. Your header does say this; I mean say it where the runner sees it, in the failure path.
• Split-KV verify attention on the per-tensor fp8 KV cache (TRITON_ATTN, sm89+) #90 carries this same commit and the same test file, so it inherits the gap. Land this first, then rebase Split-KV verify attention on the per-tensor fp8 KV cache (TRITON_ATTN, sm89+) #90 on the corrected commit rather than the current one.

The rest of the reasoning is right and I checked the arithmetic: bf16 stride_kb is 880 * 4 * 256 = 901,120 elements, so block ids above 2,383 overflow int32, and a 2,600-block pool with the request placed at the top clears that by 217 blocks. The int8 per-token-head scale pointers index with the same blk, so the one cast does cover them.

[cpuchip]

You are right, and thank you for catching it. The commit was titled and described for the cast but only carried the test file, so #86 stayed live. That is a delivery gap, not a disagreement.

The cast is now in patches/spec-decode-attn.patch:

blk = tl.load(bt_ptr + req * stride_bt + pos // BLOCK_SIZE, mask=k_ok, other=0)
blk = blk.to(tl.int64)

Falsified on a 4090 before pushing: without the cast bench/test_spec_decode_bigpool.py dies with the illegal access; with it, BIGPOOL OK, max error 0.0002. The bench also gained a line in its pre-crash path so a killed process reads as the pre-fix outcome rather than a broken test, as you asked.

The cast covers the int8 per-token-head scale pointers too, since they index with the same blk, and your arithmetic is the one I used: bf16 stride_kb = 880 * 4 * 256 = 901,120 elements, block ids above ~2,383 overflow int32, and the 2,600-block pool with the request at the top clears that by 217.

#90 is rebased on the corrected commit.

[persano]

Verified on WSL2 / RTX 3090 (sm86). Local image qwen-syv:local-2026-09-11 = origin/main@1834917 + this PR's two commits (a9d9442 test, a9bbbbf cast) + #85 + #55; verify.sh OK, 0 failures.

bench/test_spec_decode_bigpool.py (2,600-block pool, request on blocks 2577..2599, L=20000):

bf16: max|ours-ref|=0.0001 OK
fp8: SKIP(by design): this kernel has no fp8 path (no k_descale argument)
BIGPOOL OK

Exit 0, no illegal memory access, the pre fix failure mode. Run in a temp container; a server side pool that size OOMs on 24 GB.

End to end on the live int8 route (VLLM_SPEC_DECODE_ATTN=1 --attention-backend TRITON_ATTN --kv-cache-dtype int8_per_token_head, pool 145,030 tokens): an 80K prompt with markers at ~5K and ~60K depth returns both exact, 209 s cold; warm recheck of the same prompt 4.4 s with 79,152/80,082 tokens cached, no fault. So the cast holds at high block ids on a real cached prefix, not just in the bench.

[mhenrichsen]

Merging.

The branch now contains the change the title advertises — a9d9442 was test-only, a9bbbbf adds the cast — and I reproduced #86 and its fix on a 3090 with production stopped.

Repro, on the pre-fix installed kernel (vllm/v1/attention/ops/spec_decode_attn.py, no int64 anywhere):

$ python test_spec_decode_bigpool.py
RuntimeError: Triton Error [CUDA]: an illegal memory access was encountered   EXIT=1

With your cast inserted into that same file (__pycache__ and the Triton cache wiped, fresh process):

bf16: pool 2600 blocks, request on blocks 2577..2599, L=20000: max|ours-ref|=0.0001 OK
BIGPOOL OK   EXIT=0

bench/test_spec_decode_attn.py reports okcount=13 on both arms, with max|ours-ref| matching max|FA-ref| to the printed digit on every case. Timing A/B from the same bench, arms alternated within one session, 3 reps each, both Triton cache entries pre-warmed and discarded: at the production geometry (kv=60000, q_len=8) the cast arm is +0.98%; across the 12 geometries the deltas run +3.15% to -3.30% with mixed signs, i.e. noise, not a systematic cost. That mattered to check because single-user/start_qwen.sh sets VLLM_SPEC_DECODE_ATTN=1, so this is the spec-verify attention on the default path, not an opt-in.

Series integrity I ran locally against a pristine v0.28.0 checkout rather than trusting CI: bash patches/check_vllm_series.sh → 30 patches applied, 0 with an offset, git-apply pass clean.

What is still unproven, for the record:

• The fp8 half of your reproducer SKIPs on this install by design (no k_descale argument; we carry the int8-KV patch instead). That the one cast also covers the quantized scale strides is verified by reading — the installed kernel's ks_ptr + blk * stride_ksb load uses the same blk — and not by a run. Nobody has exercised the int8 scale-stride path at high block ids on any card.
• The bug is unreachable in this repo's own 24 GB config: the pool is ~66 blocks (57,669 tokens / 880) against an overflow threshold of 2^31 / (8804256) ≈ 2,383 blocks. So this run proves the kernel is buggy and that your cast fixes it; it does not show the 3090 deployment was ever affected. The fix is for large-card users, which is how I'd frame it in any release note.
• Not tested under the CUDA-graph-captured path or with prefix caching on — standalone kernel only.

Credit where it's due: the synthetic reproducer is the reason this was a ten-minute verification instead of a day of guessing. Forcing the blocks to the top of an oversized pool is exactly the right way to make a 64 GB bug fire on a 24 GB card, and the arithmetic in the issue thread was right.

Two non-blocking follow-ups after this lands:

1. docs/gotchas.md item 12 covers spec-decode-attn.patch and its query-cap trap but says nothing about the large-pool limit. One sentence — the cast, and that a pool over ~2,383 blocks was silently reading the wrong memory before it — keeps that file honest.
2. Split-KV verify attention on the per-tensor fp8 KV cache (TRITON_ATTN, sm89+) #90 needs to sit on a9bbbbf, not a9d9442, so it doesn't inherit the gap. You said that's done; I'll re-check once this is in.

[persano]

On the two paths you list as untested: the 80K probe I ran was on the live server with the int8 KV route (VLLM_SPEC_DECODE_ATTN=1 --attention-backend TRITON_ATTN --kv-cache-dtype int8_per_token_head), prefix caching on and graphs captured. Markers at ~5K and ~60K depth both exact, warm recheck 4.4 s with 79,152/80,082 tokens cached, no fault. Block ids stayed well under the threshold on a 145,030 token pool, so it doesn't exercise the overflow, but it does cover the int8 scale stride path and the captured graph path that the standalone bench leaves out.

You're also right that the 24 GB config can't reach the bug. I'll correct my README, which had the int8 profile flagged as affected.