Skip to content

[KDA] Default GLM-5-Next spec verify to the CuTe DSL KDA kernel - #42905

Merged
kpham-sgl merged 1 commit into
mainfrom
mmangkad/glm5next-kda-cutedsl-verify
Oct 7, 2026
Merged

[KDA] Default GLM-5-Next spec verify to the CuTe DSL KDA kernel#42905
kpham-sgl merged 1 commit into
mainfrom
mmangkad/glm5next-kda-cutedsl-verify

Conversation

@mmangkad

@mmangkad mmangkad commented Oct 7, 2026 •
edited by github-actions Bot
Loading

Copy link
Copy Markdown
Contributor

Summary

Defaults --linear-attn-verify-backend to nv_cutedsl for GLM-5-Next (Glm5NextForConditionalGeneration) when speculative decoding is on, the flag is unset, and the GPU is SM100-class. Target verify then runs the fused CuTe DSL KDA MTP kernel (conv1d + recurrence in one launch); shapes outside its contract use the Triton verify kernel.

Results

GLM-5.3-Flash FP8, 4x GB300, TP4, EAGLE 5/1/6 (SGLANG_SIMULATE_ACC_LEN=5).

Serving, random 8k in / 1k out, prefix cache flushed:

C Output tok/s (main -> PR) Median TPOT ms (main -> PR) Median ITL ms (main -> PR)
1 543.1 -> 550.4 (+1.3%) 1.606 -> 1.575 (-1.9%) 1.61 -> 1.57 (-2.5%)
48 3770.0 -> 3823.7 (+1.4%) 8.287 -> 8.078 (-2.5%) 5.08 -> 4.88 (-3.9%)

Target verify per layer (Triton conv1d update + recurrence vs fused CuTe DSL kernel), 16 heads, 6 tokens per request, triton.testing.do_bench_cudagraph on fixed buffers, median of 3:

Requests ReplaySSM rings: Triton -> CuTe (us) Intermediate-state snapshots: Triton -> CuTe (us)
1 14.66 -> 6.88 (-53.1%) 17.31 -> 12.53 (-27.6%)
2 15.76 -> 6.93 (-56.0%) 18.58 -> 12.58 (-32.3%)
4 16.02 -> 6.96 (-56.6%) 20.03 -> 12.64 (-36.9%)
8 19.41 -> 7.17 (-63.1%) 22.58 -> 13.17 (-41.7%)
12 20.53 -> 11.40 (-44.5%) 31.65 -> 23.31 (-26.4%)
16 22.57 -> 11.54 (-48.9%) 36.12 -> 24.71 (-31.6%)
24 38.24 -> 18.36 (-52.0%) 56.45 -> 38.09 (-32.5%)
32 41.07 -> 23.13 (-43.7%) 64.19 -> 49.05 (-23.6%)
48 57.32 -> 32.89 (-42.6%) 89.75 -> 71.73 (-20.1%)

Correctness

GSM8K (sgl-eval, all 1319 questions, temperature 1.0, top_p 0.95, max_tokens 16384):

GSM8K Avg accept length
main 0.9719 3.466
PR 0.9719 3.477

CI States

Latest PR Test (Base): ❌ Run #37571052144
Latest PR Test (Extra): ❌ Run #37571051946
Latest PR Test (AMD ROCm 10): ❌ Run #37571052093

@mmangkad

mmangkad commented Oct 7, 2026

Copy link
Copy Markdown
Contributor Author

/rerun-test test/registered/e2e/models/test_glm53_flash_b200.py

@github-actions

github-actions Bot commented Oct 7, 2026 •
edited
Loading

Copy link
Copy Markdown
Contributor

🚀 4-gpu-b200 (1 test): ✅ View workflow run 

cd test/ && python3 registered/e2e/models/test_glm53_flash_b200.py

kpham-sgl
kpham-sgl approved these changes Oct 7, 2026
View reviewed changes

@kpham-sgl kpham-sgl left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

nice catch

@kpham-sgl
kpham-sgl merged commit f47d895 into main Oct 7, 2026
102 of 121 checks passed
@kpham-sgl
kpham-sgl deleted the mmangkad/glm5next-kda-cutedsl-verify branch October 7, 2026 05:31
@JackChuang JackChuang mentioned this pull request Oct 16, 2026
Open
4 of 5 tasks
HanHan009527 pushed a commit to bytedance-iaas/sglang that referenced this pull request Oct 19, 2026
@JackChuang @HanHan009527
[Perf] [NVIDIA] KDA: Extend the CuTe DSL MTP verify kernel to SM90
b6df297 
The fused CuTe DSL KDA verify kernel for MTP / speculative decoding
(fused_kda_decode_mtp_dspark) used paired FP32 instructions
(mul_packed_f32x2 / fma_packed_f32x2) that NVVM only accepts for SM100
and newer, and the KDA backend only ran it on SM100, so Hopper always
fell back to the unfused Triton verify path.

- kda_decode_mtp.py: a USE_PACKED_FMA constexpr (capability major >= 10)
  keeps the packed path on SM100 and newer and uses scalar FP32 FMAs on
  SM90, as cutedsl_gdn_mtp_ring.py does. SM100 compiles the same code as
  before.
- kda_backend.py: the CuTe verify gate accepts SM90 as well as SM100, so
  servers that select nv_cutedsl (explicitly, or Kimi-K3 through its
  hook) run the CuTe kernel on Hopper.
