[New Model][Nvidia] Add SM12x support for DeepSeek V4 Flash with essential fixes - #41834
Open
jasl wants to merge 288 commits into
Open
[New Model][Nvidia] Add SM12x support for DeepSeek V4 Flash with essential fixes#41834jasl wants to merge 288 commits into
jasl wants to merge 288 commits into
Conversation
jasl
requested review from
22quinn,
LucasWilkinson,
MatthewBonanni,
ProExpertProg,
WoosukKwon,
hmellor,
houseroad,
mgoin,
pavanimajety,
robertgshaw2-redhat,
tlrmchlsmth,
yewentao256 and
youkaichao
as code owners
May 6, 2026 15:17
![claude[bot]](https://avatars.githubusercontent.com/in/1236702?s=60&v=4){kind=small}
Contributor
Author
|
@zyongye |
![gemini-code-assist[bot]](https://avatars.githubusercontent.com/in/956858?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
Code Review
This pull request implements support for DeepSeek V4 on SM12x (Blackwell) architectures by providing Triton-based fallbacks for DeepGEMM-dependent operations. Key enhancements include the introduction of specialized Triton kernels for sparse MLA, FP8 einsum, and MQA logits, as well as memory optimizations in the sparse attention indexer to compute top-k indices without materializing full logits. Additionally, the PR updates the model loader to support weight name filtering for skipping MTP weights and handles Blackwell-specific FP8 quantization scales. I have no feedback to provide.
💡 Codex Reviewvllm/vllm/model_executor/layers/sparse_attn_indexer.py Lines 86 to 89 in 9596dbf  Keep DeepGEMM requirement for SM120 FP4 indexer path
This helper now disables the DeepGEMM requirement for every SM120 run, but the FP4 indexer cache path still depends on DeepGEMM kernels ( vllm/vllm/model_executor/model_loader/default_loader.py Lines 236 to 240 in 9596dbf  Propagate weight_name_filter to fast safetensor loaders
The new pre-load ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
If Codex has suggestions, it will comment; otherwise it will react with 👍. Codex can also answer questions or update the PR. Try commenting "@codex address that feedback". |
jasl
changed the title
jasl
requested review from
ApostaC,
alexm-redhat,
heheda12345,
njhill,
orozery and
ywang96
as code owners
May 6, 2026 15:54
jasl
force-pushed
the
codex/ds4-sm120-min-enable
branch
from
May 6, 2026 16:26
042e366 to
df2e6f8
Compare
The guard is `env AND use_eagle`, and use_eagle only reaches managers whose group is in eagle_group_ids, so setting the flag is not evidence the guard is live for a given model. Log the per-manager count at startup so an A/B of the flag is verifiable from the serve log instead of assumed -- with the env off it must report 0 active, which is the negative control. (The first version of this used `logger` without importing it, which would have raised NameError at engine startup and broken every serve in the A/B. Caught before running by checking that the name was actually defined at module level rather than trusting that a logger exists in every module.)
The startup log added in the previous commit paid for itself immediately. With
the flag set, DeepSeek-V4 reported:
Same-step ghost-block guard: 2/5 managers active (SlidingWindowMLAManager)
Upstream gates the guard on `use_eagle`, having framed the race as a spec-decode
problem. But KVCacheCoordinator only sets use_eagle on managers whose group is
in eagle_group_ids, and its fallback to "flag every group" fires only when NO
group is flagged. On DSv4 the sliding-window groups ARE flagged, so the fallback
never runs and MLAAttentionManager -- the main attention path, and the one the
recall gate exercises -- was left completely unguarded.
The race is in block publication: a hash is committed to the shared BlockPool at
scheduling time, before the forward writes the KV. Nothing about that is
specific to speculative decoding. So the env var becomes tri-state:
0 = off, 1 = upstream semantics (use_eagle-gated), 2 = every group.
Without that log line this would have run a 45-minute A/B whose treatment arm
was 40% applied, produced no clear improvement, and invited the conclusion that
the mechanism was wrong -- when the patch simply was not enabled where it
mattered. Reading a condition back from the engine is worth more than setting it
carefully.
Setting the env default to 2 was tried and reverted after measuring what it
costs. tests/v1/core, with the guard defaulted on:
guard 0 (upstream default) -> 1 failed, 245 passed
guard 1 (upstream semantics) -> 3 failed
guard 2 (every group) -> 12 failed
The 11 extra failures are all in test_prefix_caching.py, which calls
allocate_slots repeatedly to represent SUCCESSIVE scheduling steps while calling
new_step_starts exactly once in the whole file. With the guard on, those
allocations look like one step and the hits are correctly deferred, so the tests
fail. The tests are step-agnostic rather than wrong -- but flipping the default
here would fork 11 upstream tests and break every future test written the same
way, for no benefit that setting the variable in the deployment does not already
give.
So the default stays 0 and the harness turns it on at serve time. The single
remaining failure, test_async_scheduling_pp_allows_rescheduling_with_output_
placeholders, fails identically on the pre-merge tree 4ebd1fb and is not
related to this work.
Also moves `from vllm import envs` above `vllm.utils.math_utils` so the import
block stays isort-clean.
Verified before committing, rather than assumed:
- new_step_starts is called unconditionally at the top of Scheduler.schedule and
the coordinator forwards it to every manager, so cached_blocks_this_step
cannot grow across steps or defer a request permanently.
- envs caching is enabled in the serving path (engine/core.py, multiproc_
executor.py), so the property does not do an os.getenv per scheduling call.
…hable Restores upstream's dspark -> V2 routing, which this fork removed on 2026-08-03. That removal was correct on the evidence available then: V2 lost long-context recall under concurrency, 8 samples per runner from one serve giving V1 mean 22.25 [20,24] against V2 mean 10.50 [6,13], Mann-Whitney U=0. That collapse was not a property of the runner. It was the same-step ghost-block race in the prefix cache (vllm-project#42359): a block's hash becomes visible to other requests before the forward pass writes its KV, and a serve that loses the race keeps serving from the poisoned blocks for its lifetime. V2 was the heavy casualty only because its larger KV pool admits all twelve concurrent requests at once (Waiting peaks at 0 against V1's 11), which maximises the number of racing writers. With the guard on, 4 fresh serves per runner, 3 arthur c=12 runs each, same tree, same guard mode, runner the only variable: V1 serve means 22.3 / 21.7 / 21.7 / 21.0 gate mean 21.67 min 19 V2 serve means 22.0 / 20.7 / 22.7 / 22.7 gate mean 22.00 min 20 Mann-Whitney p = 0.697; neither side has a single-digit serve and on the full suite V2 leads everywhere else it can be measured: pp2048 +3.1/+4.2/+7.1% at d8192/16384/32768, tg128 +4/+20/+28%, TTFT lower at every depth, KV +24.9%. GSM8K, issue19, multi-needle and arthur c=1 all tie. Two corrections recorded in the comment rather than quietly dropped: the "+6.6% draft acceptance" advantage does not survive re-measurement (2.772 V1 vs 2.710 V2 on the same formula and sample size -- it was measured while V2 was poisoned), and V1's own 2-4 needle shortfall was the same defect, not an inherent concurrency margin: the guard lifts V1 from 20.7 to 23.0. V1 remains fully supported. VLLM_USE_V2_MODEL_RUNNER=0 forces it, and the env check precedes every routing rule. A test pins both halves -- the default and the escape hatch -- so neither can drift silently. Regression sweep with V2 defaulted: 261 passed, 1 failed, where that one failure reproduces identically on the pre-merge tree 4ebd1fb.
The previous two commits were individually defensible and jointly wrong. Making V2 the default runner while leaving the guard's default at 0 meant a plain serve -- nothing set -- got V2 + prefix caching + DSpark with no guard, which is precisely the combination measured at arthur c=12 mean 11.5, 3 of 4 serves degraded, floor 3/24. Before this work the default was V1 without a guard at 20.7. I had made the out-of-the-box configuration worse while each change looked like an improvement on its own. Our own serve script set the variable, so our measurements never saw it; anyone following the PR would have. The fix couples the two decisions where they belong. The env var now distinguishes UNSET (None) from an explicit 0/1/2, and KVCacheCoordinator raises the default to 2 when prefix caching and speculative decoding are both on -- the only configuration where a hash can be published before its KV is written and a second request can be admitted in the same step to match it. Deliberately in the coordinator, not in envs.py: a manager constructed directly still resolves to OFF, which is what keeps the step-agnostic tests in test_prefix_caching.py passing. An explicit value always wins, so `=0` remains a real escape hatch -- pinned by a test, because a guard you cannot switch off is a guard you cannot rule out when diagnosing something else. Two eagle tests did start failing, since they build a coordinator rather than a bare manager: both allocate for req0 and then for req1 expecting a hit, with no step boundary between them. That is the race, not a cache hit. Each gains one `new_step_starts()` call, which makes them exercise what they are actually about. Sweep: 276 passed, 1 failed -- that one failing identically on the pre-merge tree 4ebd1fb. Verified end to end rather than by inspection: with nothing set, dspark routes to V2 and the coordinator turns the guard on; with `=0` set, it stays off.
Contributor
|
This pull request has merge conflicts that must be resolved before it can be |
Conflicts and how they were taken: - vllm/v1/core/sched/scheduler.py -> UPSTREAM. Upstream refactored the per-method lookahead chain into a single VllmConfig.num_lookahead_tokens property (vllm-project#51438). It is semantically identical for DSpark -- use_eagle() includes "dspark", so the property returns num_speculative_tokens exactly as our chain did -- and upstream's comment is our fork's wording, so the reasoning was absorbed rather than lost. Our side also carried a duplicated dspark branch that the refactor makes redundant. - requirements/cuda.txt, docker/{Dockerfile,versions.json} -> UPSTREAM's FlashInfer 0.6.16.post3 (we were on 0.6.16). apache-tvm-ffi stays 0.1.11, the pin we already run; the tilelang double-registration abort is triggered by moving tvm-ffi to 0.1.13.post0, not by the FlashInfer version, and 0.6.16.post1 was previously verified serving on 0.1.11. - tests/v1/kv_connector/.../test_scheduler.py -> UPSTREAM's added case. Verified after resolving, rather than assumed: all four of our changes survive (guard property, coordinator default, tri-state env, dspark->V2 routing), and Scheduler.schedule() still calls new_step_starts() unconditionally -- checked by AST for enclosing control flow, since the guard's per-step set is cleared there and a conditional call would make it grow without bound and defer permanently. Upstream commits of note for this branch: vllm-project#51438 reserves spec-decode lookahead blocks in V2 warmup, vllm-project#50365 drops atomic contention in the sparse-MLA index remap, vllm-project#48668 preserves prefix-cache stats on zero-output steps. NOT yet re-tested: the FlashInfer bump changes kernel_warmup.py and gpu_model_runner.py, and the fleet still has 0.6.16 installed. Acceptance runs after the fleet is upgraded.
The two endpoints serving the same model disagreed about its own vocabulary.
`ChatCompletionRequest.reasoning_effort` has always accepted `max` -- its
docstring even records that the tier is DeepSeek-V4-specific -- while
`ResponsesRequest.reasoning` took its type straight from the OpenAI SDK, whose
`ReasoningEffort` stops at `xhigh`. So `{"reasoning": {"effort": "max"}}` was
rejected by schema validation on /v1/responses and accepted on
/v1/chat/completions, for the same model, in the same server.
`max` is not a synonym: DeepSeek's V4 encoding ships a distinct prompt for it
(`REASONING_EFFORT_PROMPTS["max"]`), and DeepSeek's own API documents
none/low/high/max as the supported set.
No mapping is added here. `DeepSeekV4Tokenizer.apply_chat_template` already
folds every spelling onto the model's three tiers -- `none` disables thinking,
`minimal`/`medium` become `low`, anything else becomes `high` -- so widening
the schema is the whole fix. A first draft of this change added a second
normalisation table in `deepseek_v4_encoding`; it was dropped once the existing
one was found, rather than left in as a competing source of truth.
Tests pin the behaviour that had none: that both endpoints accept the same
seven spellings, that `max` specifically survives into `chat_template_kwargs`
without thinking being switched off on the way, that `none` still disables
thinking, and that the widened field still rejects a value the model has no
tier for. Against the unpatched tree three of them fail with ValidationError.
Verified end to end on a two-node TP=2 serve: `max` is accepted on the patched
tree and rejected on the unpatched one, with low/minimal/medium/high/xhigh/none
behaving identically on both.
With no thinking kwarg, `DeepSeekV4Tokenizer.apply_chat_template` defaults
thinking ON while `DeepSeekV4ReasoningParser` defaults it OFF and selects
`IdentityReasoningParser`. The model reasoned and the reasoning, with a bare
`</think>`, was returned inside `output_text` as though it were the answer --
on the default path, since omitting `reasoning` is exactly what a stock OpenAI
SDK does.
/v1/chat/completions was immune only because it normalises thinking into the
chat-template kwargs at the protocol boundary, and its docstring says why:
"so the tokenizer and reasoning parser see the same effective state".
/v1/responses never called that hook.
Rather than add a second copy of the derivation -- the duplication is how the
two endpoints came to disagree, twice now, this and `max` -- it moves to
`deepseek_v4_chat_kwargs` and both request types call it. `ChatCompletionRequest`
keeps its public methods, delegating; behaviour there is unchanged by
construction.
Tests: 26 new cases, 11 of which fail on the unpatched tree, covering both
checkpoints (DeepSeek-V4-Flash and -0731), an explicit thinking=false surviving
normalisation, `effort: "none"` still disabling thinking, and every effort
spelling being accepted. No regressions: tests/reasoning 440 passed,
tests/tokenizers_/test_deepseek_v4.py 45 passed,
test_responses_reasoning_effort.py 11 passed. The pre-existing collection
errors (`schemathesis`, `cohere_melody` absent) reproduce on the unpatched
tree.
This also retires the `--default-chat-template-kwargs '{"thinking":true}'`
workaround the deployment was carrying.
1 task
This was referenced Aug 12, 2026
Five conflicts. Three were positional -- both sides adding a different import
or an elif -- and two were real.
vllm/models/deepseek_v4/nvidia/flashmla.py: upstream moved an unconditional
`assert topk_indices_buffer is not None` into the non-swa_only branch. Kept
ours, which subsumes it: _prefill_workspace_topk_bound returns 0 for
compress_ratio <= 1 without touching the buffer, and falls back to the
indexer's topk (or 2048) when the buffer is missing rather than asserting.
vllm/v1/kv_offload/cpu/gpu_worker.py was the substantial one. Upstream added a
canonical CPU page layout while this branch carries the compact layout, and the
two collide across the constructor, the docstrings, the dispatch point, and the
stream/event lifecycle. Both are opt-in and orthogonal, so both are kept, with
three resolutions that needed more than "take both":
* Upstream dropped the `if self.gpu_to_cpu:` guard before
`stream.wait_stream(...)`, so loads wait on the compute stream too -- an
earlier transfer could otherwise be overwritten by already-queued compute
work. This branch had extracted that block into _submit_descriptors, so
taking our side would have silently dropped the fix; it is applied to the
extracted helper instead.
* The layout dispatch is now a chain -- compact short-circuits, then
canonical, then the worker view -- with an assert that the two layouts are
never combined.
* Upstream renamed kv_cache_groups_data_refs to layer_refs_per_group. Keeping
both assignments left `self.kv_cache_groups_data_refs = kv_cache_groups_data_refs`
referencing a parameter that no longer exists, which would have raised on
the first construction. The old name is gone from vllm/ entirely.
Also verified before resolving: upstream's new `mask_token_id` fallback in
llm_base_proposer sits ahead of the dspark_noise_token_id branch, but the 0731
checkpoint has no top-level mask_token_id (it uses dspark_noise_token_id
128799), so the DSpark path is unchanged.
CPU offload is not enabled in this deployment -- no flag in the serve script,
no mention in the engine log -- so the blast radius of the offload resolution
is the PR, not the running service. It still needs its own test pass.
Ahead of upstream, which is still on 0.6.16.post3. 0.6.17 carries the SM12x release content this fork exists for: the Blackwell fused-MoE refresh with its NVFP4 accuracy fix, MXFP4 W4A8/W4A16 through the unified MoE API, the SM120 gated-SwiGLU moe_gemm path, trtllm_allreduce extended to SM12x, and a native qk_rope_head_dim=0 sparse-MLA decode. Checked before taking it: the declared dependency sets of 0.6.17 and 0.6.16.post3 are byte-identical, so this is a version move rather than a new transitive-dependency surface. apache-tvm-ffi stays pinned at 0.1.11 and the reason is now written down. 0.6.17 asks only for >=0.1.6,<0.2 excluding 0.1.8 -- a range that 0.1.13.post0 also satisfies, and that version is ABI-incompatible with the tilelang build here. An unpinned resolve would take it. All three places that carry the version move together, which the comment in requirements/cuda.txt has always asked for and which is easy to half-do: requirements/cuda.txt, docker/Dockerfile's FLASHINFER_VERSION, and docker/versions.json. The remaining 0.6.16 string is a test fixture asserting that mismatched python/cubin versions are rejected, not a pin. Not yet built or run -- that follows on the two nodes still free.
…fload worker Both are in the file I hand-resolved during the 08-11 merge, and both come from the same mistake: keeping our side of a block while taking upstream's removal of what that block depended on. 1. transfer_async submitted the descriptor UPPER BOUND, not the count written. Upstream's writer rotation made the old `assert op_idx == num_copy_ops` inexact, so the resolution relaxed it to `<=` and truncated src/dst/sizes to op_idx -- then passed num_copy_ops to _submit_descriptors, which discards those views and re-slices batch_src[:num_copy_ops]. The truncation was dead code and the DMA received num_copy_ops - op_idx descriptors that were never written. The buffers come from torch.empty and are recycled through _buffer_pool, so the tail is uninitialized on a fresh buffer and holds a previous, larger transfer's live pointers on a reused one -- raw pointer copies of arbitrary length between arbitrary addresses. Reachable with canonical_layout + gpu_to_cpu + num_writers > 1, which is the replicated-MLA-under-TP shape (num_writers = tp_size // dcp_size). Upstream gates on op_idx; this now does too. 2. The handler's shutdown() read self._mmap_region, which it never assigns. Upstream moved mmap ownership to CPUOffloadingWorker (vllm-project#51622). The merge took upstream's removal of the handler's assignment but kept our handler-side cleanup, leaving an attribute that is read and never established -- AttributeError on every handler shutdown, swallowed by the caller into two spurious "Failed to shut down" tracebacks and a stuck handler_failed flag. The worker already does the real cleanup, so the block is deleted rather than the assignment restored, matching upstream's intent. Neither is reachable in our serving configuration -- KV offload is opt-in via kv_offloading_size, which we never set -- but this branch is a PR, and (1) is a memory-safety defect for anyone who does enable it. Also restores the blank line the merge ate before _submit_descriptors, so the file passes ruff-format again.
…dagraph mode Upstream vllm-project#51768 added _validate_mrv1_piecewise_cudagraph, which raises at config time for DeepseekV4ForCausalLM under Model Runner V1 with PIECEWISE cudagraphs. PIECEWISE is in the O2/O3 default (FULL_AND_PIECEWISE), so VLLM_USE_V2_MODEL_RUNNER=0 alone -- which our comment offered as the escape hatch, and which our acceptance suite exercises -- now fails to boot rather than selecting V1. The env still selects V1; it just is not sufficient by itself any more. Say so where the claim is made, and name the flag that completes it.
1 task
Every transfer-level case in this file builds mappings with num_writers=1, where the sized upper bound and the number of descriptors actually filled are equal -- so the store path that submits the bound instead of the fill is invisible to the whole suite. The one num_writers=2 test never calls transfer_async. Assert the count relationally (rotated store submits half of what an unrotated one does) rather than re-deriving the canonical page id arithmetic the test is supposed to be independent of.
Contributor
|
This pull request has merge conflicts that must be resolved before it can be |
Five conflicts, and the interesting work was in what did NOT conflict. CONFLICTS vllm/v1/attention/backends/mla/indexer.py -- each side changed a DIFFERENT argument of the same call. Upstream vllm-project#47808 relaxed require_uniform when varlen paged-MQA logits are available; ours added the sparse_short_extend_tiering gate that the fork's three other call sites use. Both kept. vllm/model_executor/layers/sparse_attn_indexer.py -- git aligned our fp8_fp4_paged_mqa_topk_indices argument tail against upstream's tail for fp8_fp4_paged_mqa_logits, a different function we call further down. Resolving "take ours" would have dropped upstream's new `indices=decode_metadata.indices` SILENTLY: the kwarg defaults to None, so nothing would have failed. Applied to the call it belongs to. vllm/v1/worker/gpu/spec_decode/dspark/speculator.py -- took upstream's load_draft_model, which restores a d2t scatter assignment the merge-base had and our line lost in an EARLIER merge while keeping its reader; _d2t_scatter_index was declared, read, and never assigned. Dead branch, not a crash, which is why nothing noticed. vllm/models/deepseek_v4/nvidia/flashmla.py -- both import blocks and both code blocks were needed: our d512-split helpers exist only here, and upstream's new SWA backend classes are referenced by a `swa_backend_cls` line that merged cleanly. Dropped `round_up`, which upstream imports for call sites this fork removed. vllm/models/deepseek_v4/nvidia/dspark.py -- six hunks. Ours flattened this file into one class, so most of upstream's side navigates a `self.model` that does not exist here. Kept our structure; ported upstream vllm-project#47808's semantics onto it: DSparkConfidenceHead, confidence-head load tracking, and the graceful-degrade null-out for checkpoints without the head. Did NOT import DSparkMarkovHead (we flattened it to markov_w1/w2) or maybe_prefix (our prefix already ends in a dot, so it would produce a double dot and a layer name that silently misses quant-config lookups). WHAT A CLEAN AUTO-MERGE HID Both sides' confidence heads were stacked at __init__ with NO conflict markers: ours constructed and then immediately overwritten by upstream's, whose line called maybe_prefix, which nothing imports -- a NameError at draft-model construction that no marker pointed at. Collapsed to one head of upstream's type built from our locals, reproducing our existing mtp.2.confidence_head.proj parameter path exactly. A LATENT BUG THIS SURFACED DSparkDeepseekV4ForCausalLM never declared draft_id_to_target_id. nn.Module's __getattr__ raises rather than returning None, so once upstream's restored load_draft_model read it, DSpark draft-model load would have died -- every DSpark serve, i.e. production. The AMD sibling has always declared it. Declared here too, and the speculator now accepts both the nested and flattened draft-model shapes instead of assuming `model.model`. Also removed a dead `has_prefilling_rows` in indexer.py: assigned, never read, and paying a blocking .item() device sync for it. Pre-existing, not from this merge. Lint: touched files carry no new ruff errors; the ones that remain tree-wide are present on the pre-merge tree too. Not yet built or run -- that is next.
KeyError: 'confidence_head.weight' at draft-model load, so the serve never came
up. Collapsing our bare ReplicatedLinear into upstream's DSparkConfidenceHead
moved the parameter to confidence_head.proj.weight -- the wrapper holds its
linear as .proj -- while the checkpoint rename map still pointed at the old name.
Listed as identity rather than dropped: an unlisted `rest` falls through to
f"layers.{layer_idx}.{rest}", which would bury a top-level head under layers.2.
The other top-level entries (norm.weight, hc_head_*) are listed the same way.
The merge audit predicted this exact KeyError and I recorded it as 'flagging,
not fixing' instead of acting on it. The prediction was right and the cost was a
full model-load cycle.
These three tests are this fork's, covering our probabilistic draft-probs and spec-step-indices behaviour. They call SpecDecodeBaseProposer.propose(), which an upstream merge made take num_speculative_tokens as a required positional argument. The tests were never updated, so they have been failing with TypeError: propose() missing 1 required positional argument since that merge -- red, unnoticed, and covering code we care about. The fourth call site in the same file already passes it, which is what made the omission easy to miss. Read from proposer.num_speculative_tokens rather than restating each test's literal, so they stay correct if the fixture's value changes. Not from the 08-13 merge: identical failures at 67f5de5 and at the merge head, verified by running these exact tests at both SHAs.
test_dspark_sequential_sampling_writes_persistent_draft_logits called DSparkSpeculator.clear_runtime_draft_logits, which no longer exists, and the fixture missed enable_adaptive_verification because object.__new__ skips __init__. The missing method is not a lost fix -- I first read it that way and was wrong. fbbc8e7 added a reuse-and-clear scheme (assign base_logits into draft_logits, clear it afterwards); 7d9970d replaced that six days later with a persistent preallocated buffer, because DSpark drafting is CUDA-graph replayed and a Python-side reassignment does not run per replay. The clearing call became a `pass` and was later deleted. Restoring it would reintroduce something removed for cause. So the invariant is still worth asserting -- the draft-logits buffer must never be replaced -- but it has to be checked against something that still exists. The test now runs a second sampling pass and asserts buffer identity across it. Both this and the three test_mtp failures were red at 67f5de5 as well as at the merge head, so neither came from the 08-13 merge. What the merge did was move this one's failure to a later line, which is what made it visible. Remaining in tests/v1/spec_decode: test_dflash_drafter_window_reserves_bonus_token, whose fix is upstream vllm-project#51256 -- in the sixteen commits not yet merged.
Zero conflicts, which is the case that needed auditing rather than the case that did not. Six overlapping files were compared three-way (base 34735ac / ours 00c6d69 / upstream / merged); five were a clean union, verified by line-count ledger and AST rather than by eye. One real defect, plus one pre-existing one the audit surfaced. WHAT THE MERGE BRINGS THAT WE WANTED 025d56a DeepGEMM pin moves to deepseek-ai nv_dev tip (SM120 support) That is the ref our own support_deep_gemm() gate has been waiting on -- its comment says "re-enable once DeepGEMM vllm-project#324 lands", and vllm-project#324 is on nv_dev. We were pinned to vllm-project's fork with a TODO; upstream got there first. Whether the SM120 scale-factor assertion still fires is now testable rather than assumed. 2ac1f68 vllm-project#51256 reserves the DFlash bonus query slot -- expected to clear the last remaining tests/v1/spec_decode failure 79f3183 vllm-project#52058 bounds KV block zeroing launch geometry (we authored nothing in that file, so upstream's version applies directly) DEFECT INTRODUCED BY THE MERGE Upstream vllm-project#51879 added data_parallel_size and data_parallel_rank_local to OffloadingParallelConfig as REQUIRED fields with no defaults, and updated all five of its own construction sites. There is a sixth: our compact-offload test, in a file that does not exist upstream, so nothing conflicted and nothing prompted a look. _make_offloading_config() raised TypeError at setup, erroring three tests that certify compact-vs-legacy CPU manager dispatch. Both fields added. DEFECT THE AUDIT SURFACED, PRE-EXISTING test_kv_block_zeroer.py's two warmup fixtures built a 7-element _meta against KVBlockZeroer's 6-name unpack -- a duplicated tensor where production has seg_block_strides then seg_page_sizes. Both warmup tests raised ValueError before asserting anything, silencing precisely the warmup gate vllm-project#52058 strengthens. Not from this merge (byte-identical at 00c6d69); found because the merge made that file worth reading. All five fixtures now match the unpack, checked by AST. Not yet built or run.
test_dflash_drafter_window_reserves_bonus_token is upstream's, byte-identical to theirs, and it builds SimpleNamespace runner stubs carrying exactly the fields upstream's _input_fits_in_drafter reads. This fork's version reads two more: self.parallel_config, for the per-rank gate-off sentinel added with the TP drafter-gate work, and self._drafter_gate_off_logged, its log counter. So our production change broke their mock -- AttributeError from a stub that is correct for upstream and incomplete for us. Their test, our behaviour, our stub update. I expected upstream vllm-project#51256 (Reserve the bonus query slot in DFlash scheduling budget) to fix this, because the PR title matches the test name. It did not and could not: the failure was never upstream's. Verified after merging it -- same AttributeError at the same line. All five tests/v1/spec_decode failures identified earlier are now green: 3 x test_mtp (propose() signature drift), test_dspark_config (assertion on a method we deliberately removed), and this one.
Contributor
|
This pull request has merge conflicts that must be resolved before it can be |
4 tasks
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
20 participants
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.
Summary
This PR enables DeepSeek V4 Flash on SM120/SM121 Blackwell client hardware by carrying the SM12x fallback and tuning stack needed for the current vLLM V1 path. It targets RTX PRO 6000 Blackwell Workstation Edition, RTX 5090-class SM120, and GB10 / DGX Spark SM121 users who cannot use SM100-only TMEM /
tcgen05kernels.The branch is reconciled on top of the merged #43477 and provides the stock-deps path: DeepSeek V4 on SM120/121 that builds and serves on released FlashInfer / DeepGEMM wheels, complementing #43477's route that needs the unreleased FlashInfer #3395 + DeepGEMM #324 dependency branches. It is kept synced onto current
upstream/main.Latest validated head: tag
sm120-pr-41834-stable-preview-20260809(aa0d513027), synced ontoupstream/mainas of 2026-08-09 (f18e10a7e1) — see Update 2026-08-09 below. The default model runner is now V2;VLLM_USE_V2_MODEL_RUNNER=0still selects V1, which stays supported.Model / speculative-decode status.
deepseek-ai/DeepSeek-V4-Flash-0731is the checkpoint this branch is validated on. It removed the MTP heads and folded the DSpark draft into the main checkpoint, so DSpark (method: "dspark",num_speculative_tokens: 5) is the speculative path; MTP is supported only for older checkpoints that still carry those weights. Running without speculation is also fully supported and validated.Change footprint — model kernels vs. core-vLLM touch points
187 files, ~+29.1k / −1.3k against
upstream/main, of which ~10.7k added lines are tests. The branch splits cleanly into model/kernel code and a small set of core-vLLM integration points:vllm/models/deepseek_v4/**plus the SM12x sparse-MLA decode / indexer / DeepGEMM kernels that live in shared dirs (v1/attention/backends/mla/sparse_mla_kernels.py,model_executor/layers/sparse_attn_indexer.py,v1/attention/backends/mla/{indexer,sparse_swa}.py,utils/deep_gemm.py,kernels/mhc/tilelang.py), the new DSv4 reasoning parser / tokenizer, and device tuning JSONs.models/deepseek_v4/sparse_mla.py, perf) —_c128a_effective_topk_widthtakes the max position from the CPU-sideCommonAttentionMetadata.max_seq_leninstead of a per-stepint(positions.max().item())device sync, dropping a launch-stream stall on every C128A metadata step. Decode is identical (max_seq_len-1 == positions.max()); only chunked prefill sees a safe, slightly-wider 128-aligned top-k.single_type_kv_cache_manager.py,kv_cache_coordinator.py,kv_cache_manager.py,sched/scheduler.py(+1)cache_blockstail-block-reuse rewritev1/spec_decode/{dspark,dspark_sampling,llm_base_proposer,dflash}.py,config/speculative.pyfused_moe.py,oracle/mxfp4.py,routed_experts.py,experts/flashinfer_cutlass_moe.py,quantization/mxfp4.py,oracle/nvfp4.pyquantization/utils/fp8_utils.py,linear/scaled_mm/{cutlass,marlin}.py,csrc/.../marlin_moe_wna16/ops.cu(the only C++)config/vllm.py,compilation/breakable_cudagraph.py,passes/utility/fix_functionalization.py,config/compilation.pychat_completion/protocol.py,serve/render/serving.py,tool_parsers/structural_tag_registry.py,chat_utils.py,engine/protocol.py,chat_completion/{serving,batch_serving}.py,reasoning/__init__.pyreasoning_content/thinkingparam / tool-call streaming (jasl#19 instruction-following)model_executor/warmup/deepseek_v4_sm12x_warmup.py(new),kernel_warmup.py(+11)kernel_warmup.pystays a two-line hook on upstream's fileweight_utils.py,default_loader.pyenvs.py,utils/flashinfer.py,utils/import_utils.py,v1/worker/{gpu_model_runner,ubatch_utils}.pyVLLM_DEEPSEEK_V4_*flags +has_cutedsl/has_flashinfer_trtllm_sparse_mlaprobesTwo notes for review:
kv_cache_coordinatorcache_blocksrewrite (affects hybrid-KV models; validated ≥ prior behavior), the proposer base-class change, and the OpenAI-entrypoint plumbing. Everything else (MoE oracle, fp8_utils, cudagraph gate, warmup, envs) is arch / quant / env-gated and inert for other models.Duplicate-work check
The nearest open/merged PRs are related but not duplicates:
42657aca65) and carries the stock-deps DSv4 SM120/121 path that runs on released wheels._prefill_workspace_topk_boundreturns early forcompress_ratio <= 1and never reaches the affected buffer.Upstream PRs whose fixes this branch previously carried as local deltas and has since retired in favour of upstream's own version: #48304, #48911, #48959 (via #49052).
Fixed preview tags
These tags are in
jasl/vllmand give users stable pins while the PR is still moving:sm120-pr-41834-stable-preview-20260809aa0d513027sm120-pr-41834-stable-preview-202608040f59188db1sm120-pr-41834-stable-preview-202608029a94c54292DeepSeek-V4-Flash-0731support, two DSpark config fixes, #49335 / #50686 absorbed. See Update 2026-08-02.sm120-pr-41834-stable-preview-20260727dd64074e6f0…-20260727,70a33886bd); DSpark VRAM work (jasl#27) merged; bounded block-table gather incompute_global_topk_indices_and_lens.sm120-pr-41834-stable-preview-20260721832775efd1sm120-pr-41834-stable-preview-20260717f63bfd3d7bsm120-pr-41834-stable-preview-20260711b5c0d43b96sm120-pr-41834-stable-preview-20260704b43470e871constexpr→runtime (stops the Triton recompile → unified-memory leak → hard-freeze) + fp8-einsumtl.multiple_of(16)(~24% decode @256k).sm120-pr-41834-stable-preview-20260703444fe3ac8bpersistent_topkfor <128 KB-smem parts.Older tags (
…-20260705back to…-20260612…) remain injasl/vllmfor history.Update 2026-08-02 —
DeepSeek-V4-Flash-0731, 234 upstream commits, two DSpark fixesValidated head
9a94c54292(tagsm120-pr-41834-stable-preview-20260802), 234 upstream commits absorbed, level withupstream/mainas of 2026-08-02.What's in it
DeepSeek-V4-Flash-0731support. The new checkpoint ships no MTP heads —enorm,hnorm,e_proj,h_projandshared_headare absent from the weight index, andmtp.{0,1,2}.*now carries the DSpark-stylemain_norm/main_projstructure (matchingdspark_target_layer_ids: [40, 41, 42]). DSpark is the speculative path going forward; the MTP code is retained for older checkpoints.num_speculative_tokensvsdspark_block_size— this rule was relaxed on2026-08-04; see Update 2026-08-04. It now errors only BELOW the block size and warns
above it. The original reasoning and measurements follow.
>=and its error message recommended exceeding it. The drafter emits exactly one block per pass, so the extra slots are structurally unreachable — measured on a prose workload, the 7th draft position accepted 0.000 in every sample (the 6th in all but one, 0.004 there), andnst=7drafts 40% more tokens per step for strictly worse acceptance:All samples are shown rather than a single figure: the probe reads whatever
SpecDecoding metricslines vLLM flushed inside its window, so a low sample means "not much steady traffic in that slice", not a worse drafter. Bothnst=7runs also hit connection errors partway through, so their spread is noisier.method: "mtp"is no longer silently rewritten to"dspark". Auto-detection preserved an explicitly requested method only foreagle/eagle3/dflash/dspark. Since 0731 putsdspark_block_sizein every DSv4 config,method: "mtp"fell through to the dspark branch, was rewritten, and then failed validation with a DSpark message the user never asked for.<|end_of_sentence|>on the defaultdrop_thinking=Truepath). Fix DSpark warmup without sparse index buffer #50693's regression test is carried; its code fix is not reachable here.Check failed: num_tokens > 64, andFLASHMLA_SPARSE_DSV4missingtile_sched. Details in this comment.Validation (GB10 SM121, 2-node TP=2,
DeepSeek-V4-Flash-0731, torch 2.13.0, FlashInfer 0.6.15.post1, nccl 2.30.7)The GSM8K difference (1.06 pp flexible / 1.21 pp strict) is within this gate's measured single-run spread (~1.1 pp). Resolved: three runs per cell were collected and the arms interleave, so it was noise.
0731is the first checkpoint where the strict and flexible extractors disagree at all; on every prior baseline they were identical.Perf — pinned llama-benchy standard (fp8 KV, prefix-cache on,
FULL_AND_PIECEWISE, mml 49152, util 0.85; C=1, 3 runs), against the full recorded range of the prior MTP2 baselines. This crosses a checkpoint boundary, so read it as a sanity band rather than a controlled A/B:Batched prefill (pp2048) is above the historical band at all three depths (+2.2% / +0.9% /
+2.0%) — the only consistent directional move here. Clearing the max of ten prior runs at all
three depths says more than any single one of those margins would: +0.9% is inside this metric's
own resolution, so read the consistency rather than the magnitudes. No sign of DSpark being
slower than MTP2 was.
One caveat reported rather than buried:
ctx_tg @ d16384sits 10.7% below its historicalminimum, the only metric outside its band. It is non-monotonic against our own neighbouring
depths (39.37 at d8192, 40.70 at d32768, where history has d16384 ≈ d8192), which points at a
single-run artifact rather than a depth-specific regression. Resolved: repeated on later
heads and it did not recur.
A measurement caveat for anyone benchmarking this branch: the
±in a benchy row is the spread of the three runs inside one invocation, and it runs 5–30× smaller than the build-to-build spread. This branch's own history spans 31% on tg128 @ d32768 and ~1.3% on ctx_pp, so anything under ~15% on tg or ~2% on ctx_pp is not resolvable this way.Update 2026-08-04 — four fixes from community reports, 35 upstream commits, and first SM120 validation
Validated head
0f59188db1(tagsm120-pr-41834-stable-preview-20260804).This is the first head validated on both SM121 and SM120. Every SM120 discrete-GPU
result on this PR up to now was a contributor's measurement we could not reproduce. We have
since rebuilt a 2× RTX PRO 6000 Blackwell box as a first-party SM120 target.
Fixes
DSpark's fused Markov sampler could emit an out-of-vocab token id (
e171c51036)._dspark_markov_probs_blocks_kernelstoresvocab_sizeas the filler for a block with noactive lane. On a fully-masked row — every candidate
-inf, which structured-outputconstraints can produce — no block has an active lane, so every block stores the filler and
the reduce kernel returns it verbatim as the sampled token. Nothing downstream bounded it: the
runner clamped
input_idswithmin=0only, and the DSv4 hash-MoE router indexestid2eid[token_id * 6 + lane]on a[vocab_size, 6]table. Result is an illegal memory accesson every TP rank.
This is the producer on the V1 path, which is this branch's default. @alexbi29's report
traced the same class of defect to the V2 samplers ([Bugfix] Bound tile-local argmax to vocab_size in samplers #50843) — a real defect, but a different
tree. Fixed by folding out-of-range to
0(matchingtorch.argmaxon such a row, so the fusedkernel stays bit-identical to the eager reference) and making the runner clamp two-sided.
Worth stating plainly for anyone with similar gates: our own gates could not have caught
this. The fused path is skipped when
all_greedy, and both our long-context recall gate andGSM8K are greedy, so they are structurally incapable of executing that kernel. The new
regression test is explicitly non-greedy.
Adopted [Bugfix][DSv4] Bound token_id before the tid2eid gather in hash-MoE routing #50844 (
3df857ba50) — boundtoken_idbefore thetid2eidgather. Defence indepth;
prompt_token_idsreach that gather directly when--skip-tokenizer-initdisables theengine's vocab check. Not taking [Bugfix] Bound tile-local argmax to vocab_size in samplers #50843 (V2-tree only, inert on our default) or [Bugfix][MoE] Bound expert_map gathers on data-derived expert ids #50845,
which has a defect reported on its own thread.
Eager scratch pool is now OFF by default (
d42b8d9f55,b1ef3033f4), opt-in viaVLLM_DEEPSEEK_V4_EAGER_SCRATCH_POOL=1. @tobymao bisected output corruption under concurrentmixed prefill+decode to it: pool active 7/7 rounds corrupt, disabled 0/2. We first removed the
cross-template aliasing (
max()→sum()sizing with per-family offsets); they tested thatcommit directly and it was still corrupt in round 1. Their diagnosis is the useful part: the
pre-pool code was race-free for free because per-call transients go through the caching
allocator, whose cross-stream reuse is event-guarded — the pool reuses memory without that
machinery, so no static partitioning fixes it. Making the cross-layer reuse safe needs
producer-waits-on-consumer events against the real stream graph; until then, off by default.
Two contributor PRs merged — sm12x: add tuned FP8 W8A8 block config for N=4096,K=12288 jasl/vllm#37 (tuned FP8 W8A8 config for
N=4096,K=12288onRTX PRO 6000) and sm12x: hoist the E8M0 block-scale upcast out of the FP8 GEMM hot path jasl/vllm#38 (hoist the E8M0 block-scale upcast out of the FP8 GEMM hot path,
13,561 kernel launches removed per 25 decode steps), both from @alexbi29.
num_speculative_tokensrule relaxed. Upstream removed its own assertion in [Bugfix] Remove bad startup assertion #50869 as"invalid". They were right that erroring above
dspark_block_sizeis wrong — two users onthis thread run
nst=7againstblock_size=5and it demonstrably works. The two directions arenot symmetric, so this branch now errors below the block size (that genuinely garbles
output) and warns above it, quoting the acceptance cost. Strictly more permissive than what
shipped before.
Validation
Full gate battery on both architectures, same branch:
--block-size 256The ~1.1 pp GSM8K difference between architectures sits inside this gate's measured single-run
spread and spans different silicon, different memory architecture and a 3× smaller KV cache
(6.25 GiB vs ~18.5 GiB). We are not claiming a difference from it.
Check failed: num_tokens > 64does not reproduce on this branch. @fuzzifikation reportedstock 0.26.0 dying there on SM120 at
--block-size 256, correctly tracing it to the DSv4 decodedispatch requiring
page_block_size == 64. On our SM120 box, at the same--block-size 256, theserve comes up and the assertion never appears — the DSv4 packed KV cache is laid out in 64-token
pages independent of vLLM's logical block size, and FlashInfer derives
page_block_sizefromtensor geometry rather than the engine config. The SM120 packed decode path is confirmed engaged
in the same run. Note the same assertion has two distinct gates (
page_block_sizeand(num_heads, topk), the latter being #50720 / flashinfer#3989), so patching one and still seeingit means checking the other.
Prefill: V1 vs V2 model runner
The 2026-08-02 V1-vs-V2 comparison never measured throughput. It has now been measured, blocked
and pre-registered — 10 blocks, both arms inside each node pair, exact sign-flip permutation test,
Holm-corrected across the six prefill cells, with the decision rule committed before any data was
collected:
All six survive Holm; both node pairs agree in direction on every cell. Decode is not resolved
in either direction —
tg128was declared unresolvable before the run (its within-build spreadequals its entire historical range) and is reported for the record only.
V1 was the default when this was written; that was reversed on 2026-08-09 — see Update 2026-08-09. The reasoning below was correct on the evidence available at the time, and the collapse it describes was real; it turned out not to be a property of the runner. Kept unedited because how the conclusion failed is the useful part. V2 is ahead on prefill, KV headroom (+4.70 GiB) and draft acceptance
(+6.6%), but its long-context recall under concurrency is unreliable in a way that is worse than a
consistent deficit: across 14 independent serves on the same build and configuration, roughly two
thirds land in a state that loses most of the needles (arthur c=12 as low as 3/24), while the rest
match V1 at 22–24/24. The mode is fixed at startup and stable within a serve, and nothing we have
found predicts or detects it. A deployment could run clean for days and restart into the bad mode.
The cause is not identified. Eliminated so far: the eager-scratch cross-template aliasing, the
upstream merges, and the eager scratch pool as a whole (pool on 2 good / 6 bad vs pool off 3 good /
3 bad over 14 serves — no effect). The startup logs of a good and a bad serve are structurally
identical, which rules out "a different code path was taken". Anyone opting into V2 with
VLLM_USE_V2_MODEL_RUNNER=1should know this.Measurement note
Two errors from our own process, since they affect how the numbers above should be read.
The n=8 sampling that originally established V2's recall deficit took eight gate runs from one
serve — it measured within-serve variance while the quantity that actually varies is
across-serve. Raising n on the wrong axis. The 14-serve figures above use the inverted design:
many serves, few gates each.
And the
±in a benchy row is the spread within one invocation; it runs 5–30× smaller than thebuild-to-build spread. The blocked design above exists because of that: a coarse range screen over
the same 10 blocks returns "no measurable difference" on all six prefill cells, while the paired
test finds all six. Had the screen been the decisive statistic, this section would have concluded
the opposite and been wrong.
Update 2026-08-09 — the V2 recall collapse was a prefix-cache race, not the runner; V2 becomes the default
70 upstream commits (to
643c125fab), and the long-standing reason this branchpinned V1 is gone: it was an unfixed upstream bug, not a property of the V2
model runner.
The defect
FullAttentionManager.cache_blocks()commits prefix block hashes to the sharedBlockPoolat scheduling time, before the forward pass writes their KV. Arequest admitted later in the same step can match those hashes and read unwritten
values.
MambaManagerhas guarded this since #29387; no other manager does.This is #42359, open and unmerged. Two more reports look
like the same triple on different models — #50188 (prefix caching + MTP spec
decode + fp8 KV, byte-identical repeat requests, RTX 5090 / Qwen3.6-27B-NVFP4)
and #43559 (closed without a merged fix). Anyone on
--enable-prefix-cachingwith speculative decoding is exposed; DeepSeek-V4 is not special here.
What makes it hard to catch: the damage persists. A serve that loses the race
keeps serving from the poisoned blocks for its lifetime, so a later serial
request fails too — which is why it looked like a per-serve "mode" rather than a
race. It is also stochastic, roughly half of cold serves.
Evidence
Same binary,
VLLM_ALLOW_SPEC_DEC_SAME_STEP_PREFIX_HITthe only variable, cachepopulated by the real gate, 4 fresh serves per arm (a single clean serve
proves nothing at ~50% incidence), 3 arthur c=12 runs each:
Mann-Whitney U, p = 0.0043. Every serve's runner and guard state was read back
from the serve log rather than assumed.
It also fixes V1, which was not expected: V1's arthur c=12 goes 20.7 → 23.0
with the guard on (24.0 with prefix caching disabled entirely). V1's own 2–4
needle shortfall was the same defect, not an inherent concurrency margin.
Runner arbitration, re-run on the fixed tree
Same tree, same guard mode, runner the only variable:
Recall: p = 0.697, neither side with a single-digit serve. GSM8K differs by
<0.4 pp against ~1.1 pp single-run noise. V2 is not behind anywhere and leads
on throughput, latency and KV headroom, so it becomes the default.
VLLM_USE_V2_MODEL_RUNNER=0still selects V1, which stays supported.Correction: the "V2 +6.6% draft acceptance" figure in Update 2026-08-04
does not survive re-measurement — 2.772 (V1) vs 2.710 (V2) on the same
formula and sample size, i.e. a tie. It was measured while V2 was poisoned.
If you are running this branch
Nothing to set — the guard is on by default where it matters.
KVCacheCoordinatorenables it whenever prefix caching and speculative decodingare both active, which is the only configuration in which a block hash can be
published before its KV is written and a second request admitted in the same
step to match it.
This correction matters: an earlier revision of this update shipped V2 as the
default while leaving the guard off by default, which would have handed a plain
serve the exact combination measured at mean 11.5 with a 3/24 floor. Both
changes looked like improvements in isolation. If you pulled
c054feedac,take
aa0d513027instead, or set the variable yourself.To turn it off (it is a real escape hatch, pinned by a test):
1is upstream's semantics, gated onuse_eagle; on DeepSeek-V4 that coversonly 2 of 5 managers and leaves the main MLA path unguarded — measured, not
assumed, via a startup log line this branch adds that reports how many managers
are actually guarded.
2covers every group and is what the engine selects.Regression sweep on the merged tree: 261 passed, 1 failed, that one failing
identically on the pre-merge tree
4ebd1fb698.The two endpoints disagreed about the same model, twice
ResponsesRequest.reasoningtook its type from the OpenAI SDK, whoseReasoningEffortstops atxhigh, so DeepSeek's documented top tiermaxwasrejected by schema validation on
/v1/responseswhile/v1/chat/completionsaccepted it.
Worse, and on the default path: with no thinking kwarg,
DeepSeekV4Tokenizer.apply_chat_templatedefaults thinking on whileDeepSeekV4ReasoningParserdefaults it off and selectsIdentityReasoningParser. The model reasoned and its reasoning, with a bare</think>, came back insideoutput_textas though it were the answer —whenever a request omitted
reasoning, which is exactly what a stock OpenAISDK sends. Chat was immune only because it normalises thinking state at the
protocol boundary, and its own docstring says why: "so the tokenizer and
reasoning parser see the same effective state". Responses never called that
hook. The derivation now lives in
deepseek_v4_chat_kwargsand both requesttypes call it, so a third endpoint cannot repeat it.
Measured on both checkpoints with no workaround flag set, 21/21 each:
DeepSeek-V4-Flash-0731DeepSeek-V4-Flash</think>in the answereffort: nonedisables thinkinghighreasons deeper thanlow26 unit cases accompany it, 11 of which fail on the unpatched tree.
tests/reasoning440 passed,tests/tokenizers_/test_deepseek_v4.py45 passed.Acceptance on the exact published SHA
Everything above was re-measured on
d44e224ab9— the commit this tag points at,after a second upstream sync (17 further commits, FlashInfer 0.6.16.post3) — not
on an ancestor assumed to be equivalent. 17 of 18 checks pass:
tests/v1/coreVLLM_USE_V2_MODEL_RUNNER=0→ V1, guard still 5/5, c=1 2/2VLLM_ALLOW_SPEC_DEC_SAME_STEP_PREFIX_HIT=0→ guard 0/5, c=1 2/2The one non-pass, and what it turned out to be.
tests/v1/spec_decodedoesnot complete on this hardware — it wedges under a 30-minute bound on this head
and on
4ebd1fb698alike. Narrowed totest_max_len.pyand measured both ways:No individual case is broken. Each stands up a full engine, and repeated
create/tear-down inside one process does not release resources fast enough on a
single-GPU unified-memory node. That also explains why both trees wedge and
why they stop at different points. It remains unverified coverage rather than a
pass; running one process per case produces a verdict instead of a hang.
test_async_scheduling_pp_allows_rescheduling_with_output_placeholdersis thesame class: it builds
pipeline_parallel_size=2, and a GB10 node has one GPU, soit fails at config construction. It is the only case in
tests/v1/corethat needsmore than one GPU; the other 509 pass.
What this arbitration does and does not cover
Everything above was measured on one configuration: 2-node TP=2,
DeepSeek-V4-Flash-0731, DSparknum_speculative_tokens: 5, fp8 KV,max_model_len131072, prefix caching on, GB10 (SM121). The default now appliesto every DSpark config, including shapes not measured here — TP=4, other
context lengths, the NVFP4 checkpoint, single-node setups.
The reasoning generalises better than the numbers do: the race is in block
publication and is not specific to a model shape, and V2's advantage comes from
KV headroom and scheduling rather than anything config-specific. But if you run
a materially different shape and see something worse,
VLLM_USE_V2_MODEL_RUNNER=0returns you to V1 and a report would be welcome — that is a gap in our coverage,
not a claim we have ruled out.
Running DSpark
DSpark is DeepSeek's self-drafting speculative-decode variant; on
0731the draft weights are carried in the main checkpoint, so no separate--speculative-modelis needed.vllm serve deepseek-ai/DeepSeek-V4-Flash-0731 \ --trust-remote-code \ --tokenizer-mode deepseek_v4 \ --tool-call-parser deepseek_v4 --enable-auto-tool-choice \ --reasoning-parser deepseek_v4 \ --tensor-parallel-size 2 \ --kv-cache-dtype fp8 \ --block-size 256 \ --max-model-len 49152 \ --max-num-seqs 64 \ --max-num-batched-tokens 8192 \ --gpu-memory-utilization 0.85 \ --enable-prefix-caching \ --speculative-config '{"method":"dspark","num_speculative_tokens":5,"draft_sample_method":"probabilistic"}'num_speculative_tokensmust equal the checkpoint'sdspark_block_size(5). Larger values are rejected: they are never accepted and only waste draft compute.--kv-cache-dtype fp8is mandatory — DSv4'sfp8_ds_mlaattention asserts an fp8 KV layout, so the defaultautofails at model construction. Not DSpark-specific.VLLM_USE_V2_MODEL_RUNNER=1opts into the V2 DSpark speculator; V2's long-context recall is correct after the fix(dspark): reserve V2 padded Q scratch jasl/vllm#26 padded-Q fix.Dependencies (stock-deps path)
Pins on the current head: torch 2.13.0 (triton 3.7.1) ·
flashinfer-python/flashinfer-cubin0.6.15.post1 · tilelang 0.1.12 ·nvidia-cutlass-dsl[cu13]4.6.0 ·quack-kernels>=0.6.1· nvidia-nccl-cu13 2.30.7 (multi-node, see below).requirements/cuda.txt(flashinfer-pythonand the GitHub-releaseflashinfer-cubin, which must be the same version); it ships the SM120 packed sparse-MLA kernels, so a stock build picks them up with no manual install dance.nvidia-nccl-cu13==2.30.7on every node. A rebuild silently reverts it to torch's bundled version, and a per-node mismatch hangs the NCCL handshake.VLLM_DEEPSEEK_V4_FLASHINFER_SM120_DECODE) and prefill (VLLM_DEEPSEEK_V4_FLASHINFER_SM120_PREFILL) FlashInfer sparse-MLA paths default on; set either=0to fall back to the FlashMLA / Triton path. Both are availability-gated, so stock installs without the kernel degrade gracefully rather than raising.Running the NVFP4 checkpoint
This branch also serves
nvidia/DeepSeek-V4-Flash-NVFP4on SM12x (RTX PRO 6000 / GB10). The NVFP4 MoE auto-selects the FlashInfer CUTLASS backend (the SwiGLU-clamp model gate accepts it), so no--moe-backendflag and no special FlashInfer build are required:Expert-parallel off (plain TP) is the supported path. Accuracy matches MXFP4 (GSM8K 8-shot ~0.96 on both SM120 and SM121). On SM12x NVFP4 is not a memory or throughput win versus MXFP4: NVFP4 weights are ~4 GiB/GPU larger, leaving less KV-cache room; single-stream prefill is marginally faster and aggregate decode marginally slower. Its value here is checkpoint availability / parity with the SM100 datacenter path — MXFP4 remains the better practical choice on consumer Blackwell.
AI assistance disclosure
AI assistants, including OpenAI Codex/GPT models and Anthropic Claude models, were used for code review, refactoring support, regression-script writing, and benchmark analysis. The branch was validated through human review plus the commands and harness artifacts listed above; every performance and accuracy number quoted was measured on real SM120/SM121 hardware.