fix(kv-cache): grow sequences by block delta, not append — prevents block leak + device block-table overrun (RTX 5090 verified)

[milosde111]

Summary

Fix KVCacheManager::allocate() to grow a sequence's block table by the delta (add = need - have) instead of unconditionally appending need blocks on every call. Prevents block leaks and device block-table row overruns when the same seq_id is re-allocated (the slot-reuse branch already anticipates this call pattern).

Correctness/robustness fix — no expected decode speedup. Callers today pre-allocate full max_seq once; the e2e bench confirms no perf regression.

The bug

KVCacheManager::allocate(seq_id, num_tokens) (runtime/src/kv_cache.cpp:64) computes need = ceil(num_tokens / block_size) as the total blocks the sequence should own, then unconditionally appends need blocks to seq_blocks[seq_id]. The seq_slot reuse branch right below it is explicitly written to handle an already-allocated seq_id (if (it != seq_slot.end()) slot = it->second;), so re-allocation is an anticipated call pattern — but the block-append path breaks it two ways:

1. Block leak. Growing a live sequence 100→200 tokens (block_size=16) allocates 13 more blocks on top of the 7 it already held (20 total) instead of the correct 6 delta — 7 blocks orphaned, unrecoverable until free().
2. Device out-of-bounds write. The only size guard is need > kMaxBlocksPerSeq; nothing bounds blocks.size(). Repeated growth pushes the vector past kMaxBlocksPerSeq (4096), and the cudaMemcpy(d_block_tables + slot*kMaxBlocksPerSeq, blocks.data(), blocks.size()*4, …) then overruns the fixed-width device table row into the next sequence's row — silent cross-sequence corruption.

Today both callers (generate, bench_decode) pre-allocate the full max_seq once and free() before re-allocating, so it's latent — but it's a memory-safety landmine that fires the moment KV blocks are allocated incrementally (prefill-then-extend), which the slot-reuse logic already presumes.

The fix

Grow by the delta only: add = need - have, append add blocks (0 when the sequence already has enough). A fresh seq_id has have == 0, so the common path is unchanged. Preconditions (free-list depth for add, a free table row only when the seq_id is new) are all checked before any free-list/slot mutation, preserving the original fail-clean behavior. blocks.size() is now always == need ≤ kMaxBlocksPerSeq, so the cudaMemcpy can never overrun the row.

---

⚠️ Tick the box only if you actually ran it on an RTX 5090. False attestation is treated as gaming.

• Tested on RTX 5090 (sm_120)

Proof of correctness (RTX 5090)

Hardware: NVIDIA GeForce RTX 5090, compute capability 12.0, 170 SMs, 1792 GB/s BW
Date: 2026-07-03 UTC
Repro command: cmake --build build --target kv_allocate_repro kv_cache_gpu_test && ./build/runtime/tests/kv_allocate_repro

KV allocate repro (runtime/tests/kv_allocate_repro)

allocate(1,100) → grow allocate(1,200) on live seq_id, block_size=16:

step	main (buggy)	this PR (patched)
after allocate(1,100)	used=7 ✓	used=7 ✓
after allocate(1,200)	used=20 (leaked 7)	used=13 ✓
re-allocate same size	unbounded growth	used=13 (idempotent) ✓
free(1)	blocks leaked	free=512 == total ✓
RESULT	FAIL	PASS

main (buggy):
after allocate(1,100): used=7 (expect 7)
after allocate(1,200): used=20 (correct=13) ← 7 blocks leaked
RESULT: FAIL (expected 13, got 20)

this PR (patched):
after allocate(1,100): used=7 (expect 7)
after allocate(1,200): used=13 (correct=13)
re-allocate same size: used=13 (expect 13)
free(1): free=512 (total=512)
GPU: NVIDIA GeForce RTX 5090 (sm_12.0)
RESULT: PASS

Additional GPU tests (RTX 5090, no regression)

[PASS] kv_cache_gpu_test on NVIDIA GeForce RTX 5090: fresh allocate, delta re-grow, no leak, seq isolation (512 blocks total)
[PASS] decode_runner_gpu_test: 2 layers x 2 seqs, output finite
[runtime] NVIDIA GeForce RTX 5090 cc=12.0 SMs=170 BW=1792 GB/s
[PASS] qwen35_gpu_test: generated 8 tokens: 136 136 136 136 136 136 136 136
[runtime] NVIDIA GeForce RTX 5090 cc=12.0 SMs=170 BW=1792 GB/s

Decode tok/s (end-to-end, from bench/scripts/bench.sh)

No expected speedup — correctness fix only. Bench confirms no regression vs main:

