[Bugfix][SM120] DSv4: pass contiguous C128A decode topk indices on SM120 by lucifer1004 · Pull Request #53574 · vllm-project/vllm

lucifer1004 · 2026-08-24T12:38:16Z

[Bugfix][SM120] DSv4: keep C128A decode topk indices contiguous on SM120

Purpose

DeepSeek-V4 Flash on SM120 (e.g. RTX 6000 Pro, DGX Spark) with the FlashInfer
sparse-MLA backend crashes during CUDA-graph capture at server startup:

Check failed: (eidx.IsContiguous()) is false: eidx must be contiguous

Root cause: since #52823 (adaptive topk width), build_c128a_topk_metadata
returns a width-narrowed slice of the persistent c128a_global_decode_buffer,
which keeps the buffer's row stride and is therefore not contiguous. The C128A
decode path in DeepseekV4FlashInferSM120Attention._forward_decode passes it
to FlashInfer as extra_sparse_indices (eidx). FlashInfer dispatches calls
with num_tokens <= 64 to the standalone decode kernels (no eidx contiguity
check) and larger calls to the paged orchestrator, whose
CHECK_INPUT_AND_TYPE(eidx, ...) rejects the tensor. Plain target-only decode
stays under the 64-token cutoff and never trips this; with speculative
decoding enabled, verification batches carry num_decodes * (1 + K) rows and
cross the cutoff, so the failure surfaces during CUDA-graph capture at startup
(confirmed independently on DGX Spark / sm_121a by @maci0).

Fix: build_c128a_topk_metadata gains a full_width_decode flag, set by the
metadata builder on SM120 only. With it, the decode view is the full-width row
slice of the persistent buffer: contiguous, at a graph-stable address (a
per-step .contiguous() copy would allocate a fresh tensor every step and
break capture), with the kernel's strided writes and the adaptive
max_compressed_tokens bound unchanged. Decode-side consumers bound reads by
the per-token topk lens, so stale columns past the active width are never
read. SM100 (FlashMLA/TRTLLM) keeps the existing active-width slice, and the
prefill view stays narrowed for all backends since its Triton consumers count
non-negative entries across the row.

Longer term, the cleaner fix is to let the FlashInfer SM120 kernel accept a
real row stride for eidx; this PR is the minimal short-term unblock.

Test plan

pytest tests/kernels/attention/test_flashmla_sparse.py —
test_deepseek_v4_c128a_adaptive_width_has_capture_stable_stride now runs
both modes: SM120 mode asserts full-width/contiguous/graph-stable decode
view, default mode keeps the adaptive-width slice. All pass.
prek run --files ... — all hooks pass (ruff, mypy-3.10, etc.).
E2E on 4x RTX 6000 Pro (SM120), DeepSeek-V4-Flash-0731 NVFP4, TP4 with
DSpark speculative decoding, aggregated serving: without this fix the server
dies during CUDA-graph capture with the eidx error above; with this fix the
server reaches readiness and an 8K/1K random benchmark completes with zero
failed requests.

Notes

AI assistance (Kimi Code / OpenAI Codex) was used in preparing this change;
every line was reviewed and tested by the submitting human.

claude

Claude Code Review

This pull request is from a fork — automated review is disabled. A repository maintainer can comment @claude review to run a one-time review.

maci0 · 2026-08-24T17:03:20Z

Confirming this root cause on a second SM120 deployment: 2x DGX Spark (GB10, sm_121a), DeepSeek-V4-Flash-0731 + DSpark k=5 speculative decoding, vLLM main e25c586b9 (2026-08-23), FlashInfer sparse-MLA backend.

The 0731 checkpoint alternates compress_ratios [..., 4, 128, 4, 128, ...] across layers, so both the C4A and C128A decode branches execute on this model. We hit the same Check failed: (eidx.IsContiguous()) is false: eidx must be contiguous crash during CUDA-graph capture at boot, and it reproduces exactly the mechanism described here: with DSpark k=5 the verification batch carries num_decodes * 6 rows and crosses the 64-token cutoff into the paged orchestrator path that checks eidx contiguity.

We shipped a consumer-side workaround (.contiguous() on extra_sparse_indices in both branches of DeepseekV4FlashInferSM120Attention._forward_decode) which boots and keeps greedy output coherent. We also verified the C4A branch is not affected: compute_global_topk_indices_and_lens does torch.empty_like() of a contiguous row slice of the topk_indices_buffer allocated in dspark.py, so its view(num_decode_tokens, 1, -1) is contiguous. This builder-side fix is the right place — thanks for the PR.

yewentao256

Thanks for the work!

I think the correct fix in long term is to let FlashInfer kernel accepts real stride(0), but let's accept the short term fix here. Small updates needed before landing

yewentao256 · 2026-08-25T15:22:04Z

@@ -322,10 +329,9 @@ def build_c128a_topk_metadata(
    )
    assert global_decode_buffer.stride(-1) == prefill_buffer.stride(-1) == 1

-    global_decode = global_decode_buffer[:num_decode_tokens, :max_compressed_tokens]
+    global_decode = global_decode_buffer[:num_decode_tokens]


Take a look, the slicing works for FlashMLA on SM100, but might not works for FlashInfer sparse-MLA on SM120, let's add a if statement here for different backend. Let's keep SM100 unchanged

yewentao256 · 2026-08-25T15:22:31Z

@@ -268,6 +268,7 @@ def _build_c128a_metadata(

        result: dict[str, torch.Tensor | None] = {}
        if num_decode_tokens > 0:
+            # Full-width view: contiguous (SM120 eidx check) and graph-safe.


Suggested change

# Full-width view: contiguous (SM120 eidx check) and graph-safe.

build_c128a_topk_metadata returned a width-narrowed slice of the persistent c128a_global_decode_buffer, so the view kept the buffer's row stride and was not contiguous. The FlashInfer SM120 sparse-MLA kernel rejects it as extra_indices (eidx must be contiguous), killing the server during CUDA-graph capture whenever a decode-side call exceeds 64 tokens (e.g. DSpark verification batches with num_decodes * (1 + K) rows) on a C128A layer. Return the full-width row slice of the buffer instead: it is contiguous and stays in the persistent buffer, so CUDA-graph capture sees a stable address. Decode-side consumers bound reads by the per-token topk lens, so stale columns past the active width are never read. The prefill view keeps the active-width slice because its Triton consumers count non-negative entries across the row and are stride-aware. Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>

lucifer1004 · 2026-08-26T09:03:25Z

Done — the decode view is now full-width only on SM120 (full_width_decode flag gated on capability.major == 12); SM100 keeps the adaptive-width slice unchanged. Also dropped the extra comment. Re-verified on RTX 6000 Pro: capture crash gone, benchmark passes.

Agreed the cleaner long-term fix is letting the FlashInfer SM120 kernel take a real row stride for eidx — noted in the PR description.

…nto inferlab-integration

yewentao256

Thanks for the work!

yewentao256 · 2026-08-26T19:29:21Z

+    if full_width_decode:
+        global_decode = global_decode_buffer[:num_decode_tokens]
+    else:
+        global_decode = global_decode_buffer[:num_decode_tokens, :max_compressed_tokens]


Could we remove the additional argument full_width_decode and put the structure jugement here?

Also, add a TODO saying we should support SM120 adaptive width

Happy to inline the capability check here. One testing concern before I do: with the platform judgment inside build_c128a_topk_metadata, the two branches are no longer directly selectable from unit tests, so the existing stride/contiguity test would need to monkeypatch current_platform.get_device_capability() to cover both the SM120 full-width and the adaptive-width paths. Is monkeypatching the platform singleton acceptable in this test suite, or would you prefer keeping a narrow seam (e.g. a keyword arg with a platform-derived default) so both paths stay directly testable?

Let's keep the production code simple, for tests we can use some trick.

Done — the capability check is inlined into build_c128a_topk_metadata with a TODO for SM120 adaptive width, and the test now monkeypatches get_device_capability to cover both branches (bdff48b).

yewentao256 · 2026-08-26T19:29:51Z

+    With full_width_decode (SM120), the decode view keeps the full buffer
+    width: full-width row slices are contiguous, which the FlashInfer SM120
+    kernel requires of eidx, and decode reads are bounded by the per-token
+    lens. Otherwise the decode view is narrowed to the active width. The
+    prefill view always stays narrowed: its Triton consumers count
+    non-negative entries and must not scan stale tail columns.


Suggested change

With full_width_decode (SM120), the decode view keeps the full buffer

width: full-width row slices are contiguous, which the FlashInfer SM120

kernel requires of eidx, and decode reads are bounded by the per-token

lens. Otherwise the decode view is narrowed to the active width. The

prefill view always stays narrowed: its Triton consumers count

non-negative entries and must not scan stale tail columns.

Drop the full_width_decode argument: build_c128a_topk_metadata now checks the device capability itself, with a TODO for SM120 adaptive-width support. The stride/contiguity test monkeypatches the platform capability to cover both branches. Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com>

…nt) into inferlab-integration

