[Bugfix] Guard slot mapping block table loads

[frank-suwen]

Purpose

Fixes #53982.

Both ComputeSlotMappingKernel and the Model Runner V2 _compute_slot_mappings_kernel() can compute a block index beyond the width of a request's block-table row when a narrow cache group's table does not span the sequence in raw token positions. This can cause an out-of-bounds GPU read.

This change adds block_indices < block_table_stride to both load masks and emits PAD_SLOT_ID for out-of-range lanes. The explicit padding also covers the Model Runner V2 CP_SIZE == 1 path, where the load's other=0 would otherwise produce a valid-looking slot in block zero. In-range behavior is unchanged.

I searched open PRs by issue number, affected symbols, behavior, and file paths. The related open PRs do not implement this bounds guard: #51017 adds broader test-only coverage for V1 slot mapping, while #50287 addresses Mamba block-table capacity and staged writes.

OpenAI Codex assisted with investigation, implementation, test design, GPU validation, and PR preparation. I reviewed and understand all changed lines, ran the local checks, and validated the GPU behavior on an NVIDIA RTX 4090.

Test Plan

.venv/bin/python -m pytest tests/v1/worker/test_gpu_block_table.py -k out_of_range_block_indices -vv
.venv/bin/python -m pytest tests/v1/worker/test_gpu_block_table.py -vv
.venv/bin/python -m pytest tests/v1/worker/test_jit_warmup_migration.py -vv
.venv/bin/pre-commit run --files tests/v1/worker/test_gpu_block_table.py vllm/v1/worker/block_table.py vllm/v1/worker/gpu/block_table.py

Test Result

Validated on an NVIDIA RTX 4090 with PyTorch 2.13.0+cu130:

• Before the production fix, both focused regression tests failed as expected: the out-of-range lane produced slot 144 from the adjacent request row instead of PAD_SLOT_ID.
• After the fix, the focused regression tests passed: 2 passed, 8 deselected.
• Full GPU block-table module: 10 passed.
• JIT warmup migration module: 5 passed.
• All applicable pre-commit hooks passed.

The issue discussion also contains independent validation of equivalent guards: @zzw09773 tested the V1 path on 8× MI325X, and @stefanskiasan tested the V2 path on 8× MI350X. Their testing during the issue investigation provides additional support for the fix. I really appreciate both contributors sharing their results and offering to validate this PR branch once it is available.

[frank-suwen]

@zzw09773 @stefanskiasan, the draft PR implementing the bounds guards for both slot-mapping paths is now available. Thank you again for sharing your MI325X and MI350X results and offering additional validation. When convenient, could you please validate this branch on your respective ROCm setups and report the results here? Thanks!

[claude]

Claude Code Review

This pull request is from a fork — automated review is disabled. A repository maintainer can comment @claude review to run a one-time review.

[drakosha]

Third-platform data point, CUDA rather than ROCm: SM90, 2x H200 NVL, TP2,
GLM-5.3-Flash NVFP4 off the #53906 branch, fp8 KV via FLASHINFER_MLA_SPARSE_SM90,
MTP k=3. The equivalent of the V2-path hunk has been in our production image since
2026-08-29.

