Repository navigation
Conversation
Add an opt-in decode executor for small decode batches of
DeepSeek-V4.1-Flash on GB300 (TP4, DSpark), enabled with
--additional-config '{"decode_backend": "latency_specialized"}' (the
default, auto, keeps the stock decode), under
vllm/models/deepseek_v41/nvidia/specialized_decode/. As in vllm-project#61867, the
model owns the executor: the model state builds it once the weights are
loaded, and FULL-graph captures of pure decode steps run their decoder
layers on it. Its paths:
- a fused Lamport all-reduce + mHC at every sublayer boundary up to 384
rows, writing the attention input as MXFP8, collecting wo_b's
published output and taking the deferred MoE finalizes;
- GEMMs for 1-8 rows: a CuTe DSL routed-expert front, route kernel and
tcgen05 GEMMs; FlashInfer's low-latency MXFP8 GEMM with pinned
tactics for the attention projections; a low-batch BF16 GEMM for the
indexer projections; a CuTe DSL LM head and the custom all-gather for
logits;
- indexer and attention: shared sparse-index tables, a persistent KV
counter for the sparse-MLA decode, cooperative top-k, side-stream
indexer work and a row-sharded dense indexer;
- attention-metadata builds replayed from CUDA graphs, and the Engram KV
projection chain on the prefetch stream.
The executor is validated on GB300, MRV2, TP4 and DSpark; any other
deployment raises at start-up on every TP rank with the requirements it
misses. With the default, the model runs as before.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: Yaoyao Ding <yaoyaod@nvidia.com>
- Rename specialized_decode/README.md to AGENTS.md, like the other package guides. - Explain why the eligibility check takes a MIN all-reduce over the TP group: NVLink multicast is probed on each rank, and ranks that disagree would hang. In AGENTS.md and at the all-reduce. - Drop test_decode.py, whose tests restate the code, and its entries in the CPU CI step. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Yaoyao Ding <yaoyaod@nvidia.com>
- vLLM main now pins flashinfer-python 0.7.2. The FlashInfer internals that the publishing wo_b GEMM, the MXFP8 layer input and the pinned MXFP8 tactics use are unchanged in 0.7.2, as is the sparse-MLA decode's KV-counter argument, so the four features run on 0.7.2 instead of falling back to the stock path. - LLBf16Gemm.dispatch now takes the output dtype (vllm-project#62042); the attention side projections' AuxProjectionGemm, which overrides it, passes it through. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Yaoyao Ding <yaoyaod@nvidia.com>
With decode_backend latency_specialized, the DSpark draft of DeepSeek-V4.1-Flash gets a decode executor of its own, which the target's model state builds on the target executor's kernels and streams. The speculator's per-step work is fused, captured in a CUDA graph or moved off the critical path: - the draft layers on the fused all-reduce + mHC, the executor's attention and the routed-expert and MXFP8 GEMMs, with a CuTe DSL draft MoE finalize; - the Markov loop as one sampling kernel per draft position, and the whole chain in one cooperative CuTe DSL kernel per rank with the Markov head sharded by vocabulary; - sampling, postprocess_sampled and the draft input preparation in a post-forward CUDA graph, which a model state offers through post_forward_cudagraphs, with the draft context work on a side stream; - all draft KV-cache groups prepared in one launch. This applies to every DFlash, DFlash2, LiLiCorr and DSpark user and writes the same outputs. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Yaoyao Ding <yaoyaod@nvidia.com>
Resolve the conflict in .buildkite/test_areas/models_basic.yaml: keep the DeepSeek-V4.1 B200 step and main's new Kimi-K3 MI355 step. Signed-off-by: Yaoyao Ding <yaoyaod@nvidia.com> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…latency-decode Follow vllm-project#61868, which merged vLLM main. Signed-off-by: Yaoyao Ding <yaoyaod@nvidia.com> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Follow main after vllm-project#61868 landed. The specialized decode files take this branch's versions (main's equal the vllm-project#61868 revision this branch carries); .buildkite keeps both main's new test and this branch's dspark source dependency. Signed-off-by: Yaoyao Ding <yaoyaod@nvidia.com> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
|
This pull request has merge conflicts that must be resolved before it can be |
5ea6ad6 to
1b01835
Compare
|
Hi! We've been looking at the same draft-token gather overhead and found that this change in When you rebase it, note that #64458 changed broadcast = self.pp_handler is not None and num_spec_tokens > 0
if broadcast or input_batch.has_structured_output_reqs:
draft_tokens = self.req_states.draft_tokens[input_batch.idx_mapping]
else:
draft_tokens = self.req_states.draft_tokens[: input_batch.num_reqs] # width only
self.draft_tokens_handler.set_draft_tokens(input_batch, draft_tokens[:, :num_spec_tokens])
if broadcast:
self.pp_handler.broadcast_drafts(draft_tokens)A couple of tests would help: check that the placeholder width matches K (0, partial, max), and that PP still sends the full width when K < max. Happy to help review. Thanks! |
1b01835 to
3f557e6
Compare
3f557e6 to
7d00395
Compare
Per review, keep this PR out of the shared spec-decode and runner code: - Remove the post-forward CUDA graph (post_forward_cudagraph.py and its Model Runner V2 hook) and the draft context side stream it drove. - Restore the DFlash speculator, including the per-group input preparation, and the DSpark speculator to main's versions. - Remove the generic fused Markov sampler (fused_sample.py). - Move the vocabulary-sharded Markov chain under the model's specialized decode package (nvidia/specialized_decode/sampling/). It is wired up again through the draft model in a later commit. Signed-off-by: Yaoyao Ding <yaoyaod@nvidia.com> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Resolve two conflicts under vllm/models/deepseek_v41/nvidia/: - dspark.py: keep main's dspark_fused_markov opt-in (vllm-project#64442) beside the draft executor's registration, and main's compute_draft_logits_local. - moe/routed_expert_gemm.py: take current_cu_stream from main's shared moe_tail_fusion primitives (vllm-project#64974) and keep the finalize kernel. Signed-off-by: Yaoyao Ding <yaoyaod@nvidia.com> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
With decode_backend=latency_specialized, the DSpark draft model now samples through the model hook that DSparkSpeculator._sample_sequential already asks first, so the speculator itself is unchanged: - The specialized decode implies dspark_fused_markov: the fused loop of ops/fused_markov.py (vllm-project#64442), under the same conditions. - Once the draft's executor is built, batches of up to 8 request rows (4 at 8 ranks) take the vocabulary-sharded Markov chain instead (ShardedMarkovSampling in specialized_decode/sampling/). Its set-up is agreed over the TP group, as before. The layers test checks the hook with the switch on and off. Signed-off-by: Yaoyao Ding <yaoyaod@nvidia.com> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…ison The routed-expert test compares the reference FMA top-k reduction with FlashInfer's finalized output and then the draft's finalize kernel with the reference. Check the finalize kernel first, so its bitwise result is still checked when the FlashInfer comparison fails: with FlashInfer 0.7.2.post1 it does, because FlashInfer now adds the top-k products in another order (vllm-project#63921 relaxes that comparison). Signed-off-by: Yaoyao Ding <yaoyaod@nvidia.com> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: Yaoyao Ding <yaoyaod@nvidia.com> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: Yaoyao Ding <yaoyaod@nvidia.com>
dc91a1d to
4a8a438
Compare
Lower the step time of the DeepSeek-V4.1 latency_specialized decode at TP4, concurrency 1, with DSpark on GB300. Every change is unconditional; there are no new env vars or config options. Boundaries and MoE: - The fused all-reduce + mHC consumers let their PDL dependents launch early (the MXFP8 consumer at its start, the plain one after its poll, the low-batch finalize-publish after its wait). - A boundary's main-stream successor is captured before the side-stream kernels forked at it. - The boundary before the MoE also writes the MoE input as MXFP8 for the shared expert's gate_up. - The persistent routed-expert GEMMs leave 16 SMs to the shared expert, and GEMM1 lets GEMM2 launch after 21 k tiles. Attention: - The CuTe DSL sparse-MLA decode ropes and quantizes the BF16 query itself, and the Q/KV norm writes the KV cache, so the attention follows wq_b in the PDL chain; both are bitwise equal to the CUDA kernels they replace. - The sparse-MLA decode loads its graph-invariant inputs before its grid wait, reads an index source's request-local top-k without a table build in between, and a side stream prefetches the KV rows of later layers into L2. - The candidate consumers' top-k runs on topk_bf16_small, JIT-built at warm-up; the candidate source's block selection moves to the candidate side stream after the token top-k. - Indexers outside the decode plan reset their top-k rows early. Signed-off-by: Bowen Wang <bowen@inferact.ai> Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
4a8a438 to
ebcc07f
Compare
Overview
Lowers the step time of the DeepSeek-V4.1-Flash
latency_specializeddecode at TP4, concurrency 1, with DSpark on GB300. All changes are unconditional: no new env vars or config options, and they also apply above concurrency 1 (see Limitations).Stacked on #61869: one commit, ebcc07f, on its head bc6b0f2, which has merged main 5415568 (with #63921). Review that commit. The two Model Runner V2 changes that apply beyond this executor (the FULL graph's sliced output copy and the skipped draft-token gather) moved to #64701. This PR no longer changes the shared runner or speculators: #61869 removed the post-forward CUDA graphs on review, and the changes here that extended them went with them. Outside the executor it adds
skip_topk_buffer_clear(default False) to the two shared sparse indexers, optional arguments toquery_quant.pyandall_reduce_mhc.pythat default to the old behaviour, and a package-data line insetup.py.Claims
topk_bf16_small) is not bitwise equal to the code it replaces (the same set up to ties at the k-th value, in another order). GSM8K reads 0.45 pt lower with this PR (0.9050 → 0.9005, four runs per arm, every PR run below every base run). The evidence for a real loss is mixed: a per-question sign-flip test gives p = 0.15, a permutation test on the run totals p = 0.029 (0.083 exchanging runs only within jobs), and the 95% interval on the difference is about −1.0 to +0.1 pt. The other changes that touch a kernel's arithmetic are bitwise equal to the code they replace; the rest moves work between streams, changes launch timing or the routed GEMMs' CTA count, or adds a read-only L2 prefetch.Validation
Setup: 1 GB300 tray (4 GPUs), TP4,
--max-num-seqs 1,--numa-bind, DSpark k=5, synthetic acceptance length 3.51. vLLM nightly image dd4e401 (FlashInfer 0.7.2.post1) with each arm's Python package mounted over it, both without #62172's Python hunks (the image predates its kernels). AgentX 1.0.12, 1500 s, seed 42.One job on one tray, in the order PR, base, PR, base; scored on the 193 requests with an inter-token latency that the base runs completed, all of which every run also completed. Mean ± sd of two runs per arm: P90 891.9 ± 3.3 → 964.0 ± 3.0 tok/s (t ≈ 23), median 921.3 ± 2.6 → 999.3 ± 0.3 tok/s (t ≈ 42). AgentX runs at temperature 1 with synthetic rejection. Apart from #62172's Python hunks, each arm's package has the same
vllm/tree as its commit. The specialized decode re-tunes its FlashInfer GEMMs at every start, which can move the step time by up to about 1%; both arms loaded the same pinned tactics before tuning through a benchmark-side hook (not part of this PR).On the previous stack (#63921 + #61869 with its post-forward graphs, nightly 3966924), an earlier revision of this PR that also extended those graphs gave 907.5 → 986.5 and 986.1 / 936.0 → 1017.0 and 1017.4. Both arms read about 1.6–2.3% lower on this base; that comparison crosses trays, nightlies and the removal of the post-forward graphs, and is not attributed further.
Accuracy (GSM8K, real DSpark acceptance, 5-shot, 1319 questions, temperature 0; one run's standard error about 0.8 pt):
Paired by question, the base and this PR differ in which questions they get right on 59–75 questions per run pair (only the base right / only this PR right: 34/25, 35/25, 34/30, 40/35), about as many as two runs of one build (base 33/28, this PR 35/34). McNemar exact p = 0.25 to 0.71.
Two more runs per arm, in earlier jobs (trays 11, 15 and 17), on an earlier merge of main and #61869 with this PR's changes (the same diff apart from one
AGENTS.mdline), read 0.9075, 0.9030 → 0.9014, 0.9014. Over all eight runs every PR run reads below every base run, 0.9050 → 0.9005 (−0.45 pt). A sign-flip test on each question's right counts gives p = 0.15 (24 fewer right answers out of 5276; 133 questions have unequal counts; paired standard error 0.30 pt, 95% interval −1.0 to +0.1 pt). A permutation test on the eight run totals gives p = 0.029 (2 of 70 splits; 0.083 exchanging runs only within jobs). The data are inconclusive: they do not establish a loss, and they do not rule out one of up to about 1 pt. Earlier revisions of this PR on other stacks read 0.896 and 0.898, with no base run on those stacks to pair them with.topk_bf16_smallis the only change here with its own selection rule: on layers 24, 28, 32 and 36 it replaces DeepSelect's kernel, and among columns equal to the k-th value it keeps the lowest-numbered ones, which need not be the ones DeepSelect keeps.decode_topkon layer 20 keeps its rule, but it breaks ties by CTA arrival order in both arms, so this PR's launch changes there can change which tied columns it keeps. Real indexer logits hold few distinct BF16 values, so the k-th value may often be tied; neither that rate nor either effect was measured. GSM8K is the only eval run on this PR.Tests (GB300):
Run on ebcc07f with nightly dd4e401's image, the branch's Python package mounted over it, without #62172's Python hunks (the image predates #62172's kernels, so it cannot run them).
test_numerics.pyran with a stand-in forray, which the image lacks and which it only imports.test_shared_query_quant_preserves_projectionfails on main on GB300 (24 cases,AttributeError: params_dtype: #64111 made the kernel read an attribute the test's hand-built layer lacks), so this PR sets it there.test_layersruns 5 layers with random weights through the executor and compares each with the native path within tolerance. The new tests compare the fused query RoPE, the KV insert in the norm and the local-index attention bitwise with the kernels they replace, andtopk_bf16_smallexactly with its specified output (torch.topk's k-th value, the selected columns ascending, ties toward the lowest column).CI. No CI agent has four GB300s. The nightly 4xB200 step runs the kernel, numerics, top-k and sparse-MLA tests, and
test_layersskips there;tests/kernels/coreruns on H200, where the new KV-insert test skips. The GB300 runs above cover that gap.Serve and benchmark commands
GSM8K used
"rejection_sample_method":"block"(real acceptance) andtests/evals/gsm8k/gsm8k_eval.py-style prompting (5-shot, 256 tokens; the client sends 4 at once and--max-num-seqs 1runs them one at a time).Details
Where the time went. An nsys trace of an earlier stack (#63921 + #61869) at 3.73 ms per step showed these costs:
wq_band the attention.Per-layer medians on this base, from one trace per arm:
Layers 2, 8 and 14 keep the stock indexer's top-k (attention half 68–78 → 61–77 µs); layer 20's block selection moves to the side stream (L20 70.4 → 61.7 µs) and layers 24, 28, 32 and 36 take
topk_bf16_small(about 56 → 48 µs). Layer 1 reads slower in the PR's trace (95.3 → 117.9 µs); the total is net of that. The MoE half moves between traces of identical code (31.6 and 34.4 µs per plain layer in two earlier traces), probably with the number of experts the window hits.The 320 µs saved across the layers (one trace per arm; identical code has moved this sum by up to about 110 µs, in the MoE half) is in line with the roughly 300 µs per step behind the AgentX ITL (0.085 ms per token at acceptance length 3.51), and a fixed 6000-token decode without the tracer reads 1.02–1.07 → 0.95–0.97 ms per token in the same runs. The traced whole step moves less (3757 → 3579 µs): in the PR's trace rank 0's GPU catches up with its host in the metadata chain before the next step, so it runs no kernel for about 216 µs per step after the last layer (59 in the base; gaps over 2 µs, copies not counted as busy), and the other ranks wait for it in the step's first all-reduce (about 170 µs, against 5–13 in the base). Outside the trace window the same request's chunk gaps read 3.76 and 3.71 ms in the base and 3.42 and 3.36 ms in the PR (inside it, 3.75 and 3.59), so that floor binds only under the tracer; whether the untraced step has a lower host floor of its own was not measured.
The changes and what each bought (measured during development, when each was still switchable; single starts without pinned tactics, so differences under about 40 µs/step are within start-to-start tuning noise)
gate_upwq_bvia PDLtopk_bf16_small)Probe = per-step median of a fixed 6000-token decode; nsys = one trace.
Third-party code.
topk_bf16_smallis adapted from SGLang (Apache-2.0), whose kernel adapts DeepSelect's init-window select (MIT); the file carries both notices. It is JIT-built withtorch.utils.cpp_extensionat warm-up (on the first start, then cached), likevllm/models/deepseek_v41/amd/ops/topk.pybut with no fallback: without nvcc and a host compiler the executor's warm-up fails (itsAGENTS.mdstates the requirement). An earlier revision of this PR also vendored a bf16 cluster top-k for the dense indexers' token top-k; it is out, to keep the vendored tree out of this PR, and those layers stay ondecode_topkor the stock indexer as in the base. On the previous stack, without #62172, it was worth about 1.5% P90 / 1.2% median in one run against two; it would need measuring again. The ahead-of-time DeepSelect build vLLM pins does not fit as is: its gather signals completion barriers in the wrong CTA (deepseek-ai/DeepSelect#16), and a row holding a NaN gets a marker the sparse-index table build does not mask.Limitations
test_layerscovers correctness up to 12 requests.test_sparse_mla_cute.py(the RoPE) andtests/kernels/core/test_fused_q_kv_rmsnorm.py(the KV write) catch a compiler change that breaks that.Not a duplicate. Related open PRs make other changes, some in the same files: #64146 (TP8 on GB300), #64144 (the fused all-reduce + mHC without NVLink multicast) and #64893 (H200). Searches for
latency_specialized, "specialized decode", "DSv4.1 TP4" and "DSv4.1 top-k DeepSelect" found no open PR with these changes.AI assistance. Claude was used to profile, design, implement and benchmark these changes and to write this description.
Pull Request Checklist
I used vLLM's
/pr-checklistskill. (Mandatory for agents, optional for humans).AI assistance was used during the creation of this PR.
Design Fit: Minimizes impact on core components, reuses existing functionality, and justifies added complexity.
Testing and Validation: Validates the change and ensures any added tests are meaningful and reliable, with CI coverage or documented CI resource constraints and validation performed outside CI.
Code Quality and Style: Keeps code and comments clear and concise, and updates relevant documentation and examples.
Pull Request Contents: Includes a brief summary and relevant links, supports claims with evidence, explains root causes and implementation trade-offs, and follows the contributing guide.
🤖 Generated with Claude Code