Skip to content

[Perf][GLM-5.3-Flash] Fill short-context decode causal top-k rows with one Triton launch - #12

Open
JaredforReal wants to merge 23 commits into
mainfrom
perf/glm53-flash-mtp-decode-cleanups
Open

JaredforReal wants to merge 23 commits into
mainfrom
perf/glm53-flash-mtp-decode-cleanups

Conversation

@JaredforReal

@JaredforReal JaredforReal commented Sep 8, 2026 •

Copy link
Copy Markdown
Owner

Purpose

When every context in a decode batch fits in index_topk (2048 tokens), the GLM-5.3-Flash kpool indexer skips top-k and fills each decode row with the exact causal indices (j for j <= position, else -1). _fill_causal_indices did this with torch ops: arange, cast, broadcast copy, compare and masked fill, i.e. 5 kernels per call, on both NVIDIA and AMD.

This PR replaces them with one Triton kernel on GPU. CPU tensors keep the torch code. The output is unchanged.

The other changes this PR used to carry are gone:

Performance

Microbenchmark on H200, 2048-wide rows. Each call is replayed 78 times in one CUDA graph, as in the decode graph, and the table shows GPU time per call:

decode rows torch (5 kernels) this PR (1 kernel)
1 5.39 us 0.88 us
8 7.14 us 1.10 us
64 7.40 us 1.22 us
256 9.36 us 1.57 us

In eager mode both are host-bound at ~33 us per call.

End-to-end serving numbers on current main: TODO.

Test Plan

  • tests/v1/attention/test_sparse_indexer_decode_seq_lens.py. The new GPU test compares the Triton fill with the torch fill for widths 8, 1000 and 2176, including positions past the row width.

Test Result

  • H200: 10 passed.
  • ROCm: not tested yet. The kernel also runs there, because rows.is_cuda is true on ROCm.

Why this is not a duplicate

_fill_causal_indices on upstream main (vllm/models/glm5next/common/sparse_indexer.py) is still the torch version. No open upstream PR touches it.

AI assistance (Claude) was used. The submitter reviewed every changed line and ran the tests above.


🤖 Generated with Claude Code

Copilot AI 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.

🟡 Changes recommended

The KDA fused projection currently bypasses LoRA/quantization logic by using raw weights, and the new Triton fast path needs an explicit HAS_TRITON guard to avoid runtime failures when Triton is disabled.

Once you've addressed the issues Copilot identified, you can request another Copilot review.

Pull request overview

This PR targets GLM-5.3-Flash MTP decode performance by reducing kernel-launch overhead in two hotspots and fixes a configuration plumbing issue so --mamba-ssm-cache-dtype is honored for the KDA recurrent state.

Changes:

  • Plumb cache_config.mamba_ssm_cache_dtype into GLM5Next KDA state dtype selection.
  • Reduce decode hot-path launches by (a) introducing a Triton kernel to fill causal top-k rows in one launch and (b) batching adjacent f_b/g_b projections into a single strided-batched GEMM.
  • Add a CUDA test ensuring the fused causal-fill matches the torch formulation.
File summaries
File Description
vllm/models/glm5next/nvidia/kda.py Pass through SSM cache dtype for KDA state; fuse f_b/g_b projection into one batched GEMM with a cached stacked weight.
vllm/model_executor/layers/sparse_attn_indexer_kpool.py Add a Triton kernel fast-path to fill causal indices in one launch for CUDA tensors.
tests/v1/attention/test_sparse_indexer_decode_seq_lens.py Add a CUDA test comparing the fused causal-fill with the torch implementation.
Review details
  • Files reviewed: 3/3 changed files
  • Comments generated: 2
  • Review effort level: Lite

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

Comment on lines +211 to +219
def _fill_causal_indices(rows: torch.Tensor, positions: torch.Tensor) -> None:
"""rows[t, j] = j for j <= positions[t], else -1 (one launch on CUDA)."""
if rows.is_cuda and rows.shape[0] > 0:
block = 1024
grid = (rows.shape[0], triton.cdiv(rows.shape[1], block))
_fill_causal_indices_kernel[grid](
rows, rows.stride(0), positions, rows.shape[1], BLOCK=block
)
return
Comment thread vllm/models/glm5next/nvidia/kda.py Outdated
Comment on lines +313 to +317
if self._fg_b_weight is None:
self._fg_b_weight = torch.stack(
[self.f_b_proj.weight.t(), self.g_b_proj.weight.t()]
).contiguous()
fg_a = projected[:, fg_start : fg_start + 2 * self.head_dim]
andakai and others added 22 commits October 13, 2026 03:05
…vllm-project#61183)

