Skip to content

[Bugfix] GLM-5.3-Flash: launch the kpool paged MQA logits in the varlen mode its schedule was built with - #55270

Open
ivanium wants to merge 4 commits into
vllm-project:mainfrom
ivanium:fix/glm-kpool-varlen-indices
Open

ivanium wants to merge 4 commits into
vllm-project:mainfrom
ivanium:fix/glm-kpool-varlen-indices

Conversation

@ivanium

@ivanium ivanium commented Sep 4, 2026

Copy link
Copy Markdown
Collaborator

Note

Update: the original hang issue is caused by a stale deepgemm 2.5.0 version. Updating to 2.6.1 resolved it. This PR is still a little bit helpful to get metadata aligned, but it did not fix anything (there is nothing to fix)

Purpose

Make the kpool indexer's DeepGEMM paged-MQA launch agree with the schedule metadata it was built with. On SM100 DeepseekV32IndexerMetadataBuilder schedules the paged MQA logits in varlen mode (per-token rows tagged with request indices, get_paged_mqa_logits_metadata(..., indices=...)), and sparse_attn_indexer.py passes indices=decode_metadata.indices to fp8_fp4_paged_mqa_logits. The kpool indexer did not, so DeepGEMM selects the non-varlen next_n=1 kernel and runs it on a varlen schedule. Under MTP a request's verify rows share one request index, so the two decompositions differ whenever a request's rows straddle a 256-state split boundary (1024 tokens at index_kpool=4).

What that mismatch does depends on the DeepGEMM revision, verified with a standalone kernel script (same fp8 cache, block table and weights; schedule built with indices; launched with and without indices; compared against a pure non-varlen reference):

DeepGEMM launch without indices on a varlen schedule
2.5.0 (891d57b, what a stale local FetchContent checkout builds) deadlock: the exact-end-match scheduler walks past an end it can never hit; the TMA warp exits and the math warps spin on the Q-stage mbarrier. Reproduced 3x end to end (GLM-5.3-Flash MTP3 GSM8K hangs within 4 min, user-triggered CUDA core dumps show the single live CTA) and in the 40-row script.
2.6.1 (8b1392b9, the repo pin) tolerated: no hang, logits bit-identical to the reference over 640 rows including 80 straddling requests; GLM MTP3 GSM8K on main runs fine (0.9704 / 0.9731).

So with the pinned DeepGEMM this is a latent contract violation rather than a live bug. The one-line fix matches the non-kpool indexer, removes the dependence on the scheduler tolerating a mismatched launch, and keeps the kpool path working on DeepGEMM revisions with the strict scheduler.

Duplicate check: #55201 touches the kpool top-k sentinel handling, not this call; no open PR passes indices here.

Test Plan

  • Standalone DeepGEMM script described above, both DeepGEMM revisions, 40-row and 640-row batches.
  • GLM-5.3-Flash on GB200, --speculative-config '{"method":"mtp","num_speculative_tokens":3}', GSM8K 5-shot x2 (2638 requests, concurrency 32), single node TP4 and TP4-to-TP4 NIXL PD.

Test Result

  • DeepGEMM 2.5.0, GSM8K MTP3: hangs on main within 4 min every time; with the fix, single node 2638/2638, 0 request errors, 0.9723 / 0.9738 (0.9727 / 0.9742 without MTP); PD 2638/2638, 0 errors, 0.9708 / 0.9719.
  • DeepGEMM 2.6.1, GSM8K MTP3 single node with the fix: eager 0.9704 / 0.9719, CUDA graphs 0.97, 0 errors. main without the fix: 0.9704 / 0.9731, 0 errors.

No unit test: the observable failure is a device deadlock on one DeepGEMM revision, and the op needs a live GPU, DeepGEMM and a populated kpool cache to reach the call.

AI assistance (Claude Code) was used for the capture, analysis and fix; the submitter reviewed the change.

…gits

On SM100 the indexer metadata builder schedules DeepGEMM's paged MQA logits
in varlen mode (per-token rows tagged with request indices), but the kpool
indexer launched the kernel without `indices`, so the non-varlen kernel ran
on a varlen schedule. The varlen scheduler pairs adjacent rows of a request
into one atom charged with the later row's context length; the non-varlen
walk treats every row as its own request and stops only on an exact
(q_atom, kv_split) end match. Once a pair straddles a 256-state split
boundary the walk steps past an end marker it can never hit, the TMA warp
stops prefetching Q and the math warps wait forever on the Q-stage barrier:
100% GPU utilization with idle host threads until the engine times out.
Single-token decode has one row per request, so both schedules coincide and
the bug never shows without speculative decoding.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Signed-off-by: Yifan Qiao <yifanqiao@inferact.ai>
@coderabbitai

