Skip to content

[Perf][DSv4.1] DeepSeek-V4.1-Flash specialized decode on Hopper (H200 TP4) - #64893

Draft
mgoin wants to merge 16 commits into
vllm-project:mainfrom
mgoin:mgoin/dsv41-hopper-specialized-decode
Draft

mgoin wants to merge 16 commits into
vllm-project:mainfrom
mgoin:mgoin/dsv41-hopper-specialized-decode

Conversation

@mgoin

@mgoin mgoin commented Oct 31, 2026 •

Copy link
Copy Markdown
Member

Overview

Add a Hopper (SM90) backend to DeepSeek-V4.1's specialized decode executor (#61868), for DeepSeek-V4.1-Flash at TP4 on H200 with DSpark k=5. It is opt-in, as on GB300, through --additional-config '{"decode_backend":"latency_specialized"}'. Decode steps of up to 192 rows run in the executor; larger steps, prefill and anything it does not serve fall back to the stock path over the same weights.

Stacked on #61869 (DSpark low-latency decode), #57770 (split-vocabulary greedy argmax) and #64892 (SM90 mHC overlap behind the fused all-reduce). Pure TP4 with Humming MoE also needs humming#177 (humming-kernels 0.1.19, #64802).

Reviewing: the Hopper executor commit is the new work, plus the merge of main on top of it, which adapts the stack to main's DraftLogitsCache and the moved CuTe DSL primitives. Everything before them comes from #61869 and #64892. #57770 is merged in main. #61869 is still open, so its changes are in this diff until it lands.

Duplicate check: there is no other open Hopper backend for the specialized decode. #62958 and #64442 only touch DSpark sampling.

Claims

  • Median ITL on H200 TP4, DSpark k=5, on this branch merged with main: −38% at c1 (8.65 → 5.34 ms) and about −5% at c8 on real prompts; −37% at c1 and −14% at c8 at 64K context. At c16 and above the gain was about 5–7% before the merge and has not been re-measured since.
  • Quality is unchanged within noise on GPQA-Diamond, AIME25 and 128K needle retrieval; greedy logprob drift is within the stock run-to-run floor.
  • No host syncs in steady-state decode (VLLM_GPU_SYNC_CHECK=error, re-run on the merged build: 0 at c1, c8 and c32).
  • The Hopper commit changes no core components (model runner, scheduler, KV cache manager); the model_runner.py changes in this diff come from [Perf][DSv4.1] DSpark draft on the specialized decode #61869. Everything in the Hopper commit is under vllm/models/deepseek_v41/nvidia/specialized_decode/, and SM100 behavior is unchanged.

Validation

Setup: 4× H200, --tensor-parallel-size 4 --max-model-len 200000 --max-num-seqs 64, DSpark k=5, Humming MoE with humming#177. "Off" is the same build with the executor disabled, so it includes every stacked PR. Runs are in ABBA order (on, off, off, on).

Speed, real prompts, output length 1024, temperature 0.7. Median ITL in ms:

conc 1 2 4 8 16 24 32
executor off 8.76 10.31 12.81 16.49 21.26 26.61 29.83
executor on 5.40 6.88 10.57 14.27 19.84 25.02 28.31
Δ −38.3% −33.3% −17.5% −13.5% −6.6% −6.0% −5.1%

Acceptance lengths match between arms (2.79–3.15).

Long context, 64K prompts, two on/off pairs. Median ITL in ms:

conc on off Δ
1 5.64, 5.65 8.96, 8.96 −37%
8 15.20, 13.58 15.59, 16.15 about −9%

Quality.

Eval on off on − off (paired)
GPQA-Diamond, 8 × 198 79.9 80.6 −0.6 ± 0.6
AIME25, 8 × 30 86.2 85.8 +0.4 ± 1.8
128K needle, single and multikey, c1 and c8 12/12 each 12/12 each
  • On the final build, GPQA-Diamond at 3 repeats scores 80.0.
  • Greedy logprobs at identical prefixes, 24 prompts:
    • on vs off: mean |Δ logprob| on the chosen token 2.2e-2;
    • off vs off (noise floor): 2.6e-2;
    • every divergence happens at a top-1 margin of at most 0.25 nats.

The concurrency sweep and the 8-repeat evals ran before a fused-indexer fix at long context. On the final build c1 and c8 measure 5.39 and 14.40 ms, unchanged.