Instrumented, not read: a one-shot log in KpoolTailMetadataBuilder.build() shows the
generic slot mapping for the kpool tail group is -1 for every token, i.e. the
out-of-row block index is clamped to PAD and the tail kernels skip their writes instead
of gathering a garbage block id. So here the unguarded read is a silent skip rather than
a fault. That skip exposed a second, unrelated bug (the tail pool was never written,
positions was not passed down; fixed separately in the #53906 thread) — the guard is
not masking it.

Load carried since: needles on 1M-token prompts 4/4, 32 concurrent 118k prompts with 0
failures and 0 preemptions, an eviction cycle of eleven 500k prompts against a 4.9M-token
KV pool with CPU offload round-trip. No illegal accesses, no engine deaths.

AI assistance was used for this work.

[frank-suwen]

Third-platform data point, CUDA rather than ROCm: SM90, 2x H200 NVL, TP2, GLM-5.3-Flash NVFP4 off the #53906 branch, fp8 KV via FLASHINFER_MLA_SPARSE_SM90, MTP k=3. The equivalent of the V2-path hunk has been in our production image since 2026-08-29.

Instrumented, not read: a one-shot log in KpoolTailMetadataBuilder.build() shows the generic slot mapping for the kpool tail group is -1 for every token, i.e. the out-of-row block index is clamped to PAD and the tail kernels skip their writes instead of gathering a garbage block id. So here the unguarded read is a silent skip rather than a fault. That skip exposed a second, unrelated bug (the tail pool was never written, positions was not passed down; fixed separately in the #53906 thread) — the guard is not masking it.

Load carried since: needles on 1M-token prompts 4/4, 32 concurrent 118k prompts with 0 failures and 0 preemptions, an eviction cycle of eleven 500k prompts against a 4.9M-token KV pool with CPU offload round-trip. No illegal accesses, no engine deaths.

AI assistance was used for this work.

Thank you for sharing this detailed CUDA/H200 validation. It is very helpful. Thanks!

[zzw09773]

@frank-suwen Confirmation on 8× MI325X (gfx942, ROCm) for #54296 (191f82d710493c9a02077d976d8cf0403ab8c96a, branch fix/slot-mapping-block-table-bounds).

What we ran

• Image: our production GLM serving image local/vllm-openai-rocm:glm53-flash-gfx942-v9-tritonspec (throwaway container, HIP_VISIBLE_DEVICES=0 only; no model weights loaded).
• Tests: the two focused cases from this PR —
  test_v1_slot_mapping_masks_out_of_range_block_indices,
  test_v2_slot_mapping_masks_out_of_range_block_indices
  (python -m pytest … -k out_of_range_block_indices -vv).

Static note vs our image

• Full-file overlay of the PR kernels onto this image is not import-compatible (V1 needs triton_scalar_specialization_rep from current main; V2 kernel expects kernel_block_sizes which this image’s launcher does not pass).
• Our production image already carries an equivalent V1 guard (in_range + PAD_SLOT_ID in vllm/v1/worker/block_table.py).
• Our production V2 path (vllm/v1/worker/gpu/block_table.py) did not yet have the bounds mask (GLM-5.3-Flash uses the V2 runner).

Results (gfx942)

1. Stock image (no overlay): V1 PASSED, V2 FAILED — got tensor([80, 144]) instead of [80, PAD_SLOT_ID]; 144 is the neighboring row’s block_id*block_size, i.e. classic OOB load across the row stride.
2. Same image with only the V2 in_range / PAD_ID hunk applied (PR semantics, adapted to this image’s kernel signature): 2 passed in 6.71s.

So: the PR’s V1 fix matches what we already ship; the V2 fix is necessary and the focused tests pass on MI325X/gfx942 once applied. We are baking that equivalent V2 hunk into the serving image now. Happy to pull the PR branch into a clean tree if you want a verbatim checkout re-run as well.

[frank-suwen]

Thank you @zzw09773 for the detailed MI325X validation! The results you shared really provide strong confirmation of the fix in this PR. Yes, if convenient, I would also appreciate a rerun from a clean checkout of the exact PR branch for additional confirmation. Thanks!

[zzw09773]

Follow-up: the remaining gfx942 GLM-5.3-Flash abort (Memory access fault after hours of MTP serving) is not this V2 slot-mapping OOB. Crash B was already on an image with the V2 in_range guard and still died. Separate write-up: #55280

[zzw09773]

@frank-suwen Follow-up: clean checkout of this branch at 191f82d710493c9a02077d976d8cf0403ab8c96a on the same 8× MI325X (gfx942, ROCm).

Checkout. git clone --depth 1 --branch fix/slot-mapping-block-table-bounds from frank-suwen/vllm → 191f82d (unchanged).

The two focused tests from this PR (test_v1_slot_mapping_masks_out_of_range_block_indices, test_v2_slot_mapping_masks_out_of_range_block_indices):

• pytest -k out_of_range_block_indices is skipped on ROCm by the file’s skipif(not current_platform.is_cuda()) (RocmPlatform.is_cuda() is False). We invoked the two test functions directly.
• Whole-tree PYTHONPATH of this checkout: V1 PASSED. V2 cannot construct BlockTables on this runtime — this branch’s buffer_utils.py calls torch.ops._C.get_cuda_view_from_cpu_tensor, which the ROCm image does not ship.
• Overlaying this PR’s vllm/v1/worker/gpu/block_table.py (the file at 191f82d, not an adapted hunk) onto the image’s matching V2 launcher: V1 PASSED, V2 PASSED (got [80, PAD_SLOT_ID]).

So the PR’s V2 kernel, run as-is from this branch, masks the OOR row-stride load on gfx942. Happy to rerun if the branch moves.

[frank-suwen]

Hi @zzw09773, thank you for the follow-up and for clearly documenting the checkout and runtime limitations. Your confirmation and detailed results are very helpful. This looks good to me, and yes, I do not think another rerun is needed at this point. Thanks!

[lkeld]

@frank-suwen @zzw09773 adding the NVIDIA data the issue asked for ("we only have gfx942 data"). measured parts i have labeled and the rest is from reading the tree

environment

• vllm/vllm-openai:glm53-flash-x86_64-cu130 @ sha256:0674a3ea3c971706cea310bca71daadff68ea1717c758d2f54affa5def001b86 = vLLM v0.28.1rc1.dev580+g385dce36b (the [Model] add GLM-5.3-Flash support #53906 tree), torch 2.13.0+cu130, triton 3.7.1, CUDA 13.0 image. Using V2 Model Runner (the tree's default).

• Serving that crashed: 1x B300 and, separately, 2x H200 TP=2, with [Bugfix] GLM-5.3-Flash: fp8 plan dtype on SM90 sparse MLA, and right-size the indexer prefill workspace #55222 applied on top. Sanitizer reproduction below: 1x B300 SXM6 (sm_103), driver 595.84, pristine day-0 image. End-to-end below: 2x H200, driver 590.48.01.

• orcarouter/GLM-5.3-Flash-Uncensored-NVFP4 (text_config identical to zai-org/GLM-5.3-Flash: index_kpool=4, index_topk=2048).

• --max-model-len 1048576 --max-num-seqs 64 --kv-cache-dtype bfloat16 --moe-backend marlin, chunked prefill, no speculative decoding. Boot: attention block 4352 (hybrid alignment), indexer page 64, kv cache group sizes [4352, 4, 4352, 4352, 4352, 4352] (the 4 is the kpool tail group).

symptom

CUDA error: an illegal memory access was encountered, async, surfacing at the next launch/sync: seen in indexer.py build_prefill_chunk_metadata (Triton launch), in the KDA chunk_kda_with_fused_gate_fwd -> prepare_chunk_indices -> .tolist(), and once in the indexer decode path. only with random-word prompts (repeated text hash-dedupes and never reaches the position). the position it dies at moves with knobs that only change the allocation layout, which is what #53982 predicts for a layout-dependent overrun:

config	result
prefix caching on, --max-num-batched-tokens 8192 (4352-token chunks)	dies on the first prefill chunk past ~493K; below that, 90 min / ~300M tokens of 64-session load at up to 351K was clean
prefix caching on, --max-num-batched-tokens 32768 (26112-token chunks)	dies at ~351K
prefix caching off, 2000-token chunks	565K and 1,040,000-token prompts complete
prefix caching on + PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True (plus an unrelated local indexer patch)	dies on the first request, a 58K prompt

we initially read this as rows x context int32 overflow in the indexer kernels; it is not (those use 64-bit row offsets and the builder caps each logits tensor at VLLM_SPARSE_INDEXER_MAX_LOGITS_MB).

mechanism on this tree (from reading 385dce36b).

• KpoolTailSpec.max_num_blocks_per_req() returns 1; vllm/v1/worker/gpu/model_runner.py widens it with get_block_table_width(1, 4) (128-token alignment) to 32 columns and keeps slot_mapping_enabled=True for it, since only CircularBufferSpec is excluded and KpoolTailSpec is a SlidingWindowSpec. Same width on the V1 runner path (SlotMappingMode.NONE is Mamba-only).
• _compute_slot_mappings_kernel (and ComputeSlotMappingKernel.kernel) then loads block_table[req, pos // 4] for every scheduled token: past the request's row at position 128, past the whole [max_num_seqs, 32] int32 tensor (8 KB at 64 seqs) at ~8K, about 1 byte further per token of position. KpoolTailMetadataBuilder.build() recomputes the circular mapping into its own buffer, so the read is the only effect, consistent with the crash being the only symptom.

model-free reproduction on B300

slotmap_oob_repro.py (collapsed below) builds the GLM tail group exactly as the runner does (get_block_table_width(1, 4) = 32 columns, 2 requests, i.e. a 256-byte int32 block table) with the runner's own BlockTables / BlockTable, and runs both kernels for positions [0, 3, 130, 8192, 100000, 493447, 1048575]. no weights; seconds. perfect day 0 container (neither kernel has the guard; verified this before running).

PYTORCH_NO_CUDA_MEMORY_CACHING=1 compute-sanitizer --tool memcheck --print-limit 12 python3 slotmap_oob_repro.py

Unpatched (threads 3-6 are positions 8192, 100000, 493447, 1048575; the distance past the table grows one byte per token of position):

========= Invalid __global__ read of size 4 bytes
=========     at _compute_slot_mappings_kernel+0x1fc0 in block_table.py:344
=========     by thread (3,0,0) in block (0,0,0)
=========     Access to 0x777e35602000 is out of bounds
=========     and is 1993 bytes after the nearest allocation at 0x777e35601800 of size 56 bytes
=========         Host Frame: compute_slot_mappings in block_table.py:204
=========         Host Frame: run_v2 in slotmap_oob_repro.py:63
========= Invalid __global__ read of size 4 bytes
=========     at _compute_slot_mappings_kernel+0x1fc0 in block_table.py:344
=========     by thread (4,0,0) in block (0,0,0)
=========     and is 93801 bytes after the nearest allocation at 0x777e35601800 of size 56 bytes
========= Invalid __global__ read of size 4 bytes
=========     by thread (5,0,0) in block (0,0,0)
=========     and is 487245 bytes after the nearest allocation at 0x777e35601800 of size 56 bytes
========= Invalid __global__ read of size 4 bytes
=========     by thread (6,0,0) in block (0,0,0)
=========     and is 1042373 bytes after the nearest allocation at 0x777e35601800 of size 56 bytes
========= Program hit cudaErrorLaunchFailure (error 719) due to "unspecified launch failure" on CUDA API call to cudaDeviceSynchronize.
torch.AcceleratorError: CUDA error: unspecified launch failure
========= ERROR SUMMARY: 6 errors

Same tree with this PR's diff applied (patch -p1 of 191f82d; both hunks apply as-is):

[V2 _compute_slot_mappings_kernel (default runner)] tail block table (2, 32) int32 = 256 bytes (width 32 columns)
   pos=        0  col=      0  in row                       -> slot=28
   pos=        3  col=      0  in row                       -> slot=31
   pos=      130  col=     32  row overrun (+4 B)           -> PAD
   pos=     8192  col=   2048  OUTSIDE tensor (+7940 B)     -> PAD
   pos=   100000  col=  25000  OUTSIDE tensor (+99748 B)    -> PAD
   pos=   493447  col= 123361  OUTSIDE tensor (+493192 B)   -> PAD
   pos=  1048575  col= 262143  OUTSIDE tensor (+1048320 B)  -> PAD
[V1 ComputeSlotMappingKernel] tail block table (2, 32) int32 = 256 bytes (width 32 columns)
   (identical: in-range slots 28 / 31, PAD for every out-of-range position)
========= ERROR SUMMARY: 0 errors

This PR's own regression tests (pytest -k out_of_range_block_indices tests/v1/worker/test_gpu_block_table.py, file taken from 191f82d) on the same box: unpatched 2 failed, both returning tensor([80, 144]) where [80, -1] is expected (144 is the adjacent row's block id, i.e. the overrun); with the diff 2 passed.

end to end

Same image (385dce36b + #55222) with this PR's guard, on 2x H200 (driver 590.48.01), TP=2: Using V2 Model Runner, FLASHINFER_MLA_SPARSE_SM90, Marlin NVFP4 MoE, prefix caching on (Mamba cache mode is set to 'align'), --max-model-len 1048576 --max-num-seqs 64 --max-num-batched-tokens 8192 --kv-cache-dtype bfloat16. On this TP=2 layout the hybrid attention block is 2176 (kv cache group sizes [2176, 4, 2176, 2176, 2176, 2176]), GPU KV cache size: 2,735,694 tokens, Maximum concurrency for 1,048,576 tokens per request: 2.61x. Random-word prompts with a per-request nonce (prefix cache hit rate 0.0%), max_tokens 16-32:

• 599,171-token prompt: completed, 72.5 s prefill+decode, answered OK.
• 1,038,183-token prompt: completed, 116.6 s.
• 64 concurrent sessions for 30 min, prompt lengths 2,105-297,194 tokens (median 167,632): 208 requests, 0 failures, 33.7M prompt tokens; the engine log has 0 illegal memory access / CUDA error / worker deaths and /health is OK afterwards. (Latency there is queueing: the default max_num_partial_prefills=1 serializes long prefills; prompt throughput bursts at ~24-27K tokens/s.)

Without the guard the same image died as in the table above. For completeness, that image also carries #55450, #55270 and a stride fix for _kpool_tail_seed_kernel (separate PR, refs #56037/#56059); none of them touch the slot-mapping path.

slotmap_oob_repro.py

import os, sys, torch
from vllm.v1.attention.backends.utils import PAD_SLOT_ID
from vllm.v1.worker.block_table import BlockTable, get_block_table_width

KPOOL = 4                      # index_kpool == KpoolTailSpec.block_size
MAX_NUM_REQS = 2
POSITIONS = [0, 3, 130, 8_192, 100_000, 493_447, 1_048_575]
DEV = torch.device("cuda")

def describe(width, table, out, label):
    nbytes = table.numel() * table.element_size()
    print(f"[{label}] tail block table {tuple(table.shape)} int32 = {nbytes} bytes (width {width})")
    for pos, slot in zip(POSITIONS, out.tolist()):
        col = pos // KPOOL
        past_row = (col - width + 1) * 4
        past_tensor = (col - MAX_NUM_REQS * width + 1) * 4
        where = "in row" if col < width else (f"row overrun (+{past_row} B)" if past_tensor <= 0 else f"OUTSIDE tensor (+{past_tensor} B)")
        print(f"   pos={pos:>9}  col={col:>7}  {where:<28} -> {'PAD' if slot == PAD_SLOT_ID else f'slot={slot}'}")

def run_v2():
    from vllm.v1.worker.gpu.block_table import BlockTables
    width = get_block_table_width(1, KPOOL)          # what gpu/model_runner.py passes for KpoolTailSpec
    bt = BlockTables(block_sizes=[KPOOL], max_num_reqs=MAX_NUM_REQS, max_num_batched_tokens=len(POSITIONS),
                     max_num_blocks_per_group=[width], device=DEV, kernel_block_sizes=[KPOOL])
    bt.append_block_ids(req_index=0, new_block_ids=([7],), overwrite=True)
    bt.append_block_ids(req_index=1, new_block_ids=([9],), overwrite=True)
    bt.apply_staged_writes()
    idx_mapping = torch.zeros(1, dtype=torch.int32, device=DEV)
    qsl = torch.tensor([0, len(POSITIONS)], dtype=torch.int32, device=DEV)
    pos = torch.tensor(POSITIONS, dtype=torch.int64, device=DEV)
    out = bt.compute_slot_mappings(idx_mapping, qsl, pos, num_tokens_padded=len(POSITIONS))[0]
    torch.cuda.synchronize()
    describe(width, bt.block_tables[0].gpu, out, "V2 _compute_slot_mappings_kernel (default runner)")

def run_v1():
    width = get_block_table_width(1, KPOOL)
    bt = BlockTable(block_size=KPOOL, max_num_reqs=MAX_NUM_REQS, max_num_blocks_per_req=width,
                    max_num_batched_tokens=len(POSITIONS), pin_memory=False, device=DEV,
                    kernel_block_size=KPOOL, cp_kv_cache_interleave_size=1)
    bt.add_row([7], row_idx=0); bt.add_row([9], row_idx=1); bt.commit_block_table(num_reqs=MAX_NUM_REQS)
    qsl = torch.tensor([0, len(POSITIONS)], dtype=torch.int32, device=DEV)
    pos = torch.tensor(POSITIONS, dtype=torch.int64, device=DEV)
    bt.compute_slot_mapping(num_reqs=1, query_start_loc=qsl, positions=pos)
    torch.cuda.synchronize()
    describe(width, bt.block_table.gpu, bt.slot_mapping.gpu[:len(POSITIONS)], "V1 ComputeSlotMappingKernel")

if __name__ == "__main__":
    import vllm
    print(f"vllm {vllm.__version__}  torch {torch.__version__}  gpu {torch.cuda.get_device_name(0)}")
    run_v2(); run_v1()

a GLM-only alternative would be to disable slot mapping for KpoolTailSpec the way the V2 runner already does for CircularBufferSpec (or #55219's CircularBufferSpec tail), but the generic bound here is the right fix since it protects every future narrow group. i am more than happy to re-run anything on either GPU.

ai assistance was used for this work

[drakosha]

Operational data point on top of @lkeld's sanitizer run, same guard, SM90:
2x H200 NVL, TP2, GLM-5.3-Flash NVFP4, fp8 KV, MTP k=3, prefix caching plus a
256 GiB CPU offload region, max_num_seqs 32, max_model_len 1048576. Eight days
up, 93k completed requests, 1.93e9 prompt tokens, zero preemptions, zero
illegal memory access in the engine log.

Correcting my Sep 2 comment here: the -1 slot mapping I reported came from a
second bug (positions not reaching the tail builder), fixed and merged with
#53906, and this PR's guard has been in our image since Sep 3. So we never ran
the faulting combination, and "silent skip, not a fault" should not be read as
"CUDA does not fault here".

AI assistance was used for this comment.

[frank-suwen]

Thank you @lkeld and @drakosha for the detailed NVIDIA validation and clarification. The compute-sanitizer reproduction, focused test results, and longer-running H200 data provide strong additional evidence for this fix.

@WoosukKwon, when convenient, could you please help to take a look at this PR? The guard is also referenced as a prerequisite by #56437, #56438, and #56439. Please let me know if any changes or additional validation would be helpful. Thank you again for your time!