- glm5_next_hook.py: extend the GLM-5-Next spec default
  --linear-attn-verify-backend nv_cutedsl (sgl-project#42905) to SM90 when
  --enable-linear-replayssm-spec is on; SM100 is unchanged.
- Tests: the ReplaySSM ring parity test also runs on the 1-gpu-large
  runner (H100); a new CuTe-vs-Triton verify parity test; CPU tests for
  the backend gate (SM90 and SM100 admitted, SM80 and SM120 rejected)
  and for the GLM-5-Next hook (SM90 with and without the ring).
JackChuang added a commit to JackChuang/sglang that referenced this pull request Oct 22, 2026
@JackChuang @HanHan009527
[Perf] [NVIDIA] KDA: Extend the CuTe DSL MTP verify kernel to SM90
c3cfc38 
The fused CuTe DSL KDA verify kernel for MTP / speculative decoding
(fused_kda_decode_mtp_dspark) used paired FP32 instructions
(mul_packed_f32x2 / fma_packed_f32x2) that NVVM only accepts for SM100
and newer, and the KDA backend only ran it on SM100, so Hopper always
fell back to the unfused Triton verify path.

- kda_decode_mtp.py: a USE_PACKED_FMA constexpr (capability major >= 10)
  keeps the packed path on SM100 and newer and uses scalar FP32 FMAs on
  SM90, as cutedsl_gdn_mtp_ring.py does. SM100 compiles the same code as
  before.
- kda_backend.py: the CuTe verify gate accepts SM90 as well as SM100, so
  servers that select nv_cutedsl (explicitly, or Kimi-K3 through its
  hook) run the CuTe kernel on Hopper.
- glm5_next_hook.py: extend the GLM-5-Next spec default
  --linear-attn-verify-backend nv_cutedsl (sgl-project#42905) to SM90 when
  --enable-linear-replayssm-spec is on; SM100 is unchanged.
- Tests: the ReplaySSM ring parity test also runs on the 1-gpu-large
  runner (H100); a new CuTe-vs-Triton verify parity test; CPU tests for
  the backend gate (SM90 and SM100 admitted, SM80 and SM120 rejected)
  and for the GLM-5-Next hook (SM90 with and without the ring).

Co-authored-by: Han Han <hanhan.hank@bytedance.com>
JackChuang added a commit to JackChuang/sglang that referenced this pull request Oct 22, 2026
@JackChuang @HanHan009527
[Perf] [NVIDIA] KDA: Extend the CuTe DSL MTP verify kernel to SM90
b13f513 
The fused CuTe DSL KDA verify kernel for MTP / speculative decoding
(fused_kda_decode_mtp_dspark) used paired FP32 instructions
(mul_packed_f32x2 / fma_packed_f32x2) that NVVM only accepts for SM100
and newer, and the KDA backend only ran it on SM100, so Hopper always
fell back to the unfused Triton verify path.

- kda_decode_mtp.py: a USE_PACKED_FMA constexpr (capability major >= 10)
  keeps the packed path on SM100 and newer and uses scalar FP32 FMAs on
  SM90, as cutedsl_gdn_mtp_ring.py does. SM100 compiles the same code as
  before.
- kda_backend.py: the CuTe verify gate accepts SM90 as well as SM100, so
  servers that select nv_cutedsl (explicitly, or Kimi-K3 through its
  hook) run the CuTe kernel on Hopper.
- glm5_next_hook.py: extend the GLM-5-Next spec default
  --linear-attn-verify-backend nv_cutedsl (sgl-project#42905) to SM90 when
  --enable-linear-replayssm-spec is on; SM100 is unchanged.
- Tests: the ReplaySSM ring parity test also runs on the 1-gpu-large
  runner (H100); a new CuTe-vs-Triton verify parity test; CPU tests for
  the backend gate (SM90 and SM100 admitted, SM80 and SM120 rejected)
  and for the GLM-5-Next hook (SM90 with and without the ring).

Co-authored-by: Han Han <hanhan7630@outlook.com>
JackChuang added a commit to JackChuang/sglang that referenced this pull request Oct 22, 2026
@JackChuang @HanHan009527
[Perf] [NVIDIA] KDA: Extend the CuTe DSL MTP verify kernel to SM90
f82665d 
The fused CuTe DSL KDA verify kernel for MTP / speculative decoding
(fused_kda_decode_mtp_dspark) used paired FP32 instructions
(mul_packed_f32x2 / fma_packed_f32x2) that NVVM only accepts for SM100
and newer, and the KDA backend only ran it on SM100, so Hopper always
fell back to the unfused Triton verify path.

- kda_decode_mtp.py: a USE_PACKED_FMA constexpr (capability major >= 10)
  keeps the packed path on SM100 and newer and uses scalar FP32 FMAs on
  SM90, as cutedsl_gdn_mtp_ring.py does. SM100 compiles the same code as
  before.
- kda_backend.py: the CuTe verify gate accepts SM90 as well as SM100, so
  servers that select nv_cutedsl (explicitly, or Kimi-K3 through its
  hook) run the CuTe kernel on Hopper.
- glm5_next_hook.py: extend the GLM-5-Next spec default
  --linear-attn-verify-backend nv_cutedsl (sgl-project#42905) to SM90 when
  --enable-linear-replayssm-spec is on; SM100 is unchanged.
- Tests: the ReplaySSM ring parity test also runs on the 1-gpu-large
  runner (H100); a new CuTe-vs-Triton verify parity test; CPU tests for
  the backend gate (SM90 and SM100 admitted, SM80 and SM120 rejected)
  and for the GLM-5-Next hook (SM90 with and without the ring).

Co-authored-by: Han Han <hanhan7630@outlook.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@kpham-sgl kpham-sgl kpham-sgl approved these changes

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

2 participants

@mmangkad @kpham-sgl