ctx	before (main)	after (this PR)	Δ
128	482.49 tok/s	482.74 tok/s	+0.05%
512	462.05 tok/s	462.21 tok/s	+0.03%
2048	407.16 tok/s	407.16 tok/s	0.00%
4096	349.60 tok/s	349.26 tok/s	−0.10%
16384	268.55 tok/s	268.64 tok/s	+0.03%

RTX 5090, Qwen3-30B-A3B-Q4_K_M.gguf, n=128 decode tokens, bs=1

before: main + cmake --build build --target qwen3_gguf_bench

after: fix/kv-allocate-delta-blocks + cmake --build build --target qwen3_gguf_bench

ctx=128:
before (main): 482.49 tok/s
after (this PR): 482.74 tok/s

ctx=512:
before (main): 462.05 tok/s
after (this PR): 462.21 tok/s

ctx=2048:
before (main): 407.16 tok/s
after (this PR): 407.16 tok/s

ctx=4096:
before (main): 349.60 tok/s
after (this PR): 349.26 tok/s

ctx=16384:
before (main): 268.55 tok/s
after (this PR): 268.64 tok/s

Build

cmake --build build --target sparkinfer_runtime
links clean

Fresh-allocate decode path is byte-for-byte unchanged (identical block accounting on the first allocate).

Notes for reviewers

• Fix: single hunk in runtime/src/kv_cache.cpp:64 on branch fix/kv-allocate-delta-blocks
• Tests added: runtime/tests/kv_allocate_repro.cpp, runtime/tests/kv_cache_gpu_test.cpp
• Correctness/robustness fix — no perf delta expected; request manual merge for correctness

[ai-hpc]

⏳ Needs a benchmark to be evaluated

You ticked Tested on RTX 5090 but the decode before → after tok/s table is still empty / placeholder (or shows no gain). The on-device eval won't run until it shows a real improvement.

Fill it from the end-to-end decode bench (not an isolated-kernel microbench):

bench/scripts/bench.sh --download            # baseline (before)
bench/scripts/bench.sh --download            # your branch (after)

Then the bot greenlights it on the next poll and evaluates it on an RTX 5090.

[milosde111]

@ai-hpc Filled the end-to-end decode bench on RTX 5090. This is a correctness fix, not a speedup PR — flat before/after is expected (callers only allocate() once today; decode hot path unchanged).

Decode tok/s (bench/scripts/bench.sh)

RTX 5090 · Qwen3-30B-A3B-Q4_K_M · n=128 · bs=1

ctx	before (main)	after (this PR)
128	482.49 tok/s	482.74 tok/s
512	462.05 tok/s	462.21 tok/s
2048	407.16 tok/s	407.16 tok/s
4096	349.60 tok/s	349.26 tok/s
16384	268.55 tok/s	268.64 tok/s

~0% across all contexts.

# RTX 5090, Qwen3-30B-A3B-Q4_K_M.gguf, n=128 decode tokens, bs=1
# before: main + cmake --build build --target qwen3_gguf_bench
# after:  fix/kv-allocate-delta-blocks + cmake --build build --target qwen3_gguf_bench

ctx=128:
before (main):    482.49 tok/s
after (this PR):  482.74 tok/s

ctx=512:
before (main):    462.05 tok/s
after (this PR):  462.21 tok/s

ctx=2048:
before (main):    407.16 tok/s
after (this PR):  407.16 tok/s

ctx=4096:
before (main):    349.60 tok/s
after (this PR):  349.26 tok/s

ctx=16384:
before (main):    268.55 tok/s
after (this PR):  268.64 tok/s

Correctness proof (kv_allocate_repro)

main (buggy):

after allocate(1,100):  used=7  (expect 7)
after allocate(1,200):  used=20  (correct=13)   ← 7 blocks leaked
RESULT: FAIL (expected 13, got 20)

this PR (patched):

after allocate(1,100):  used=7  (expect 7)
after allocate(1,200):  used=13  (correct=13)
re-allocate same size:  used=13  (expect 13)
free(1):                free=512  (total=512)
GPU: NVIDIA GeForce RTX 5090 (sm_12.0)
RESULT: PASS

Additional GPU tests (no regression)

[PASS] kv_cache_gpu_test on NVIDIA GeForce RTX 5090: fresh allocate, delta re-grow, no leak, seq isolation (512 blocks total)
[PASS] decode_runner_gpu_test: 2 layers x 2 seqs, output finite
[PASS] qwen35_gpu_test: generated 8 tokens: 136 136 136 136 136 136 136 136

Requesting manual merge on correctness (block leak + device block-table OOB) — not seeking frontier / speedup eval for this PR.