Temporarily Remove GLM-5.3 Flash decode CP support - #37487
Merged
kpham-sgl merged 2 commits intoSep 3, 2026
Merged
Conversation
kpham-sgl
force-pushed
the
codex/remove-glm53-decode-cp
branch
from
September 2, 2026 05:54
cf3fae4 to
821335d
Compare
kpham-sgl
marked this pull request as ready for review
September 2, 2026 21:17
kpham-sgl
requested review from
BBuf,
DarkSharpness,
Edwardf0t1,
Fridge003,
HaiShaw,
HydraQYH,
Qiaolin-Yu,
Ying1123,
alphabetc1,
celve,
ch-wan,
hanming-lu,
hebiao064,
hnyls2002,
huangtingwei9988,
hzh0425,
ispobock,
merrymercy,
xiezhq-hermann,
yizhang2077 and
yuan-luo
as code owners
September 2, 2026 21:17
Collaborator
Author
|
/rerun-test test/registered/models_e2e/test_glm53_flash_b200.py |
Contributor
|
Results for 🚀 |
Fridge003
approved these changes
Sep 2, 2026
Follow-up to #36507, companion to #37484 (prefill CP). The DSA backend has the decode-CP merge half but not the partition half: nothing under attention/dsa/ calls get_dcp_lens or update_local_kv_lens_for_dcp, unlike every other DCP-capable MLA backend, and no CI case exercises DSA x DCP. Take the merge half back out of the day-0 PR and re-land the axis whole in a PR based on main. Removed: - dsa_backend: _should_return_dsa_dcp_lse and its three call sites (tilelang decode, tilelang verify, trtllm decode), plus the _forward_tilelang return_lse branch that squeezed LSE for dcp_pack_a2a_send. - tilelang_kernel: the return_lse factory flag on sparse_attention_fwd_kernel_v1/v2 and tilelang_sparse_fwd, the LSE prim_func argument, and the unconditional per-call LSE allocation the always-present signature forced on the non-DCP decode path. The file is now byte-identical to its pre-f040cc72e6 state; the NoPE (tail_dim == 0) kernel selection and the topk % 64 relaxation are untouched. - dcp/comm: the Optional prefix_k_pe accommodation for NoPE models. - eager_runner: the target-verify skip in front of prepare_context_parallel_metadata_for_dcp. That gate was model-agnostic and also changed DeepSeek-V2, Kimi-K2.5, Kimi-K3 and Kimi-Linear under DCP + spec; main's behavior is restored for them. - glm5_next: prepare_context_parallel_metadata_for_dcp, the hook a model defines to opt into DCP. - test_dsa_trtllm_fp8_cp: the two _should_return_dsa_dcp_lse cases. The prefill-CP cases for _should_all_gather_dsa_trtllm_fp8_kv stay, matching what #37484 keeps. dcp/comm.py and eager_runner.py are now byte-identical to main. Main's DSA backend never returned LSE, so DSA x DCP is exactly as unsupported as it is on main -- no regression for GLM-5.2, DeepSeek-V3.2 or any other DSA model. Kept deliberately: the `not get_parallel().dcp_enabled` clause in _ingraph_verify_metadata_eligible (a defensive exclusion on a new feature, not DCP support), and the DCP memory-sizing fixes in pool_configurator / kv_cache_configurator, which are model-agnostic and no-ops at attn_dcp_size == 1. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Second half of the decode-CP removal: the sizing code that exists only because DCP exists. None of it names a GLM model, but every line of it is a no-op at attn_dcp_size == 1 and is reachable only when DCP is on, so it is part of the feature, not of the day-0 model support. - pool_configurator: the draft-pool budget multipliers in DefaultPoolConfigurator (draft_kv_size, draft_indexer_size and the non-DSA draft_num_layers) and in HybridSWAPoolConfigurator (_draft_swa_layers_num, _draft_swa_full_layers_num, _draft_full_layers_num), all keyed on kvc.ps.attn_dcp_size because draft pools are DCP-replicated rather than sharded. - kv_cache_configurator: _dsa_pool_geometry, which held the DSA physical page at 64 while growing the token pool by pool_page_size - physical_page_size. It differs from main only when loc_space_scale > 1, i.e. a draft worker under DCP; _build_dsa_kv_pool now passes page_size=self.pool_page_size like every other pool build site. - the tests that pinned those two behaviors: the dcp_size sweep and test_hybrid_swa_scales_draft_budget_with_dcp in test_pool_configurator, and test_dsa_draft_pool_preserves_physical_page_and_backs_virtual_tail in test_dcp_layout_unit. test_dcp_layout_unit.py is now byte-identical to main. test_pool_configurator.py differs from main only by the two non-DCP fixture lines #36507 needs (hf_config.model_type, is_draft_model), which are retained. After this commit the only DCP-touching lines left versus main are exclusion guards on new features, which stay by design: the `not get_parallel().dcp_enabled` clause in _ingraph_verify_metadata_eligible, and the two `attn_dcp_metadata is None` guards that keep the new host-side indptr fast path away from DCP batches in flashinfer_mla_backend. Those keep new code out of DCP configurations rather than implementing DCP. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
kpham-sgl
force-pushed
the
codex/remove-glm53-decode-cp
branch
from
September 2, 2026 22:39
821335d to
43a1922
Compare
Collaborator
Author
|
/rerun-test test/registered/attention/unittests/dsa/test_dsa.py test/registered/models_e2e/test_glm53_flash_b200.py test/registered/dcp/test_kimi_linear_dcp_dspark4.py test/registered/dcp/test_kimi_linear_dcp4.py |
Collaborator
Author
|
/rerun-test registered/attention/unittests/dsa/test_dsa.py |
Contributor
|
Results for 🚀 🚀 |
Contributor
|
Results for 🚀 🚀 |
Collaborator
Author
|
/rerun-test test/registered/models_e2e/test_glm53_flash_b200.py test/registered/dcp/test_kimi_linear_dcp_dspark4.py test/registered/dcp/test_kimi_linear_dcp4.py |
Contributor
|
Results for 🚀 |
Collaborator
Author
|
/rerun-test test/registered/models_e2e/test_glm53_flash_b200.py test/registered/dcp/test_kimi_linear_dcp_dspark4.py test/registered/dcp/test_kimi_linear_dcp4.py |
Contributor
|
Results for 🚀 |
Collaborator
Author
|
Should be safe, no regression to old DCP codepaths, non-dcp glm53 and non-dcp glm53 flash |
kpham-sgl
merged commit Sep 3, 2026
491cb22
into
xinyuan/glm-5.3-flash-support
105 of 117 checks passed
Leoyzen
added a commit
to Leoyzen/sglang
that referenced
this pull request
Sep 7, 2026
…5.3-Flash (SM90) Port of sgl-project#36904 (beastllama, closed unmerged) re-based onto main, plus the tilelang return_lse removed by sgl-project#37487's GLM-5.3-Flash decode-CP removal, plus DCP integration for glm-dcp-0907. - tilelang_kernel: dispatch sparse_mla_fwd_decode_partial_fp8 on CUDA (block_I=32/threads=128 tiles sized for the fp8 smem class, SM count via _cuda_sm_count); restore return_lse on sparse_attention_fwd_kernel_v1/v2 (caller-allocated base-2 LSE kernel arg, written only when traced) and extend sparse_mla_fwd_decode_combine with a global base-2 LSE output so the partial+combine fp8 path can also return LSE. - overrides (_check_tilelang_dsa_fp8_kv): allow CUDA + tilelang + fp8_e4m3 (SM89+ required, both DSA backends must be tilelang); DCP allowed since the raw-fp8 writer now applies the DCP owner filter. - kv_cache_configurator (calculate_mla_kv_cache_dim): route the raw 512B/token MLA pool layout on CUDA when both DSA backends are tilelang, keyed on the backend pair instead of the platform. - memory_pool (_write_mla_kv_buffer): key the raw fused-quant write branch on layout (not _is_hip); integrate with our DCP-aware write dispatch — DCP-widened locs go through set_mla_kv_buffer_dcp_sharded_triton_fp8_quant. - mla_buffer: add DCP_RANK/DCP_WORLD_SIZE owner filter to set_mla_kv_buffer_fp8_quant_kernel (mirrors the bf16 kernel) and a set_mla_kv_buffer_dcp_sharded_triton_fp8_quant wrapper. This closes the gap that made sgl-project#36904 refuse tilelang-fp8 + DCP. - dsa_backend: tilelang/tilelang added to the DCP backend-pair whitelist (SM80+, bf16 or raw-fp8 KV; SM89+ belt-and-braces check for fp8); _forward_tilelang gains return_lse with the DCP contract (base-2 LSE, [T,H,D]/[T,H] fp32, fixup_zero_kv_rows for owner-less rows — same contract as flashmla_kv/trtllm); supports_mha_one_shot disabled for the CUDA raw-fp8 tilelang layout (sgl-project#36904). - test: port of sgl-project#36904's CUDA raw-fp8 tilelang probe (one-hot exactness + spread-case fp8 budget with scrambled-index negative control). - test(hicache): register the pre-existing kv_cache_dtype storage-key unit test with CI and fix its bare pytest.main (pre-commit requirement). Provenance: internal deployment branch glm-dcp-0907; refs sgl-project#36904, sgl-project#37487 (removal reverted for tilelang), sgl-project#36989 (c-KV localization pattern).
5 tasks
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Follow-up to #36507, companion to #37484. Decode CP re-lands in a PR based on main after the day-0 PR merges.
Why. DSA × DCP here is the merge half without the partition half — nothing under
attention/dsa/callsget_dcp_lensorupdate_local_kv_lens_for_dcp, unlike every other DCP-capable MLA backend, and no CI case covers the combination.Commit 1 — data path
_should_return_dsa_dcp_lseand its three call sites (tilelang decode/verify, trtllm decode), plus the_forward_tilelangLSE branchreturn_lseonsparse_attention_fwd_kernel_v1/v2andtilelang_sparse_fwd, and the per-call LSE allocation it forced on the non-DCP pathOptionalprefix_k_peNoPE accommodation inall_gather_kv_cache_for_dcpeager_runnertarget-verify skip — model-agnostic, so this restores main for DeepSeek-V2, Kimi-K2.5/K3 and Kimi-LinearGlm5NextForConditionalGeneration.prepare_context_parallel_metadata_for_dcp, the DCP opt-in hook_should_return_dsa_dcp_lsetestsCommit 2 — DCP-only pool sizing (no-ops at
attn_dcp_size == 1, reachable only under DCP)× attn_dcp_sizemultipliers inDefaultPoolConfiguratorandHybridSWAPoolConfiguratorkv_cache_configurator._dsa_pool_geometry, and the tests pinning bothDFLASH/DSPARK are unaffected: both blocks sit under
is_eagle() or is_standalone(), and dflash's DCP factor lives in_dflash_draft_cell_size/scale_kv_cell_size_per_token_for_dflash, untouched.Kept. Prefill CP, KPool/DSA core, NoPE MLA, generic DCP for models that already had it, and the GLM-5.3 B200 e2e. Plus the DCP exclusion guards, which keep new code out of DCP rather than implement it:
not dcp_enabledin_ingraph_verify_metadata_eligible, and the twoattn_dcp_metadata is Noneguards inflashinfer_mla_backend.No regression vs main. Every removed line is new in #36507 —
return_lse,_should_return_dsa_dcp_lseand_dsa_pool_geometryare all absent fromorigin/main, so DSA × DCP stays exactly as unsupported as it is there.dcp/comm.py,eager_runner.pyandtest_dcp_layout_unit.pyland byte-identical to main, andtilelang_kernel.pytof040cc72e6^, leaving NoPE (tail_dim == 0) kernel selection and thetopk % 64relaxation untouched. Models defining the DCP hook go 5 → 4, dropping onlyglm5_next.py.return_lsewas a trace-time flag, so DCP-off kernels never contained the LSE epilogue: output is bit-identical, minus one allocation per tilelang decode call.Validation. compileall and pre-commit clean; no dangling reference to any removed symbol. Rebased onto
545bd6f839after #37484 merged. No GPU run yet — hence draft.CI States
Latest PR Test (Base): ❌ Run #33691459781
Latest PR Test (Extra): ❌ Run #33691459691
Latest PR Test (AMD ROCm 7.2): ❌ Run #33691459760