Repository navigation
Conversation
The cuBLAS decode route of the SM70 indexer gathers the keys once and scores them into one [rows * heads, keys] FP32 tile, admitted only when both fit the indexer workspace budget (indexer_prefill_tile_mb, default 192 MiB). Beyond that the decode falls back to the paged Triton kernel, which at 140k tokens of context took 1.27 ms per call on a V100 and 6.9 ms on a Quadro RTX 8000, against about 0.11 ms for the GEMM of the 131,072-token bucket. For six verifier rows of 64 heads the default budget admits about 112,000 compressed keys; the 64 MiB a PP5 deployment uses to keep its V100 prefill stages clear of OOM admits 32,768 keys (131,072 tokens). The gathered keys stay whole and the scores are computed in key tiles that fit next to them. A key bound whose scores fit keeps its single tile, so the route and its result are unchanged wherever it was admitted before. Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01GJxVXkjSHhfq2D71DoYSK7
…the model length The automatic buckets stopped at 64 times the short bound (131,072 tokens for DeepSeek-V4). Longer contexts replayed the full-length graph, where the indexer stays on its paged kernel by design. A last bucket at the maximum model length keeps those contexts on the cuBLAS decode route. Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01GJxVXkjSHhfq2D71DoYSK7
The dispatcher tests build VllmConfig as a MagicMock and set the configs the code under test reads. CUDAGraphWrapper now also builds a decode event tracer from observability_config.runtime_trace; on the mock that policy was a MagicMock, so tracing ran and compared the elapsed time against a MagicMock threshold. Three tests failed with "'<' not supported between instances of 'float' and 'MagicMock'" (TestCUDAGraphWrapper:: test_capture_and_replay, test_capture_replay_bypass_logic, test_nested_wrappers). The mock now carries a default ObservabilityConfig, as other config tests do. Signed-off-by: Peuqui <43776522+Peuqui@users.noreply.github.com> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01GJxVXkjSHhfq2D71DoYSK7
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Purpose
DeepSeek-V4 decode on SM70/SM75 slows down sharply once the context passes 131,072 tokens. On our PP5 rig (three Tesla V100, two Quadro RTX 8000) the decode step went from 80 to 154 ms at the moment a request crossed 131,072 tokens and grew to 250 ms at 283k. V1 and MRV2 behave the same (178,479 tokens: 180.7-180.9 ms on V1, 174.8-175.1 ms on MRV2).
An nsys capture of the same decode at 124,029 and at 140,029 tokens shows the cause. Below the threshold the indexer scores with the cuBLAS GEMM of its decode route (volta_s884gemm, about 0.11 ms per call on a V100). Above it,
_paged_index_logits_relu_kernelruns instead: 1.27 ms per call on a V100 and 6.9 ms on a Quadro RTX 8000, about 73 ms per decode step in total. Two gates keep the cuBLAS route off there:_decode_cublas_blockerkeeps the indexer on the paged kernel by design.indexer_prefill_tile_mb). For six verifier rows of 64 heads the default 192 MiB admits about 112,000 compressed keys; the 64 MiB a PP5 deployment uses to keep its V100 prefill stages clear of OOM admits 32,768 keys (131,072 tokens). A diagnostic boot confirmed this gate for the 262,144 and 307,200 buckets. Raising the budget is no way out: the same field bounds the prefill score tile, and 512 MiB ran a V100 stage out of memory in the prefill.This PR closes both gates, in two commits:
With the default budget, the second commit alone restores the route up to a 449k-token bound; the first is needed for smaller budgets and for longer models (a 1M-token bound means 262,144 keys).
A third commit repairs three tests in tests/v1/cudagraph/test_cudagraph_dispatch.py, the file the second commit extends, which fail on main: their MagicMock VllmConfig had no observability config, so the decode event tracer of CUDAGraphWrapper compared the elapsed time against a MagicMock threshold ("'<' not supported between instances of 'float' and 'MagicMock'"). The mock now carries a default ObservabilityConfig, as other config tests do.
Not a duplicate: gh pr list --repo 1CatAI/1Cat-vLLM --state all --search "indexer decode", "cuBLAS indexer", "workspace budget", "graph_contract", "DeepSeek-V4 long context" and "131072" find no PR that changes this route or the bucket list; #874 introduced the buckets, #714 (closed) proposed the cuBLAS decode route. Open PRs touching the same files (#1190, #1192) change other parts (codec ids in the indexer kernels, tests) and merge without conflicts.
AI assistance was used for this change. I reviewed every line and ran the tests and measurements below.
Test Plan
New tests:
test_decode_cublas_score_tiles_match_one_tile(a 3 MiB budget splits 6,144 keys into six tiles; logits within FP32 rounding and identical top-512 keys per row against one tile, also captured and replayed as a CUDA graph; a budget without room for one tile falls back to the paged kernel) andtest_dsv4_decode_context_buckets_end_at_the_model_length. Two dispatcher tests now expect the model-length bucket.Serving on our fork (main 6f3557c plus our commits, MRV2 with the DSpark speculator of #1221): DeepSeek-V4-Flash MXFP4/FP8 with DSpark (num_speculative_tokens 5), PP5 over three Tesla V100 and two Quadro RTX 8000, FP8 KV cache, max_num_seqs 1,
VLLM_SM70_INDEXER_PREFILL_TILE_MB=64. Arms: fix (this PR on the fork, PYTHONPATH on the fork tree, own compile cache) against the same entry without it; boots fix cold, fix, base, fix, plus two earlier base boots with the same probe. Per boot: one warm-up request (32,668 tokens), then 178,479 and 287,077 tokens from repository documents and sources, greedy, 400 output tokens each. Step time = decode time / engine decode steps (vllm:spec_decode_num_drafts_total), independent of the acceptance length.Path coverage
Acceleration and benchmark contract (required for performance changes)
Test Result