coderabbitai Bot commented Sep 4, 2026

Copy link
Copy Markdown
Contributor

Review Change Stack

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info
⚙️ Run configuration

Configuration used: Repository UI

Review profile: CHILL

Plan: Team

Run ID: 49bdc319-b2b7-45ba-8775-5bcd4d599a47

📥 Commits

Reviewing files that changed from the base of the PR and between d9e2b52 and d207fb9.

📒 Files selected for processing (1)
  • vllm/model_executor/layers/sparse_attn_indexer_kpool.py

Included review availability: Your plan provides up to 10 included reviews per hour; 9 remain after this review.


Walkthrough

The decode path now passes decode_metadata.indices to the non-ROCm fp8_fp4_paged_mqa_logits operation.

Changes

Decode attention index propagation

Layer / File(s) Summary
Forward decode indices
vllm/model_executor/layers/sparse_attn_indexer_kpool.py
The decode-path fp8_fp4_paged_mqa_logits call now receives decode_metadata.indices.

Estimated code review effort: 1 (Trivial) | ~2 minutes

Suggested reviewers: zjy0516

Merge Risk: ⚪ Minimal · up to d207f

The decode path now supplies varlen indices to paged MQA logits, preventing the affected speculative-decoding hang. The change is narrowly scoped and has no remaining merge-readiness risk.

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name Status Explanation Resolution
Docstring Coverage ⚠️ Warning Docstring coverage is 0.00% which is insufficient. The required threshold is 80.00%. Docstring coverage is scoped to functions touched by this diff. Analyzed 1 functions across 1 files. Write docstrings for the functions missing them to satisfy the coverage threshold.
✅ Passed checks (4 passed)
Check name Status Explanation
Linked Issues check ✅ Passed Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check ✅ Passed Check skipped because no linked issues were found for this pull request.
Title check ✅ Passed The title clearly identifies the GLM-5.3-Flash bugfix and states that the kpool paged MQA logits launch must use the schedule's varlen mode. It accurately summarizes the primary change.
Description check ✅ Passed The description directly explains the metadata and kernel launch mismatch, the affected DeepGEMM behavior, the one-line fix, and the validation results. It is fully related to the changeset.

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

Comment @coderabbitai help to get the list of available commands.

@claude claude 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.

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.

@JaredforReal JaredforReal left a comment

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.

while we understand we should align with ds's pattern passing indices=decode_metadata.indices down to kernel, the mtp hanging mechanism need a deepdive and repro to comfirm

@ZJY0516

ZJY0516 commented Sep 4, 2026

Copy link
Copy Markdown
Member

/ci run

@ZJY0516 ZJY0516 added the ready ONLY add when PR is ready to merge/full CI is needed label Sep 4, 2026
@github-actions

github-actions Bot commented Sep 4, 2026

Copy link
Copy Markdown

✅ Triggered Buildkite CI #87214 for commit d207fb901458.

@ivanium ivanium changed the title [Bugfix] GLM-5.3-Flash: pass varlen indices to the kpool paged MQA logits (fixes MTP decode hang) [Bugfix] GLM-5.3-Flash: launch the kpool paged MQA logits in the varlen mode its schedule was built with Sep 4, 2026
@ivanium

ivanium commented Sep 4, 2026

Copy link
Copy Markdown
Collaborator Author

Update: the hang issue is caused by a stale deepgemm 2.5.0 version. Updating to 2.6.1 resolved it. This PR is still a little bit helpful to get metadata aligned, but it did not fix anything (there is nothing to fix)

@ivanium
ivanium enabled auto-merge (squash) September 12, 2026 05:05
@wongsingfo

Copy link
Copy Markdown

One additional data point: this mismatch is not only relevant to the hang behavior seen with older DeepGEMM revisions; it also has a noticeable performance impact with the current varlen-capable path.

On GLM-5.3-Flash with context length 120K, batch size 40, and speculative decoding steps = 5, passing decode_metadata.indices changes sm100_paged_mqa_logits from the non-varlen instantiation (is_varlen=false, next_n=1) to the varlen instantiation (is_varlen=true). In an nsys capture, kernel latency dropped from ~214.3 us to ~51.8 us, about a 4.14x speedup for this kernel.

Reference: #56811

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

Labels

bug Something isn't working glm ready ONLY add when PR is ready to merge/full CI is needed

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants