[Kernel] GLM 5.3 Flash related kernels (ported from #36507) by kpham-sgl · Pull Request #37477 · sgl-project/sglang

kpham-sgl · 2026-09-01T22:13:04Z

Ported from #36507.

No behavior change on main: new files with no callers, plus additive parameters that default to existing behavior.

CI States

Latest PR Test (Base): 🚫 Run #33571827221
Latest PR Test (Extra): ✅ Run #33571826901
Latest PR Test (AMD ROCm 7.2): ⏳ Run #33571827220

Fridge003 · 2026-09-01T22:25:16Z

    return q_out, k_nope_out, k_rope_out


+def mla_quantize_for_fp8_no_rope(


Looks like duplication of mla_quantize_without_rope_for_fp8,
needs to be cleaned later

Fridge003 · 2026-09-01T22:26:30Z

        tl.extra.cuda.gdc_launch_dependents()


+@triton.jit


Looks like duplication of _set_mla_kv_buffer_impl, might be cleaned later

Fridge003 · 2026-09-01T22:26:58Z

    )


+@triton.jit


Looks like duplication of get_mla_kv_buffer_triton, might clean later

Fridge003 · 2026-09-01T22:27:54Z

+    return mhc_post(x, residual, post_layer_mix, comb_res_mix)
+
+
+def hc_pre(


hc_pre / hc_post looks like duplication of mhc_pre/mhc_post, might be cleaned later

Ported from #36507. No behavior change on main: new files with no callers, plus additive params defaulting to existing behavior. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

kpham-sgl · 2026-09-01T23:42:38Z

/rerun-test test_kda_*.py test_mhc_kernels.py test_hc_combine.py test_set_mla_kv_buffer.py test_kimi_k3_prerequisite_ops.py test_dsa_kpool_multi_pool.py test_dsa.py test_unified_mla_gpu_parity.py test_deepseek_v4_amd_fused_mhc.py test_gdn_noncontiguous_stride.py registered/attention/unittests/kda/test_triton.py test_kimi_linear_models.py test_kimi_k3_b300.py test_dsa_glm52_tp_mtp.py test_deepseek_v4_flash_fp8_h200.py

github-actions · 2026-09-01T23:43:01Z

Results for /rerun-test test_kda_*.py test_mhc_kernels.py test_hc_combine.py test_set_mla_kv_buffer.py test_kimi_k3_prerequisite_ops.py test_dsa_kpool_multi_pool.py test_dsa.py test_unified_mla_gpu_parity.py test_deepseek_v4_amd_fused_mhc.py test_gdn_noncontiguous_stride.py registered/attention/unittests/kda/test_triton.py test_kimi_linear_models.py test_kimi_k3_b300.py test_dsa_glm52_tp_mtp.py test_deepseek_v4_flash_fp8_h200.py:

🚀 4-gpu-b200 (6 tests): ✅ View workflow run

cd test/ && python3 registered/attention/test_kda_decode_flashinfer.py
cd test/ && python3 registered/attention/test_kda_prefill_cutedsl.py
cd test/ && python3 registered/kernels/ops/attention/test_kda_prefill.py
cd test/ && python3 registered/kernels/test_kda_mtp_cutedsl_replayssm_ring.py
cd test/ && python3 registered/attention/unittests/dsa/test_dsa.py
cd test/ && python3 registered/attention/unittests/kda/test_triton.py

🚀 1-gpu-h100 (17 tests): ❌ View workflow run

cd test/ && python3 registered/attention/test_kda_kernels.py
cd test/ && python3 registered/attention/test_kda_prefill_flashkda.py
cd test/ && python3 registered/attention/unittests/hybrid_linear/test_kda_fused_accept_indices.py
cd test/ && python3 registered/kernels/ops/attention/test_kda_fused_decode.py
cd test/ && python3 registered/kernels/ops/attention/test_kda_helion.py
cd test/ && python3 registered/kernels/test_kda_replayssm_fold.py
cd test/ && python3 registered/kernels/test_kda_replayssm_fold_batched.py
cd test/ && python3 registered/kernels/test_kda_replayssm_ring_fused.py
cd test/ && python3 registered/kernels/test_kda_replayssm_ring_ragged.py
cd test/ && python3 registered/kernels/ops/layernorm/test_mhc_kernels.py
cd test/ && python3 registered/kernels/ops/layernorm/test_hc_combine.py
cd test/ && python3 registered/kernels/ops/kvcache/test_set_mla_kv_buffer.py
cd test/ && python3 registered/kernels/ops/test_kimi_k3_prerequisite_ops.py
cd test/ && python3 registered/kernels/test_dsa_kpool_multi_pool.py
cd test/ && python3 registered/attention/unittests/dsa/test_dsa.py
cd test/ && python3 registered/attention/test_gdn_noncontiguous_stride.py
cd test/ && python3 registered/attention/unittests/kda/test_triton.py

🚀 4-gpu-gb300 (1 test): ⏳ View workflow run

cd test/ && python3 registered/kernels/ops/attention/test_kda_prefill.py

🚀 ubuntu-latest (4 tests): ✅ View workflow run

cd test/ && python3 registered/unit/layers/attention/linear/kernels/test_kda_nvidia.py
cd test/ && python3 registered/unit/layers/attention/test_kda_helion_dispatcher.py
cd test/ && python3 registered/unit/mem_cache/test_kda_fused_decode_strided_state.py
cd test/ && python3 registered/unit/models/test_deepseek_v4_amd_fused_mhc.py

🚀 1-gpu-5090 (2 tests): ✅ View workflow run

cd test/ && python3 registered/unit/layers/test_kda_decode_mtp_slot_stride.py
cd test/ && python3 registered/unit/mem_cache/test_unified_mla_gpu_parity.py

🚀 2-gpu-h100 (1 test): ✅ View workflow run

cd test/ && python3 registered/models_e2e/test_kimi_linear_models.py

🚀 8-gpu-b300 (1 test): ✅ View workflow run

cd test/ && python3 registered/models_e2e/test_kimi_k3_b300.py

🚀 8-gpu-h200 (2 tests): ✅ View workflow run

cd test/ && python3 registered/models_e2e/test_dsa_glm52_tp_mtp.py
cd test/ && python3 registered/models_e2e/test_deepseek_v4_flash_fp8_h200.py

kpham-sgl · 2026-09-02T01:16:16Z

/rerun-test test_kda_kernels.py test_kda_prefill_cutedsl.py test_kda_fused_accept_indices.py test_kda_fused_decode.py test_kda_helion.py test_kda_prefill.py test_kda_replayssm_fold.py test_kda_replayssm_fold_batched.py test_kda_replayssm_ring_fused.py test_kda_replayssm_ring_ragged.py test_kda_mtp_cutedsl_replayssm_ring.py test_kda_nvidia.py test_kda_helion_dispatcher.py test_kda_decode_mtp_slot_stride.py test_kda_fused_decode_strided_state.py registered/attention/unittests/kda/test_triton.py test_mhc_kernels.py test_hc_combine.py test_set_mla_kv_buffer.py test_kimi_k3_prerequisite_ops.py test_dsa_kpool_multi_pool.py test_dsa.py test_unified_mla_gpu_parity.py test_deepseek_v4_amd_fused_mhc.py test_gdn_noncontiguous_stride.py test_kimi_linear_models.py test_kimi_k3_b300.py test_dsa_glm52_tp_mtp.py test_deepseek_v4_flash_fp8_h200.py

github-actions · 2026-09-02T01:16:40Z

Results for /rerun-test test_kda_kernels.py test_kda_prefill_cutedsl.py test_kda_fused_accept_indices.py test_kda_fused_decode.py test_kda_helion.py test_kda_prefill.py test_kda_replayssm_fold.py test_kda_replayssm_fold_batched.py test_kda_replayssm_ring_fused.py test_kda_replayssm_ring_ragged.py test_kda_mtp_cutedsl_replayssm_ring.py test_kda_nvidia.py test_kda_helion_dispatcher.py test_kda_decode_mtp_slot_stride.py test_kda_fused_decode_strided_state.py registered/attention/unittests/kda/test_triton.py test_mhc_kernels.py test_hc_combine.py test_set_mla_kv_buffer.py test_kimi_k3_prerequisite_ops.py test_dsa_kpool_multi_pool.py test_dsa.py test_unified_mla_gpu_parity.py test_deepseek_v4_amd_fused_mhc.py test_gdn_noncontiguous_stride.py test_kimi_linear_models.py test_kimi_k3_b300.py test_dsa_glm52_tp_mtp.py test_deepseek_v4_flash_fp8_h200.py:

🚀 1-gpu-h100 (16 tests): ✅ View workflow run

cd test/ && python3 registered/attention/test_kda_kernels.py
cd test/ && python3 registered/attention/unittests/hybrid_linear/test_kda_fused_accept_indices.py
cd test/ && python3 registered/kernels/ops/attention/test_kda_fused_decode.py
cd test/ && python3 registered/kernels/ops/attention/test_kda_helion.py
cd test/ && python3 registered/kernels/test_kda_replayssm_fold.py
cd test/ && python3 registered/kernels/test_kda_replayssm_fold_batched.py
cd test/ && python3 registered/kernels/test_kda_replayssm_ring_fused.py
cd test/ && python3 registered/kernels/test_kda_replayssm_ring_ragged.py
cd test/ && python3 registered/attention/unittests/kda/test_triton.py
cd test/ && python3 registered/kernels/ops/layernorm/test_mhc_kernels.py
cd test/ && python3 registered/kernels/ops/layernorm/test_hc_combine.py
cd test/ && python3 registered/kernels/ops/kvcache/test_set_mla_kv_buffer.py
cd test/ && python3 registered/kernels/ops/test_kimi_k3_prerequisite_ops.py
cd test/ && python3 registered/kernels/test_dsa_kpool_multi_pool.py
cd test/ && python3 registered/attention/unittests/dsa/test_dsa.py
cd test/ && python3 registered/attention/test_gdn_noncontiguous_stride.py

🚀 4-gpu-b200 (5 tests): ✅ View workflow run

cd test/ && python3 registered/attention/test_kda_prefill_cutedsl.py
cd test/ && python3 registered/kernels/ops/attention/test_kda_prefill.py
cd test/ && python3 registered/kernels/test_kda_mtp_cutedsl_replayssm_ring.py
cd test/ && python3 registered/attention/unittests/kda/test_triton.py
cd test/ && python3 registered/attention/unittests/dsa/test_dsa.py

🚀 4-gpu-gb300 (1 test): ⏳ View workflow run

cd test/ && python3 registered/kernels/ops/attention/test_kda_prefill.py

🚀 ubuntu-latest (4 tests): ✅ View workflow run

cd test/ && python3 registered/unit/layers/attention/linear/kernels/test_kda_nvidia.py
cd test/ && python3 registered/unit/layers/attention/test_kda_helion_dispatcher.py
cd test/ && python3 registered/unit/mem_cache/test_kda_fused_decode_strided_state.py
cd test/ && python3 registered/unit/models/test_deepseek_v4_amd_fused_mhc.py

🚀 1-gpu-5090 (2 tests): ✅ View workflow run

cd test/ && python3 registered/unit/layers/test_kda_decode_mtp_slot_stride.py
cd test/ && python3 registered/unit/mem_cache/test_unified_mla_gpu_parity.py

🚀 2-gpu-h100 (1 test): ✅ View workflow run

cd test/ && python3 registered/models_e2e/test_kimi_linear_models.py

🚀 8-gpu-b300 (1 test): ✅ View workflow run

cd test/ && python3 registered/models_e2e/test_kimi_k3_b300.py

🚀 8-gpu-h200 (2 tests): ✅ View workflow run

cd test/ && python3 registered/models_e2e/test_dsa_glm52_tp_mtp.py
cd test/ && python3 registered/models_e2e/test_deepseek_v4_flash_fp8_h200.py

kpham-sgl · 2026-09-02T05:17:01Z

All the potentially affected tests have passed. Safe to merge

sgl-project#37477) Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>

kpham-sgl requested review from 1am9trash, BBuf, DarkSharpness, Fridge003, HaiShaw, HydraQYH, Qiaolin-Yu, YAMY1234, celve, hebiao064, hubertlu-tw, ispobock, kkHuang-amd, merrymercy, rainj-me, yizhang2077 and yuan-luo as code owners September 1, 2026 22:13

github-actions Bot added the jit-kernel label Sep 1, 2026

Fridge003 reviewed Sep 1, 2026

View reviewed changes

Fridge003 approved these changes Sep 1, 2026

View reviewed changes

kpham-sgl changed the title ~~[Kernel] KPool FP8 index kernels, pooled top-k transform, and mHC kernel ops (ported from #36507)~~ [Kernel] GLM 5.3 Flash related kernels (ported from #36507) Sep 1, 2026

kpham-sgl added run-ci run-ci-extra bypass-fastfail high priority labels Sep 1, 2026

kpham-sgl force-pushed the kpham/port-36507-kpool-mhc-kernels branch from 1fa726d to 0a91e3f Compare September 1, 2026 22:53

kpham-sgl removed high priority bypass-fastfail labels Sep 1, 2026

Add KPool FP8 index kernels, pooled top-k transform, and mHC kernel ops

cf0f761

Ported from #36507. No behavior change on main: new files with no callers, plus additive params defaulting to existing behavior. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

kpham-sgl force-pushed the kpham/port-36507-kpool-mhc-kernels branch from 0a91e3f to cf0f761 Compare September 1, 2026 23:37

kpham-sgl merged commit c66a285 into main Sep 2, 2026
213 of 254 checks passed

kpham-sgl deleted the kpham/port-36507-kpool-mhc-kernels branch September 2, 2026 05:17

bold84 mentioned this pull request Sep 2, 2026

GLM-5.3-Flash support #36507

Merged

kpham-sgl mentioned this pull request Sep 2, 2026

Temporarily Remove GLM-5.3 Flash decode CP support #37487

Merged

bold84 mentioned this pull request Sep 2, 2026

v0.1.0-rc.72: DSA radix top-k correctness and crash fix (root-caused from the issue #3 coredump) ormandj/sglang-glm53-flash-sm120#4

Merged

StevenChenSE pushed a commit to StevenChenSE/sglang that referenced this pull request Sep 6, 2026

[Kernel] GLM 5.3 Flash related kernels (ported from sgl-project#36507) (

0752a39

sgl-project#37477) Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>

This was referenced Sep 14, 2026

[CI Monitor] Daily Report - 2026-09-14 bingxche/sglang-ci-bot#195

Open

[CI Monitor] Daily Report - 2026-09-15 bingxche/sglang-ci-bot#196

Open

hippothewild mentioned this pull request Sep 15, 2026

GLM-5.3-Flash W4AFP8 patch stack on upstream main 4df5df91 (2026-09-08) vessl-ai/sglang#71

Open

Navigation Menu

Uh oh!

[Kernel] GLM 5.3 Flash related kernels (ported from #36507) - #37477

[Kernel] GLM 5.3 Flash related kernels (ported from #36507)#37477
kpham-sgl merged 1 commit into
mainfrom
kpham/port-36507-kpool-mhc-kernels

kpham-sgl commented Sep 1, 2026 •

edited by github-actions Bot

Loading

Uh oh!

Fridge003 Sep 1, 2026

Uh oh!

Fridge003 Sep 1, 2026

Uh oh!

Fridge003 Sep 1, 2026

Uh oh!

Fridge003 Sep 1, 2026

Uh oh!

kpham-sgl commented Sep 1, 2026

Uh oh!

github-actions Bot commented Sep 1, 2026 •

edited

Loading

Uh oh!

kpham-sgl commented Sep 2, 2026

Uh oh!

github-actions Bot commented Sep 2, 2026 •

edited

Loading

Uh oh!

kpham-sgl commented Sep 2, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

		return q_out, k_nope_out, k_rope_out


		def mla_quantize_for_fp8_no_rope(

		return mhc_post(x, residual, post_layer_mix, comb_res_mix)


		def hc_pre(

		tl.extra.cuda.gdc_launch_dependents()


		@triton.jit

		)


		@triton.jit

Uh oh!

Conversation

kpham-sgl commented Sep 1, 2026 • edited by github-actions Bot Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

CI States

Uh oh!

Fridge003 Sep 1, 2026

Choose a reason for hiding this comment

Uh oh!

Fridge003 Sep 1, 2026

Choose a reason for hiding this comment

Uh oh!

Fridge003 Sep 1, 2026

Choose a reason for hiding this comment

Uh oh!

Fridge003 Sep 1, 2026

Choose a reason for hiding this comment

Uh oh!

kpham-sgl commented Sep 1, 2026

Uh oh!

github-actions Bot commented Sep 1, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

kpham-sgl commented Sep 2, 2026

Uh oh!

github-actions Bot commented Sep 2, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

kpham-sgl commented Sep 2, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

kpham-sgl commented Sep 1, 2026 •

edited by github-actions Bot

Loading

github-actions Bot commented Sep 1, 2026 •

edited

Loading

github-actions Bot commented Sep 2, 2026 •

edited

Loading