Signed-off-by: Dakai An <dakaian108@gmail.com>
Signed-off-by: Yifan Qiao <yifanqiao@inferact.ai>
Co-authored-by: Yifan Qiao <yifanqiao@inferact.ai>
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
…hat message parsing (vllm-project#56129)

Signed-off-by: pctablet505 <55033230+pctablet505@users.noreply.github.com>
Co-authored-by: Claude Sonnet 5.5 <noreply@anthropic.com>
…vllm-project#61299)

Signed-off-by: Yoni Gozlan <74535834+yonigozlan@users.noreply.github.com>
Signed-off-by: Bugen Zhao <i@bugenzhao.com>
Co-authored-by: Bugen Zhao <i@bugenzhao.com>
…olution (vllm-project#61272)

Signed-off-by: Yuan Li <dskwelmcy@163.com>
Co-authored-by: Cyrus Leung <tlleungac@connect.ust.hk>
… response parser (vllm-project#61297)

Signed-off-by: Yoni Gozlan <74535834+yonigozlan@users.noreply.github.com>
Signed-off-by: Bugen Zhao <i@bugenzhao.com>
Co-authored-by: Bugen Zhao <i@bugenzhao.com>
…ed pixels (vllm-project#61104)

Signed-off-by: godmook <cmoh4135@naver.com>
Co-authored-by: cjackal <44624812+cjackal@users.noreply.github.com>
…-project#57911)

Signed-off-by: Chaojun Zhang <chaojun.zhang@intel.com>
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Co-authored-by: xiangdong <40376367+zxd1997066@users.noreply.github.com>
…ct#60322)

Signed-off-by: Divakar Verma <divakar.verma@amd.com>
Co-authored-by: Claude Sonnet 5 <noreply@anthropic.com>
)

Signed-off-by: Sriram Kumar <sriramkumar.kishorekumar@amd.com>
Co-authored-by: Douglas Lehr <91553416+dllehr-amd@users.noreply.github.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
…llm-project#51309)

Signed-off-by: Sriram Kumar <sriramkumar.kishorekumar@amd.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
Co-authored-by: Douglas Lehr <91553416+dllehr-amd@users.noreply.github.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>
Co-authored-by: Simon Danielsson <70206058+simondanielsson@users.noreply.github.com>
Signed-off-by: Tianyu Guo <guoty@inferact.ai>
Signed-off-by: Isotr0py <Isotr0py@outlook.com>
Co-authored-by: Codex <noreply@openai.com>
Co-authored-by: Isotr0py <Isotr0py@outlook.com>
…lm-project#61285)

Signed-off-by: NickLucche <nicolo.lucchesi@mistral.ai>
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
…parse indexer (vllm-project#61149)

Signed-off-by: Yongye Zhu <zyy1102000@gmail.com>
Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…TP (vllm-project#61235)

Signed-off-by: aoshen524 <aoshen524@gmail.com>
Co-authored-by: aoshen524 <aoshen524@gmail.com>
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
Co-authored-by: linitra24 <renshuang.zhou@daocloud.io>
…atically (vllm-project#61202)

Signed-off-by: khluu <khluu000@gmail.com>
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
…m-project#60602)

Signed-off-by: kimseunghyun <hellonypcteam@gmail.com>
Signed-off-by: kimseunghyun <incongnito12@gmail.com>
Signed-off-by: Thien Tran <gau.nernst@yahoo.com.sg>
Co-authored-by: kimseunghyun <hellonypcteam@gmail.com>
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
Co-authored-by: Thien Tran <gau.nernst@yahoo.com.sg>
…t#60043)

Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>
Signed-off-by: Jiangyun Zhu <zhujiangyun@inferact.ai>
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
…okups in ECExampleConnector (vllm-project#60826)

Signed-off-by: zhouyou9505 <zhouyou9505@gmail.com>
vllm-project#61350)

Signed-off-by: Stefan Koncarevic <stefan.koncarevic@amd.com>
Signed-off-by: Stefan Koncarevic <Stefan.Koncarevic@amd.com>
…llm-project#57410)

Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>
Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
…end is marlin (vllm-project#60219)

Signed-off-by: Vadim Gimpelson <vadim.gimpelson@gmail.com>
Co-authored-by: Claude <noreply@anthropic.com>
@JaredforReal

Copy link
Copy Markdown
Owner Author

Closing: superseded upstream.

…h one Triton launch

The kpool indexer's short-context decode path (max_seq_len <= index_topk)
filled the causal top-k rows with five kernels per call. Replace them with
a single Triton kernel on GPU; CPU tensors keep the torch formulation.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Signed-off-by: Jared Wen <jaredwen@inferact.ai>
@JaredforReal JaredforReal reopened this Oct 13, 2026
@JaredforReal
JaredforReal force-pushed the perf/glm53-flash-mtp-decode-cleanups branch from 5066d2c to b1a0898 Compare October 13, 2026 08:59
@JaredforReal JaredforReal changed the title [Perf][GLM-5.3-Flash] MTP decode launch-count cleanups + honor --mamba-ssm-cache-dtype [Perf][GLM-5.3-Flash] Fill short-context decode causal top-k rows with one Triton launch Oct 13, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.