Skip to content

[Attention][DSA] support dcp for FLASHINFER_MLA_SPARSE - #46076

Merged
LucasWilkinson merged 27 commits into
vllm-project:mainfrom
ZJY0516:glm52-dcp-support
Jul 1, 2026
Merged

LucasWilkinson merged 27 commits into
vllm-project:mainfrom
ZJY0516:glm52-dcp-support

Conversation

@ZJY0516

@ZJY0516 ZJY0516 commented Jun 18, 2026 •

Copy link
Copy Markdown
Member

Purpose

add dcp support for GLM 5.2

Test Plan

VLLM_FLASHINFER_WORKSPACE_BUFFER_SIZE=1073741824
vllm serve zai-org/GLM-5.2-FP8 \
    --host 0.0.0.0 \
    --port 8000 \
    --kv-cache-dtype fp8_e4m3 \
    --tensor-parallel-size 8 \
    --decode-context-parallel-size 4 \
    --tool-call-parser glm47 \
    --enable-auto-tool-choice \
    --reasoning-parser glm45

logs

DCP 4, TP 8
(EngineCore pid=88967) INFO 06-18 17:02:00 [kv_cache_utils.py:2078] GPU KV cache size: 10,853,376 tokens
(EngineCore pid=88967) INFO 06-18 17:02:00 [kv_cache_utils.py:2079] Maximum concurrency for 1,048,576 tokens per request: 10.35x

TP 8
(EngineCore pid=93451) INFO 06-18 17:27:37 [kv_cache_utils.py:2078] GPU KV cache size: 2,765,568 tokens
(EngineCore pid=93451) INFO 06-18 17:27:37 [kv_cache_utils.py:2079] Maximum concurrency for 1,048,576 tokens per request: 2.64x
lm_eval --model local-completions --model_args "model=zai-org/GLM-5.2-FP8,base_url=http://0.0.0.0:8000/v1/completions,tokenized_requests=False,tokenizer_backend=None,num_concurrent=128" --tasks gsm8k --num_fewshot 5

Test Result

Tasks Version Filter n-shot Metric Value Stderr
gsm8k 3 flexible-extract 5 exact_match ↑ 0.9340 ± 0.0068
strict-match 5 exact_match ↑ 0.9333 ± 0.0069

Perf test

GB200

vllm serve lukealonso/GLM-5.2-NVFP4 --kv-cache-dtype fp8_e4m3 -tp 4 [-dcp 4]

concurrency 256 decode performance

dcp 4
Avg prompt throughput: 0.0 tokens/s, Avg generation throughput: 4071.3 tokens/s, Running: 241 reqs, Waiting: 15 reqs, GPU KV cache usage: 99.7%, Prefix cache hit rate: 0.6% 
tp 4
Avg prompt throughput: 0.0 tokens/s, Avg generation throughput: 2508.7 tokens/s, Running: 63 reqs, Waiting: 193 reqs, GPU KV cache usage: 99.2%, Prefix cache hit rate: 0.0% 

But dcp prefill speed is only 1/2 of tp, so the e2e throughput is not ideal


Essential Elements of an Effective PR Description Checklist
  • The purpose of the PR, such as "Fix some issue (link existing issues this PR will resolve)".
  • The test plan, such as providing test command.
  • The test results, such as pasting the results comparison before and after, or e2e results
  • (Optional) The necessary documentation update, such as updating supported_models.md and examples for a new model.

Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>
@mergify

mergify Bot commented Jun 18, 2026

Copy link
Copy Markdown
Contributor

Documentation preview: https://vllm--46076.org.readthedocs.build/en/46076/

@mergify mergify Bot added documentation Improvements or additions to documentation nvidia v1 labels Jun 18, 2026

@chatgpt-codex-connector chatgpt-codex-connector Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: 981b6e075b

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

  • Open a pull request for review
  • Mark a draft as ready
  • Comment "@codex review".

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".

ZJY0516 added 2 commits June 18, 2026 17:14
Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>
Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>
Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>
@mergify

mergify Bot commented Jun 22, 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, @ZJY0516.

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



@triton.jit
def _convert_req_index_to_global_index_dcp_kernel(

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

do we need a whole new kernel? can we just add support to triton_convert_req_index_to_global_index?

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done in 30f6fec — folded into _convert_req_index_to_global_index_kernel via DCP_SIZE/DCP_RANK constexprs.

Co-authored-by: OpenAI Codex <codex@openai.com>

Signed-off-by: Jingyi Yang <girasoleyang@gmail.com>
@drakosha

Copy link
Copy Markdown
Contributor

Hi @ZJY0516 @GirasoleY @LucasWilkinson — following up on the suggestion to consolidate sparse MLA DCP support here.

We've been working on #45426 (DCP for FLASHMLA_SPARSE). Rather than keep a duplicate exact-merge path, we'd like to build on this PR's shared indexer merge (_merge_dcp_topk_global + the radix stable_top_k_from_candidates kernel) and add the FlashMLA-specific layer on top.

Coverage this PR doesn't reach yet. On Hopper / SM90 (e.g. H200), the sparse MLA backend the selector picks today is FLASHMLA_SPARSE — the generic FLASHINFER_MLA_SPARSE path is currently in the SM10 priority block in vllm/platforms/cuda.py. The production path we need is also fp8_ds_mla KV cache via FlashMLA, while this PR validates the FlashInfer sparse path with fp8_e4m3. So #46076 gives the shared DCP/indexer infrastructure, but it doesn't yet cover the H200 + fp8_ds_mla sparse path — where any Hopper deployment of these models (e.g. GLM-5.2 at TP4/DCP4) sits today.

What we'd add on top of your shared contract:

  • FLASHMLA_SPARSE DCP localization: consume the indexer's global ids → this rank's physical fp8_ds_mla slots.
  • FlashMLA decode plumbing: base-e LSE (FlashMLA already returns base-e, matching vLLM's DCP merge convention) + the fp8-kernel out/LSE handling.
  • MTP + DCP + full cudagraphs (supports_dcp_with_varlen).
  • GLM-5.2 validation on 4×H200.

(Plus fail-closed guards for combinations we won't support under DCP, e.g. compress_ratio > 1. We don't plan to upstream the union/approximate path from #45426 — only the exact merge, which preserves the original top-k set.)

The intended end state, if this sounds right: one shared exact top-k merge + radix kernel for sparse MLA, with FlashInfer and FlashMLA differing only in the backend-specific global-id → physical-slot localization and LSE/kernel plumbing.

Our preference is to open this as a stacked PR on top of #46076 — so this PR keeps moving while the FlashMLA-specific validation stays isolated, landing after #46076 merges. If you'd rather fold it in here after the rebase, we're happy to do that.

One contract question before we move code: can we treat "the indexer emits global ids, and each backend localizes them" as the intended shared contract? #45426 currently returns local shard indices from the indexer; we'll switch to the global-id contract so both backends share one merge.

We'll also leave two small inline notes separately: the fp64 composite fallback isn't strictly equivalent to the radix key for arbitrary fp32 scores, and the localizer asserts cp_kv_cache_interleave_size == 1 while the merge helper already handles general interleave.

Comment on lines +175 to +176
local_cu_seqlen_ks: torch.Tensor | None
local_cu_seqlen_ke: torch.Tensor | None

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

are these needed? cant we just make cu_seqlen_ks and cu_seqlen_ke local when DCP is enabled?

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done in 3b1b408 — kernel writes local row bounds; the extra local_cu_seqlen fields are gone.

GirasoleY and others added 4 commits June 30, 2026 21:43
Address review comment: the second `_can_use_cutedsl_candidate_topk` check
looked redundant with the first. Remove it and the separate fallback CuteDSL
selector branch, leaving one guard (`_can_use_cutedsl_dcp_merge`) that gates the
whole CuteDSL path -- the Triton candidate pack kernel plus the CuteDSL
stable-topk selector.

Drop the `logits.shape[1] > 0` clause from the guard so it now also covers the
empty-local-shard case the second check used to allow: the Triton pack kernel
already masks empty shards (all topk_indices == -1 -> score=-inf, global_id=-1),
so no second guard is needed. When the guard is not satisfied, fall back to the
pure-PyTorch pack + deterministic fp64 stable top-k.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: Jingyi Yang <girasoleyang@gmail.com>
…ack)

Address review comment: for the first cut (Blackwell/Hopper) the DCP sparse
indexer merge should require the CuteDSL path rather than silently falling back
to an unoptimized PyTorch implementation.

- Drop the `rows >= 256` gate (a perf heuristic, not a kernel limit -- the
  CuteDSL selector grid is one CTA per row, so it works at any batch size).
- Replace `_can_use_cutedsl_dcp_merge` (+ the PyTorch pack / fp64 stable-topk /
  local->global helpers) with `_assert_cutedsl_dcp_merge_supported`, which raises
  a clear error when CuteDSL is unavailable, tensors aren't on CUDA, dtypes are
  wrong, or index_topk is not in (512, 1024, 2048) (the selector's radix sizing).
  In practice index_topk is always 2048 for supported DSA checkpoints, and the
  sibling FLASHINFER_MLA_SPARSE_SM120 backend already hard-requires it.
- `_merge_dcp_topk_global` now runs the Triton pack + CuteDSL selector directly.

Remove the now-obsolete small-k merge unit tests (they exercised the deleted
fallback at k=4/6, which the CuteDSL kernel cannot run); kernel-level pack/select
correctness stays covered by test_cutedsl_dcp_candidate_pack_and_select_matches_reference.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: Jingyi Yang <girasoleyang@gmail.com>
Address review comment: instead of a separate argsort+gather pass
(_compact_valid_to_front), pack each row's valid (rank-owned) slots to a
contiguous prefix inside _convert_req_index_to_global_index_kernel using
valid_count_ptr as an atomic slot allocator.

Add a COMPACT_TO_FRONT path: a per-tile exclusive prefix sum gives each valid
lane a distinct local offset, and one atomic add of the tile's valid count
reserves a contiguous base across the racing column tiles. The out buffer is
pre-filled with -1 so the unwritten tail stays -1. Prefix order is unspecified
(only the selected set matters), matching the kernel's existing contract. The
non-DCP / non-compact path (input column == output column) is unchanged.

DCP filtering marks non-owned slots -1, creating interior gaps; the trtllm-gen
sparse kernel reads the first valid_count entries of each row, so the valid
slots must form a contiguous prefix. Drop _compact_valid_to_front.

Add test_dcp_filter_compaction_matches_reference: wide multi-tile rows (so the
cross-tile atomic allocation is exercised) with interior -1 gaps, checked per
row against a reference filter for both ranks and interleave 1/2 -- valid_count,
the packed prefix set, and the -1 tail. Existing single-row and end-to-end
physical-attention compaction tests still pass.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: Jingyi Yang <girasoleyang@gmail.com>
Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>
@LucasWilkinson
LucasWilkinson merged commit c5200d3 into vllm-project:main Jul 1, 2026
105 checks passed
@github-project-automation github-project-automation Bot moved this from Ready to Done in NVIDIA Jul 1, 2026
@ZJY0516
ZJY0516 deleted the glm52-dcp-support branch July 1, 2026 04:41
alexbi29 pushed a commit to alexbi29/vllm that referenced this pull request Jul 1, 2026
124 upstream commits. Conflicts resolved (4 files), preserving the SM12x stack
while absorbing upstream DSv4 work:
- sparse_attn_indexer.py: absorb vllm-project#46076 dcp (inert unless dcp_world_size>1),
  keep SM120 short-row/persistent decode dispatch, dedupe the vllm-project#47164 family-120
  cooperative-topk gate (upstream independently landed our fix).
- sparse_swa.py: absorb vllm-project#46995 DSpark non-causal path, keep our
  _init_reorder_batch_threshold(supports_spec_as_decode) MTP threshold.
- indexer.py: union has_prefilling_rows + dcp_local_seq_lens.
- protocol.py: union DSv4 thinking-mode sampling override + vllm-project#35076 stop_token_ids.

Signed-off-by: jasl <jasl9187@hotmail.com>
noooop pushed a commit to noooop/vllm that referenced this pull request Jul 9, 2026
…46076)

Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>
Signed-off-by: Jingyi Yang <girasoleyang@gmail.com>
Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Signed-off-by: GirasoleY <girasoleyang@gmail.com>
Co-authored-by: Jingyi Yang <girasoleyang@gmail.com>
Co-authored-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Signed-off-by: wang.yuqi <yuqi.wang@daocloud.io>
drakosha added a commit to drakosha/vllm that referenced this pull request Jul 18, 2026
…ath + MTP

Rebase of vllm-project#46514 onto current main. The sparse-indexer DCP machinery the
original series carried (exact global top-k merge, DCP index filtering,
decode seq-len localization) has since landed upstream via vllm-project#46076 and
follow-ups, so this now contains only the FlashMLA-sparse backend side,
rebased onto the SparseMLACommon* split from vllm-project#47327 (DCP plumbing such as
dcp_world_size/dcp_rank and cp_kv_cache_interleave_size now comes from the
shared base builder/metadata instead of backend-local declarations):

- flashmla_sparse: enable DCP on the mixed-batch fp8 path; return decode
  LSE for the cross-rank merge; neutralize empty rows (out=0, lse=-inf);
  localize the indexer's global top-k ids via
  triton_filter_and_convert_dcp_index (compact_valid_to_front=False, the
  fp8 kernel masks scattered -1 natively); MTP spec-as-decode under DCP;
  keep the head-count reorder threshold under DCP via
  supports_dcp_with_varlen when cp_kv_cache_interleave_size == 1;
  reject configs where local and DCP-gathered head counts pad to
  different fp8 kernel envelopes; reject bf16 kv-cache under DCP
  (untested) and non-ag_rs comm backends.
- flashmla (dense): flatten decode LSE [batch, heads, seq] ->
  [tokens, heads] for the DCP reducer under spec-decode.
- tests: fp8 DCP head-envelope guard cases.

Validated on 4xH200 with GLM-5.2-NVFP4 (TP4/DCP4/EP + MTP,
kv_cache_dtype=fp8_ds_mla): cudagraph capture (FULL_AND_PIECEWISE),
coherent output, exact 198k-token needle retrieval, 78 tok/s N=1 /
854 tok/s aggregate N=32.

Co-authored-by: Claude
Signed-off-by: Mikhail Kostryukov <mike@triptrack.net>
philippesic pushed a commit to philippesic/vllm-semantic-cache that referenced this pull request Jul 19, 2026
…46076)

Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>
Signed-off-by: Jingyi Yang <girasoleyang@gmail.com>
Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Signed-off-by: GirasoleY <girasoleyang@gmail.com>
Co-authored-by: Jingyi Yang <girasoleyang@gmail.com>
Co-authored-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
michael-gpu pushed a commit to practical-inference/vllm-ampere-extended that referenced this pull request Oct 20, 2026
…46076)

Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>
Signed-off-by: Jingyi Yang <girasoleyang@gmail.com>
Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Signed-off-by: GirasoleY <girasoleyang@gmail.com>
Co-authored-by: Jingyi Yang <girasoleyang@gmail.com>
Co-authored-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

documentation Improvements or additions to documentation nvidia ready ONLY add when PR is ready to merge/full CI is needed v1

Projects

Status: Done

Development

Successfully merging this pull request may close these issues.

4 participants