Signed-off-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com>

yewentao256

LGTM, thanks for the iterations!

yewentao256 · 2026-08-27T15:04:43Z

/ci run

github-actions · 2026-08-27T15:04:52Z

✅ @lucifer1004, CI is now available for this PR.

/ci run starts upstream CI; /amd-ci run starts AMD CI only.
/ci retry retries failed jobs in the CI build for the current PR head. If the current head has no CI build, it starts a new CI build for the current head containing only jobs that failed in the latest earlier CI build for this PR.
/amd-ci retry retries failed jobs in AMD CI for the current PR head. Use /amd-ci run when the current head has no AMD CI build.
/ci cancel cancels scheduled or running CI builds for this PR branch; /amd-ci cancel does the same for AMD CI only.

github-actions · 2026-08-27T15:05:02Z

✅ Triggered Buildkite CI #85828 for commit 097681d294d7.

lucifer1004 · 2026-08-28T01:49:37Z

The failing step Multimodal Models (Standard) 4: other + whisper (H200) looks like an infrastructure flake rather than a real failure: the job hit its 75-minute timeout while hanging in test_moss_audio_generation_smoke (MOSS-Audio engine init went silent for ~71 min), then Buildkite cancelled it (exit status -1). No test actually failed. This PR only touches deepseek_v4/sparse_mla.py (SM120 decode path), which these multimodal tests on SM90 don't exercise. Could a maintainer retry that step? Thanks!

zyongye · 2026-08-31T08:04:20Z

/ci run

github-actions · 2026-08-31T08:04:38Z

✅ Triggered Buildkite CI #86319 for commit f4e853008ab2.