After merging main (f1df4a3ae9). Median ITL in ms, on a new 160-prompt set from RedHatAI/speculator_benchmarks:

Point on off Δ
real prompts, c1 5.34 8.64, 8.66 −38%
real prompts, c8 14.36 14.64, 15.65 about −5%
64K context, c1 5.63 8.89 −37%
64K context, c8 13.63 15.93 −14%
Check on off
128K needle, single and multikey, c1 and c8 12/12 each 12/12 each
GPQA-Diamond, 3 repeats 79.3 (80.3 / 79.3 / 78.3) not re-run

Only one on run per point is clean. The tables above this one are from before the merge and used an earlier prompt set, so compare numbers only within a table. The c8 gain is smaller here than the −13.5% in the sweep above. The cause isn't established: main, the prompt set and acceptance length (3.25 on, 3.18–3.50 off) all changed, and the runs are few. A clean c1–c32 sweep on the merged build is still to do; the attempt so far overlapped other work on the GPUs and was discarded.

Tests. pytest tests/models/deepseek_v41/nvidia/specialized_decode/test_hopper.py on H200: 19 passed.

  • The tests compare each Hopper kernel and step against the stock path or an FP32 reference.
  • The fused indexer has a long-context case: one indexer shared across ratio-1 and ratio-2 layers, rows longer than the candidate coverage, and a NaN-poisoned output.
  • test_hopper is wired into Buildkite (models_basic.yaml).
  • CI has no 4× H200 runner for the end-to-end numbers above.

Reproduce:

vllm serve deepseek-ai/DeepSeek-V4.1-Flash --tensor-parallel-size 4 \
  --max-model-len 200000 --max-num-seqs 64 --tokenizer-mode deepseek_v41 \
  --speculative-config '{"method":"dspark","num_speculative_tokens":5}'
vllm bench serve --backend vllm --endpoint /v1/completions --dataset-name custom \
  --dataset-path real_prompts.jsonl --skip-chat-template --custom-output-len 1024 \
  --temperature 0.7 --max-concurrency 1   # likewise 2, 4, 8, 16, 24, 32

The executor-on arm adds --additional-config '{"decode_backend":"latency_specialized"}' to vllm serve. The executor-off arm is the default.

Details

Where the time goes. At c1 the stock step is mostly fixed per-kernel cost, not bandwidth:

  • an isolated ~10 MB GEMV reaches only about 50% of HBM;
  • each kernel boundary costs ~0.7–1.8 µs.

So the executor:

  • chains kernels with PDL;
  • replaces the small-batch MoE, dense FP8 linears, sparse MLA and indexer with SM90 CuTe DSL kernels written for these shapes;
  • fuses the small ops around them.

The MoE runs near its streaming limit (about 4 TB/s).

Compatibility. The kernels read the weights in stock's layouts:

  • Humming MXFP4 × block-FP8 for experts;
  • Humming W8A16 for dense layers.

So steps the executor doesn't serve fall back to stock with one copy of the weights. The exception is an FP8 copy of wo_a, because stock SM90 keeps it only as BF16. It costs about 0.36 GB per GPU (1.44 GB in FP8 across 43 layers, sharded TP4).

Size. About 10.8K lines under hopper/, mostly kernels:

Area Lines
MoE 4.4K
sparse MLA 2.3K
dense GEMM 1.6K
indexer 0.9K
attention-chain fusions 0.7K

Plus +342/−85 in the shared specialized-decode package, which lets the executor select its backend by architecture and records metadata builds for the Hopper builders. Plus 1.0K lines of tests.

It mirrors the GB300 backend's structure, so the two share the executor runtime, step schedule and fallbacks.

Limitations.

  • TP4 on H200 only.
  • Steps above 192 rows use stock.
  • The gain shrinks with batch size, as the per-kernel overhead is amortized.

Pull Request Checklist
  • I used vLLM's /pr-checklist skill. (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

Juntian777 and others added 15 commits September 20, 2026 08:25
…edy draft

The DSpark greedy drafter selects a token per speculative step with a
full-vocabulary torch.argmax. At the shapes the drafter runs at (a few
rows, 129280 columns) that kernel is launch- and latency-bound, costing
~43 us per call regardless of the row count, and the five steps are
dependency-chained so the cost lands five times on every decode step.

Replace it with a two-stage split-vocabulary reduction: per-block partial
argmax, then a small final reduction. The comparator defines a total
order (NaN highest, then value, then lowest index), so the result is
independent of the reduction order and bit-exact with torch.argmax,
including first-index ties and NaN propagation. Unsupported inputs fall
back to torch.argmax.

GB200, BF16, vocabulary=129280, CUDA Graph + CUPTI, cold L2: 43.4 -> 4.4 us
at M=1 and 44.4 -> 8.0 us at M=32.

Signed-off-by: Juntian Liu <Juntianl777@gmail.com>
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Drop the CUDA Graph replay and fallback cases, and add random rows plus a
row whose first maximum lies after the first block with ties in later
blocks, so the cross-block first-index tie-break is exercised.

Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Signed-off-by: Juntian Liu <Juntianl777@gmail.com>
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>
…nto main 51d2a85

Signed-off-by: mgoin <mike.goin12@gmail.com>
…elSpeculator

Address review: both reduction stages use tl.argmax with NaN mapped to
-inf (as in the rejection sampler), and the generic DraftModelSpeculator
greedy path uses greedy_argmax too.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: mgoin <mgoin64@gmail.com>
Signed-off-by: mgoin <mike.goin12@gmail.com>
…n SM90

Enable the mHC coefficient overlap on Hopper when the fused Lamport
all-reduce + mHC kernel applies (TP4, NVLink multicast). On H200 TP4 the
fused path wins through 144 tokens, against 16 on GB200, so the token cap
becomes a per-architecture mhc_overlap_max_tokens().

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: mgoin <mike.goin12@gmail.com>
… TP4)

Run decode_backend=latency_specialized on SM90: HopperSpecializedDecode
reuses the executor's plan, layer loop, boundaries and side streams over
the stock FlashMLA layers, for FULL-graph decode steps of up to 192 rows
(32 DSpark requests), with Hopper kernels for the rows they serve:

- dense FP8 linears on a low-batch CuTe DSL GEMM (X multicast in
  clusters at 6-16 rows) or tuned Humming, wo_b publishing into the
  fused all-reduce + mHC mailbox;
- a persistent low-batch MoE kernel on the stock Humming W4A8 experts
  (routing, both expert GEMMs and the shared expert), finalized by the
  next boundary, up to 144 rows;
- sparse MLA sharing FlashMLA plans and top-k tables across layers, and
  small-batch MLA kernels with early-trigger q/kv norm, query RoPE and
  combine + inverse RoPE up to 12 rows;
- a fused FP8 indexer (logits + candidate bitmap, then the stock
  cooperative top-k) up to 48 rows;
- the DSpark draft on HopperDraftDecode.

Boundaries above the fused all-reduce + mHC cap take the stock
all-reduce and mHC. GB300 behavior is unchanged.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: mgoin <mike.goin12@gmail.com>
@mergify mergify Bot added ci/build deepseek Related to DeepSeek models DSv4.1 Related to DeepSeek-V4.1 models nvidia speculative-decoding labels Oct 31, 2026
@mergify mergify Bot added dflash mrv2 Model Runner V2 specific labels Oct 31, 2026
@mergify

mergify Bot commented Oct 31, 2026

Copy link
Copy Markdown
Contributor

This pull request has merge conflicts that must be resolved before it can be
merged. Please rebase the PR, @mgoin.

https://docs.github.com/en/pull-requests/collaborating-with-pull-requests/working-with-forks/syncing-a-fork

@mergify mergify Bot added needs-rebase agents Related to information for agents labels Oct 31, 2026
Follow main in the stacked DSpark and Hopper code:
- read the draft logits through DraftLogitsCache, and keep the fused Markov
  sampler to the replicated cache it indexes by request state;
- import the CuTe DSL primitives from vllm.models.common.ops.moe_tail_fusion.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: mgoin <mgoin64@gmail.com>
@mergify mergify Bot removed the needs-rebase label Nov 2, 2026
@mergify

mergify Bot commented Nov 3, 2026

Copy link
Copy Markdown
Contributor

This pull request has merge conflicts that must be resolved before it can be
merged. Please rebase the PR, @mgoin.

https://docs.github.com/en/pull-requests/collaborating-with-pull-requests/working-with-forks/syncing-a-fork

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

agents Related to information for agents ci/build deepseek Related to DeepSeek models dflash DSv4.1 Related to DeepSeek-V4.1 models mrv2 Model Runner V2 specific needs-rebase nvidia speculative-decoding

Projects

Status: No status
Status: Backlog

Development

Successfully merging this pull request may close these issues.

4 participants