…120 (vllm-project#53574) Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com> Signed-off-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com> Co-authored-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>

## Summary Consolidated SM120 sparse-MLA rework, decode + prefill. All numbers on RTX PRO 6000 (SM120). - **Faster decode, small T**: T=1 14.9 → **11.3µs** (−24%, graph replay, dual-cache 18K context); bitwise-identical outputs. - **Calibrated dispatch**: an analytical `chunks_per_block` model + measured decode/prefill crossover replace the per-shape autotune sweep and the hard `T ≤ 64` cutoff (up to −61% on rerouted configs; tables below). - **Continuous envelopes**: decode serves any `num_heads ∈ [1,128]` and any `topk ≥ min_topk`; prefill serves any `T ≥ 1` and any `topk % 64 == 0` width. Runtime-topk prefill drops instantiations **75 → 55**. - **swapAB prefill** carried from #4751 behind a per-call `prefill_impl` override; independently re-benched at **1.12–2.37×** over MG. - **Two new model types**: `GLM53_NOPE` (carried from #4791) and `DOTS3_SWA` (sliding-window MLA, d_qk=1088, d_v=1024, 1160 B/token footer-scale, padded-row KV support) — the latter also fixing five latent bugs along the way (rope writeback overrun at D_V==D_NOPE, flat-vs-paged addressing keyed on the wrong trait, a Python chunk-width hardcode, an undersized amax scratch at 4 math warps, a vestigial SG register array). - **Public runner**: `flashinfer.mla.SparseMLASm120Wrapper` — one persistent instance, memoized dispatch, CUDA-graph-safe (decode scratch is routing-aware and instance-owned). - Merged current main, incl. the #4732 SM121 prefill-hang fix. - Also: row-strided `indices` (unblocks vllm-project/vllm#53574's persistent-buffer narrowing) and row-strided `out_lse`; T=0 decode returns empty instead of aborting; decode bindings now validate `out_lse`/index dtypes/dim0. Carries (authorship preserved): #4461 zero-token decode (rewritten; XingSong), #4551 dispatch diagnostics + `supported_sparse_mla_sm120_configs()` (Sam Mausberg), #4751 swapAB (Lemon7-UP), #4791 GLM53_NOPE (lucamotz; extended with H=64/TP1 decode, swapAB@2176, calibration coverage). Supersedes #4683: its per-shape sweep profiles L2-resident synthetic indices, which distorts cpb when production caches are DRAM-resident (observed on 5070 Ti) — this PR removes the sweep instead (thanks Sam for the original analysis). ## Performance vs main (adc49a8) Same GPU, fixed-seed identical inputs, CUDA-graph replay GPU-only, both sides out-of-box (no tactic cache / no calibrated constants). Only surfaces present on both sides listed. | shape | main | PR | speedup | |---|---|---|---| | dsv4-dual-h64 (topk 128+512), T=1 | 14.80µs | 11.40µs | 1.30x | | dsv4-dual-h64 (topk 128+512), T=8 | 19.80µs | 16.14µs | 1.23x | | dsv4-dual-h64 (topk 128+512), T=16 | 36.66µs | 29.46µs | 1.24x | | dsv4-dual-h64 (topk 128+512), T=64 | 97.19µs | 89.79µs | 1.08x | | dsv4-h128 (topk 1024), T=1 | 14.70µs | 11.44µs | 1.29x | | dsv4-h128 (topk 1024), T=64 | 231.60µs | 219.79µs | 1.05x | | dsv3_2-h64 (topk 2048), T=1 | 14.08µs | 10.68µs | 1.32x | | dsv3_2-h64 (topk 2048), T=64 | 216.78µs | 220.23µs | 0.98x | | dsv3_2-h128 (topk 2048), T=1 | 16.57µs | 14.08µs | 1.18x | | dsv3_2-h128 (topk 2048), T=64 | 324.19µs | 323.53µs | 1.00x | | dsv4-prefill-h128 (topk 1024), T=128 | 293.82µs | 266.49µs | 1.10x | | dsv4-prefill-h128 (topk 1024), T=2048 | 4486.89µs | 3928.68µs | 1.14x | | dsv4-prefill-dual-h64 (topk 128+512), T=128 | 124.66µs | 124.64µs | 1.00x | | dsv4-prefill-dual-h64 (topk 128+512), T=2048 | 1559.56µs | 1559.35µs | 1.00x | Decode gains concentrate at small T (launch-bound); the two decode commits behind them: `quantize_q_to_smem` rewritten as a vectorized single pass (3 `bar.sync` → 1), and the decode-dsv4 IO gather reads each candidate's index once instead of twice. T=64 decode and dual-cache prefill are unchanged within noise. ## swapAB prefill (#4751) Re-benched on the PRO 6000 (#4751's table was measured on a PRO 5000), same grid, MG↔swapAB cross-checked at 5e-2 on identical inputs, `auto` bitwise-identical to forced swapAB: | shape | MG | swapAB | speedup | |---|---|---|---| | H=64, T=128 | 250.8µs | 159.7µs | 1.57× | | H=64, T=512 | 798.7µs | 565.6µs | 1.41× | | H=64, T=2048 | 2948.1µs | 2158.6µs | 1.37× | | H=64, T=8192 | 11673.6µs | 8607.7µs | 1.36× | | H=128, T=128 | 349.6µs | 267.9µs | 1.30× | | H=128, T=512 | 1348.2µs | 840.0µs | 1.60× | | H=128, T=2048 | 5330.9µs | 3011.6µs | 1.77× | | H=128, T=8192 | 21156.9µs | 11847.7µs | 1.79× | Wins everywhere; the H=64 large-T plateau (~1.4×, one CTA per token saturates ~1280 GB/s vs ~1860 at H=128) is a flat asymptote out to T=32768, so no dispatch range limit. KV layout and all parameters unchanged; both scale formats, sinks, and variable `topk_length` supported. `prefill_impl`: `"auto"` (default) / `"swapab"` / `"mg"`; forcing swapab at an ineligible shape raises. ## Dispatch: cpb model + crossover **cpb model** — analytical pick over gather bandwidth/latency, per-block overhead, and the exact list-scheduling makespan of the split grid, with an L2-footprint guard rail (at topk=1024+2176 dual the heuristic picks a single 50-chunk block at 2.7× L2 — ncu: L2 hit 69.7% vs 86.8%, costing 33%; the guard recovers it to 1.02×). Calibrated once per device inside `autotune()` tuning mode (6 fixed measurements over a ~2 GiB pool, timed as queued batches over rotating fresh index sets — launch latency overlaps execution, and the batch length keeps each set's reuse distance past an L2 turnover; small numpy LM fit; any failure = silent fallback to the C++ heuristic, so the new path can't be worse than status quo). Offline pick error vs exhaustive sweep (DRAM-cold protocol): **mean 1.011× / max 1.061×**; beats the heuristic by up to **1.37×** at mid shapes. A GPU accuracy-guard test fails loudly if a future kernel change breaks the model's assumptions, measured with the same protocol the calibration runs. Host cost ~8µs/call, memoized; zero per-replay under CUDA graphs. **Per-shape refinement** — the model's residual pick error concentrates at mid-T wave-quantization shapes (measured up to **1.35×**, e.g. DOTS3_SWA T=32: 78.0µs → 57.8µs). tuning-mode decode-form calls time the model pick ±6 candidates with the calibration protocol and persist the measured best as a per-shape override in the same tuning cache; `_resolve_cpb` consults overrides first, then the model. Across 12 production bucket shapes (T=16..64, three families, two-pass re-timing): **never worse than the model (12/12), closes every pocket to ≤1.03×**. Shapes never warmed (off-graph calls, arbitrary T, dual-cache) stay on the model. Capture-time calls only read the table/model and freeze — no measurement ever runs under graph capture or in serving. **Crossover** — per-config `decode_max_tokens` measured during the same tuning pass (probe T ∈ {4..64}, both paths, DRAM-faithful fresh indices; decode wins iff ≤ 0.95× prefill). Uncalibrated behavior is unchanged. Measured examples: | config | `decode_max_tokens` | Σ T∈{24,32,48,64}: old policy → calibrated | |---|---|---| | DSv3.2 H=128 topk=2048 (swapAB side) | 8 | 1271.4 → 494.0 µs (−61%) | | DSv4 H=64 topk=512 | 24 | 292.3 → 216.7 µs (−26%) | | DSv4 H=64 topk=128 | 16 | 132.3 → 96.7 µs (−27%) | | DSv4 H=8 topk=1024 | 64 (decode dominates) | no rerouting | Full per-probe data for all 71 calibrated configs: kernel-bench `crossover-v5` baseline. A public `calibrate_sparse_mla_sm120(device, heads=, topks=, families=, force=)` makes any envelope shape tunable outside tuning mode (idempotent skip-existing; `force=True` re-measures). ## Runtime envelopes (head counts and topk widths) - **Decode**: any H ∈ [1,128] — dedicated instantiations on the production grid (0.9–2.5% faster), one runtime-H instance otherwise, **40/40 bitwise-identical** between the two. Any `topk ≥ min_topk` (1; 513 for DOTS3_SWA so the window fits). The `_DECODE_*_DISPATCH` objects vLLM probes are membership predicates with exactly this meaning; `supported_sparse_mla_sm120_configs()` exposes the envelopes for init-time validation. Off-grid example: H=80 T=16 is 1.14× faster than the pad-to-128 workaround callers needed before. - **Prefill**: same topk rule across SG / MG / dual / swapAB. One deliberate residual asymmetry: **decode serves ragged widths (partial tail chunk, tested at topk=500); prefill requires whole 64-wide index tiles** — all production topk widths qualify, tail support needs predicated gathers + tail masking across the IO and math paths, and is deferred until a model needs it. This is safe at the routing layer: a ragged decode-form call has no prefill envelope and simply stays on decode (no crossover), and a ragged T>64 call fails loudly at the binding. 50-config parity vs the pinned build: worst **+0.94%**. One variant needed kernel-side help: DOTS3_SWA SG's BI=32 tiles are too short to cover the index→rope address-chain latency once the compile-time trip count disappeared (+24% `long_scoreboard` in NCU). The SG loop now stages the three per-tile index reads one tile ahead in registers, `if constexpr`-scoped to short tiles (unconditional staging taxed BI=64 SG +2.3%). Net: **374.6µs vs the pinned build's 380.7µs** at H=64/T=256, registers flat, `long_scoreboard` back to parity. ## Plan layer All dispatch policy lives in one memoized Python planner (`_sparse_mla_sm120_plan.py`): each variant declares its envelope once, `plan()` picks by envelope + crossover + `prefill_impl`. The C++ side is a policy-free launcher registry (the old `dispatch_v32` chain is deleted). Single-sourcing surfaced two latent upstream bugs, fixed here: prefill launchers never checked `page_block_size` against the compiled 64 (silent wrong-stride launch), and dual-cache decode-form DSv3.2-family calls silently ignored the secondary cache. ## Runner and CUDA graphs `SparseMLASm120Wrapper` holds buffers persistently: LSE pre-sized at construction, decode split-K scratch allocated only when the call actually routes to decode and cached for the instance's lifetime (a per-call temporary's freed block can be recycled into a later capture while an older graph replays into it). Capture contract: construct and warm up every captured shape before capture (or pass `out_lse`/scratch explicitly); replay is pure graph replay with zero Python. Both routing variants are correct for any T, so a crossover inside a padding bucket is at worst suboptimal, never wrong. GPU tests pin capture/replay for crossover dispatch and for runner-internal scratch. ## Compatibility Public Python API: unchanged except additive kwargs; `flashinfer.mla` exports purely additive; no-constants path behaves exactly as today. Deliberate behavior changes: - Per-shape tactic caches (`sparse_mla_sm120_decode_dsv{4,3_2}.json`) are ignored; the new calibration file is schema-versioned (v1), unrecognized versions treated as absent and recalibrated. - `autotune(True)` runs a one-time-per-device calibration (~2 GiB transient pool) instead of profiling each new shape; honors `skip_ops={"sparse_mla_sm120"}`; refuses to run under CUDA graph capture; cache writes serialized with a FileLock. - With calibration present, decode-form calls beyond the measured crossover route to prefill (the point of the feature). - T ≤ 64 shapes outside the old fixed grid now take the runtime decode instantiation instead of raising. - Prefill serves any `topk % 64 == 0` (≥ 513 for DOTS3_SWA); ragged widths fail at the binding. - Inline-scale (DSv3.2/GLM) KV caches must be contiguous through the paged entry (prefill flat-addresses the cache and crossover makes routing dynamic); contiguous padded-row caches remain decode-served and fail loudly only if prefill-routed. - `indices`/`out_lse` may be row-strided views (widening); the decode binding previously corrupted a strided `out_lse` silently. - C++ launcher entries gained row-stride parameters — internal to the JIT module, no stable ABI consumers. Out of scope (tracked follow-ups): H=64 swapAB bandwidth at large T; a pinned-topk fast path à la decode-H for DOTS3_SWA SG (locked clocks show ~2% there, boost clocks show nothing — not worth the instantiation axis on current evidence). ## Test plan All on RTX PRO 6000: **658 passed** across `test_sparse_mla_sm120{,_dispatch,_cpb_model}.py` and `test_autotuner_core.py`, pre-commit clean — including the 68-config small-T prefill matrix vs the reference (T ∈ {1..64} × SG/MG/swapAB/dual × sink/truncation), 27 C++⟺Python envelope-consistency probes, runtime-H/topk parity gates (bitwise where required), crossover routing + CUDA-graph capture/replay tests, runner scratch routing/lifetime tests, and the review-round regression tests (row-strided `out_lse`, cpb save/publish/FileLock, grid-completeness gating, padded-cache rejection, skip_ops/capture guards). This PR was prepared with AI assistance; all changes reviewed and tested locally by the submitter. --------- Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com> Co-authored-by: XingSong <sunwenhan@xfusion.com> Co-authored-by: Sam Mausberg <samuelmausberg@gmail.com> Co-authored-by: Lemon7-UP <fearless192@163.com> Co-authored-by: Luca Motz <321921718+lucamotz@users.noreply.github.com> Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com> Co-authored-by: Brian K. Ryu <bryu@nvidia.com>

…120 (vllm-project#53574) Signed-off-by: Zihua Wu <13583761+lucifer1004@users.noreply.github.com> Signed-off-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com> Co-authored-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> (cherry picked from commit 699e180)

lucifer1004 requested a review from zyongye as a code owner August 24, 2026 12:38

claude Bot reviewed Aug 24, 2026

View reviewed changes

mergify Bot added deepseek Related to DeepSeek models DSv4 bug Something isn't working labels Aug 24, 2026

This was referenced Aug 24, 2026

[DSpark] Support pipeline-parallel targets in disaggregated and aggregated serving (+ padded graph batch safety) #53577

Closed

[Bugfix] DSv4 MXFP4 selector: stop narrowing explicit aliases to their BF16 variant #53586

Merged

lucifer1004 force-pushed the pr/dsv4-c128a-decode-eidx-contiguous branch from b091bee to 3254eda Compare August 25, 2026 10:08

lucifer1004 requested review from AndreasKaratzas, WoosukKwon, mgoin, tlrmchlsmth and yewentao256 as code owners August 25, 2026 10:08

yewentao256 requested changes Aug 25, 2026

View reviewed changes

lucifer1004 force-pushed the pr/dsv4-c128a-decode-eidx-contiguous branch 2 times, most recently from e1486a7 to 0e65602 Compare August 26, 2026 08:56

lucifer1004 added a commit to lucifer1004/vllm that referenced this pull request Aug 26, 2026

Merge PR vllm-project#53574 (contiguous C128A decode topk on SM120) i…

ec77b97

…nto inferlab-integration

yewentao256 reviewed Aug 26, 2026

View reviewed changes

lucifer1004 added a commit to lucifer1004/vllm that referenced this pull request Aug 27, 2026

Merge updated PR vllm-project#53574 (inline SM120 decode-width judgme…

d4dee36

…nt) into inferlab-integration

yewentao256 reviewed Aug 27, 2026

View reviewed changes

Comment thread vllm/models/deepseek_v4/sparse_mla.py Outdated

Update vllm/models/deepseek_v4/sparse_mla.py

9d06edf

Signed-off-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com>

yewentao256 approved these changes Aug 27, 2026

View reviewed changes

yewentao256 added the ready ONLY add when PR is ready to merge/full CI is needed label Aug 27, 2026

Merge branch 'main' into pr/dsv4-c128a-decode-eidx-contiguous

097681d

lucifer1004 mentioned this pull request Aug 28, 2026

Refactor Sparse MLA SM120 flashinfer-ai/flashinfer#4802

Merged

Merge branch 'main' into pr/dsv4-c128a-decode-eidx-contiguous

f4e8530

zyongye enabled auto-merge (squash) August 31, 2026 08:04

zyongye merged commit 699e180 into vllm-project:main Aug 31, 2026
66 checks passed

tacos8me mentioned this pull request Aug 31, 2026

[New model][Multimodal] Add DeepSeek-V4-Flash-Vision-Exp support #54566

Merged

4 tasks

lucifer1004 mentioned this pull request Sep 3, 2026

[Bugfix][Hardware][NVIDIA] Fix DSV4 sparse MLA spec-decode shapes on SM120 FlashInfer path #52499

Open

Uh oh!

Conversation

lucifer1004 commented Aug 24, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

[Bugfix][SM120] DSv4: keep C128A decode topk indices contiguous on SM120

Purpose

Test plan

Notes

Uh oh!

claude Bot left a comment

Choose a reason for hiding this comment

Claude Code Review

Uh oh!

maci0 commented Aug 24, 2026

Uh oh!

yewentao256 left a comment

Choose a reason for hiding this comment

Uh oh!

yewentao256 Aug 25, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

yewentao256 Aug 25, 2026

Choose a reason for hiding this comment

Uh oh!

lucifer1004 commented Aug 26, 2026

Uh oh!

yewentao256 left a comment

Choose a reason for hiding this comment

Uh oh!

yewentao256 Aug 26, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

lucifer1004 Aug 26, 2026

Choose a reason for hiding this comment

Uh oh!

yewentao256 Aug 27, 2026

Choose a reason for hiding this comment

Uh oh!

lucifer1004 Aug 27, 2026

Choose a reason for hiding this comment

Uh oh!

yewentao256 Aug 26, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

yewentao256 left a comment

Choose a reason for hiding this comment

Uh oh!

yewentao256 commented Aug 27, 2026

Uh oh!

github-actions Bot commented Aug 27, 2026

Uh oh!

github-actions Bot commented Aug 27, 2026

Uh oh!

lucifer1004 commented Aug 28, 2026

Uh oh!

zyongye commented Aug 31, 2026

Uh oh!

github-actions Bot commented Aug 31, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

lucifer1004 commented Aug 24, 2026 •

edited

Loading

yewentao256 Aug 25, 2026 •

edited

Loading

yewentao256 Aug 26, 2026 •

edited

Loading