Repository navigation
[TRTLLM-17096][feat] modeling_v2: GLM-5.2 NVFP4 low-latency decode target for GB300 - #20524
Conversation
Add RuntimeHooks (pyexecutor/runtime_hooks.py), a per-model seam through which a modeling_v2 target can take over parts of the decode step without changing the engine's behaviour for other models: - step-input staging, so a target can write the next step's inputs from its own buffers; - the speculative sampler's store writer; - executor events and an execution-stream loop; - a decode prologue that runs before the model forward. The executor, CUDA graph runner, decoder runner, KV cache manager V2, attention backends and the speculative modules read the hooks at their call sites. Every hook is off unless a model installs it, and the default behaviour is unchanged. Signed-off-by: Xiao Wang <24860335+xwang233@users.noreply.github.com>
Add model-specific CUDA kernels under cpp/tensorrt_llm/kernels/glmMoeDsa/ and their torch ops under cpp/tensorrt_llm/thop/glmMoeDsa/, one <stem>Op.cpp per kernel <stem>.cu, registered as trtllm::glm_moe_dsa_<what>. They cover the decode path of GLM-5.2: the post-attention hop and its routing, the RoPE KV append, the sparse MLA decode, the cross-rank indexer top-k, the dense MLP, and the draft model's routed MoE, live-row hop and tail. The kernels build as one OBJECT library for the SM100 family and are used only by the GLM-5.2 modeling_v2 targets. Signed-off-by: Xiao Wang <24860335+xwang233@users.noreply.github.com>
Add the glm_moe_dsa family to modeling_v2 with three targets, each claiming its own parallel topology through the family's route(): - glm_5_2_nvfp4__sm_103__tp8: tensor-parallel attention and MoE on 8 GPUs; - glm_5_2_nvfp4__sm_103__tep: tensor-parallel attention with expert-parallel experts on 4 or 8 GPUs; - glm_5_2_nvfp4__sm_103__dep: attention data parallelism with expert-parallel experts on 8 or 16 GPUs. Code the targets share (kernel packages, dispatch, attention, the layer composition, the speculative worker and the runtime-hook wiring) lives in models/glm_moe_dsa/shared/. Each target selects the fast paths it serves through its own plan; the paths default on inside the target and can be turned off with TRTLLM_STAIRCASE_GLM_LL_DISABLE or per-path TRTLLM_STAIRCASE_GLM_LL_* variables. Any configuration that no target claims runs the built-in model. Signed-off-by: Xiao Wang <24860335+xwang233@users.noreply.github.com>
Unit tests of the glm_moe_dsa family: routing and target claims, isolation from other models, the per-target plans, the adopted layers and the runtime hooks, plus kernel tests against the native implementations. Tests that need the GLM-5.2 checkpoint find it under llm_models_root() and skip when it is not available. Signed-off-by: Xiao Wang <24860335+xwang233@users.noreply.github.com>
Add the glm_moe_dsa family and its three targets to the modeling_v2 README. Signed-off-by: Xiao Wang <24860335+xwang233@users.noreply.github.com>
|
Attention assignment inactive Checked 2026-10-29T17:09:06.492797+00:00. The changed files no longer require There is no outstanding attention team request. Individual review requests retained at this check:
Previous attention decision (inactive history)@PerkzZheng, could you help to review this PR as trt-llm-torch-attention-devs? Thanks~ <details> <summary>Assignment explanation</summary> <table> <tr><th>Components</th><td>kernels/sparse (50%); kernels/auxiliary (45%); integration/build-config (5%)</td></tr> <tr><th>Tags</th><td>algorithm:sparse (100%); specific-algorithm:dsa (100%); representation:mla (100%); phase:decode (100%); features:rope (40%); features:quantization (40%); features:speculative-decoding (20%); context:llm (100%)</td></tr> <tr><th>PR effort</th><td>13 points — Provisionally 13 for interacting sparse-attention kernels, early dependent launches and cross-rank transport contracts, independently of total file count. Classified as feature: the added weights.py documentation explicitly retains the built-in model's module tree and loader, supporting optimized targets for an existing architecture.</td></tr> <tr><th>Importance</th><td>Important — The visible contracts introduce clustered sparse-MLA softmax reduction, change RoPE/KV producer-consumer launch ordering, and fuse ordered cross-rank attention-output reduction. These require sensitive numerical, synchronization and performance review.</td></tr> <tr><th>Retained GitHub review requests</th><td>litaotju: preserve direct code-owner review request</td></tr> <tr><th>Affinity (0–1)</th><td>0.467 component/tag match; 0.350 after PIC preference</td></tr> <tr><th>Outstanding points at selection</th><td>57 → 70 (+13)</td></tr> <tr><th>Points already reserved on this PR</th><td>0</td></tr> <tr><th>Effective weight</th><td>1</td></tr> <tr><th>Completed review points (preceding 168 hours)</th><td>17</td></tr> <tr><th>Combined points after assignment</th><td>87</td></tr> <tr><th>Demand-derived target</th><td>81.00 points (soft share)</td></tr> <tr><th>Additional fair share before assignment</th><td>7.00 points</td></tr> <tr><th>Share explanation</th><td>Whole PR / qualification or coverage constraint exceeds a soft share; assignment is allowed.</td></tr> <tr><th>Selection rule</th><td>affinity with 0.3-share guard</td></tr> <tr><th>Assignment scope (PerkzZheng)</th><td>all</td></tr> </table> <p>Affinity and workload are snapshots at selection, not live counters. Existing reservations on this PR are counted once.</p> </details> |
…use the TP8 fast paths Each of the two targets gets a short README: the fast paths were designed and tuned for the TP8 target, these targets run the subset that applies at their topology, and both should be revisited with fast paths of their own. Signed-off-by: Xiao Wang <24860335+xwang233@users.noreply.github.com>
clang-format on five C++ files and one test .cu (whitespace only), an unused local and a blank line flagged by ruff, two CuTe tensor names that codespell reads as misspellings, and the prep packet's pinned slots allocated with prefer_pinned(), as the pinned-memory policy requires. The stager declines when prefer_pinned() is False, so a packet is only built when it is True. Signed-off-by: Xiao Wang <24860335+xwang233@users.noreply.github.com>
|
PR_Github #81990 [ run ] completed with state
|
|
/bot run --disable-fail-fast --extra-stage "GB300-8_GPUs-2_Nodes-PyTorch-SingleNvlinkDomain-Post-Merge-1" |
|
PR_Github #82023 [ run ] triggered by Bot. Commit: |
|
PR_Github #82023 [ run ] completed with state
|
The 2-D CuTe DSL RMSNormKernel, its compile helper and rmsnorm_cute are unchanged between FlashInfer 0.6.18 and 0.7.2, so draft_front still mirrors the native RMSNorm on 0.7.2. Signed-off-by: Xiao Wang <24860335+xwang233@users.noreply.github.com>
|
/bot run --disable-fail-fast --extra-stage "GB300-8_GPUs-2_Nodes-PyTorch-SingleNvlinkDomain-Post-Merge-1" |
|
PR_Github #82035 [ run ] triggered by Bot. Commit: |
brnguyen2
left a comment
There was a problem hiding this comment.
Ran a deep correctness review of this PR (glue layer plus a full read of the kernel/CUDA code, split by component). Posting the findings that survived verification against the code at head d0d2bf1; style/structure items were deliberately left out since those passes already happened.
Headline: the branch's spec_worker.py predates the _forward_draft_loop refactor that the last main merge brought in (eagle3.py:817 now requires accepted_tokens/pre_draft_seq_lens positionally and returns a (next_draft_tokens, next_new_tokens) tuple; the kimi_k3_vl spec workers were ported to the new contract, this one wasn't). As a result any speculative decode step at head raises TypeError before reaching the kernels — likely the root cause of the GB300 stage failures in today's pipeline runs. Details in the two inline comments on spec_worker.py; once fixed, the GSM8K gate timing in the test-db comment should be re-measured since it predates the skew.
Worth saying explicitly: the heavy kernel code came back clean. The cross-rank top-k Lamport/epoch protocols and determinism claims, the routed-MoE split-K ownership and NVFP4 scale indexing (including the stock 32-row shuffle decode), the sparse-MLA online-softmax split merge, the RoPE/KV-append kernels (checked line-for-line against mlaKernels.cu/dsv3RopeOp.cpp), and the draft-tail/draft-MoE exchange protocols all verified sound, as did the Op wrapper shape checks and mbarrier byte accounting throughout.
…p contract - GlmMtpWorker passes accepted_tokens and pre_draft_seq_lens to _forward_draft_loop and unpacks its (next_draft_tokens, next_new_tokens) result; the _forward_linear_draft_loop override takes the same arguments, forwards them to the base loop, and returns the same tuple. - Both return paths of _forward_impl set logits_vocab_shard when the target logits are this rank's vocabulary shard. - The tp8 route requires an fp8 latent KV cache, like the other GLM targets. - Slice the draft tokens to token_count only on the padded LM-head path. - Parse TRTLLM_ENABLE_PDL the same way in every kernel (on unless "0"). - Drop compile-time prints from the routed MoE kernel and the stale cross-rank agreement wording in draft_live_rows. Signed-off-by: Xiao Wang <24860335+xwang233@users.noreply.github.com>
|
/bot run --disable-fail-fast --extra-stage "GB300-8_GPUs-2_Nodes-PyTorch-SingleNvlinkDomain-Post-Merge-1" |
|
PR_Github #82091 [ run ] triggered by Bot. Commit: |
|
PR_Github #82035 [ run ] completed with state |
|
PR_Github #82091 [ run ] completed with state
|
|
/bot run --disable-fail-fast |
|
PR_Github #82189 [ run ] triggered by Bot. Commit: |
|
PR_Github #82189 [ run ] completed with state
|
|
/bot run --disable-fail-fast |
|
PR_Github #82227 [ run ] triggered by Bot. Commit: |
|
PR_Github #82227 [ run ] completed with state |
Dev Engineer Review
modeling_v2targets for SM103: TP8, TEP, and DEP.RuntimeHooks; default hooks remain disabled. Existing executor paths retain native behavior when hooks are absent or ineligible.QA Engineer Review
tests/integration/defs/accuracy/modeling_v2/test_glm_moe_dsa.pyfor TP8, TEP4, and DEP8. Tests are restricted to SM103 and require the applicable GPU and MPI topology.tests/integration/test_lists/test-db/l0_gb300_multi_gpus.ymltests/integration/test_lists/test-db/l0_gb300_multi_nodes_node2_gpu8.ymlPer-File QA Perspective
cpp/tensorrt_llm/CMakeLists.txt,cpp/tensorrt_llm/kernels/CMakeLists.txt,cpp/tensorrt_llm/kernels/glmMoeDsa/CMakeLists.txt,cpp/tensorrt_llm/thop/CMakeLists.txt, andcpp/tensorrt_llm/thop/glmMoeDsa/CMakeLists.txt: Verify target inclusion, architecture selection, symbol resolution, and Torch operator linkage.cpp/tensorrt_llm/kernels/glmMoeDsa/*: Verify SM103 support, validation errors, deterministic reductions, cross-rank synchronization, timeout behavior, PDL ordering, and invalid-input handling.cpp/tensorrt_llm/thop/glmMoeDsa/*: Verify operator schemas, tensor validation, workspace requirements, stream behavior, fallback errors, and output shapes.tensorrt_llm/_torch/_experimental/modeling_v2/models/glm_moe_dsa/*: Verify routing, TP8/TEP/DEP target selection, weight loading, fast-path eligibility, native fallback, dispatch installation, speculative decoding, and rank-uniform decisions.tensorrt_llm/_torch/_experimental/modeling_v2/models/glm_moe_dsa/shared/*: Verify attention, top-k, KV append, sparse MLA, MoE, deferred normalization, draft-loop, workspace, health-check, and host-page-retention contracts.tensorrt_llm/_torch/pyexecutor/runtime_hooks.py: Verify that all hooks remain disabled by default and that invalid hook declarations fail clearly.tensorrt_llm/_torch/pyexecutor/engine/runners/decoder/runner.py,py_executor.py,cuda_graph_runner.py,kv_cache/kv_cache_manager_v2.py, and speculative-decoding files: Verify native behavior when hooks are absent, staging cleanup after exceptions, CUDA-graph replay, stream synchronization, and sampler-store integration.tensorrt_llm/_torch/attention/backends/*: Verify staged metadata updates, FMHA library precedence, reused MTP rows, and unchanged native paths.tensorrt_llm/_torch/_experimental/modeling_v2/README.md,_router_index.py, and package initializers: Verify documented target names, shared-package import rules, and resolved-mode behavior.tests/integration/defs/accuracy/modeling_v2/test_glm_moe_dsa.py: Covers whole-model GSM8K accuracy for TP8, TEP4, and DEP8. It is listed in the GB300 CI test lists.tests/unittest/_torch/executor/test_runtime_hooks.py: Covers runtime-hook protocols, default-off isolation, staged/native paths, and executor lifecycle integration. No integration test-list entry is expected for this unit test.tests/unittest/_torch/attention/sparse/dsa/test_dsa_indexer.py: Covers reused top-k rows, decode filling, and MTP shared-row storage. No integration test-list entry is expected.tests/unittest/_torch/attention/test_fmha_manager.py: Covers prepended FMHA library ordering and unavailable-library fallback. No integration test-list entry is expected.tests/unittest/_torch/executor/test_gpu_keepalive.pyandtests/unittest/_torch/executor/test_py_executor.py: Update test stubs for new executor-hook fields. No integration test-list entry is expected.tests/unittest/_torch/modeling_v2/test_modeling_v2_claims.py: Covers family-shared import rules. No integration test-list entry is expected.tests/integration/test_lists/test-db/*.yml: Adds GB300 CI entries for the GLM GSM8K gates with extended timeouts. Verify that the selected topology matches each test’s required rank count.Description
This PR adds a GLM-5.2 NVFP4 low-latency decode target for GB300 (SM103) to
modeling_v2:models/glm_moe_dsa/glm_5_2_nvfp4__sm_103__tp8(ModelingV2Glm52Nvfp4Sm103Tp8).The optimization work targets one layout: attention and MoE tensor-parallel over 8 GPUs (TP8),
no attention DP, with one-model MTP-Eagle as a linear chain of five draft tokens. The family's
route()sends only that deployment to this target. Other GLM-5.2 deployments go to the otherglm_moe_dsatargets or to the built-in model:glm_5_2_nvfp4__sm_103__tep4from [TRTLLM-16980][feat] modeling_v2: GLM-5.2-NVFP4 sm_103 tep4 target #20662. An earlierrevision of this PR also served TEP and DEP with a subset of its TP8 fast paths; it no longer
does, so the two targets do not overlap.
glm_5_2_nvfp4__sm_107__tp16from [None][feat] Staircase: GLM-5.2 NVFP4 low-latency decode target for Rubin (modeling_v2, TP16) #20517.routing.pykeeps main's TP16 and TEP4 branches and adds the TP8 one. The NoTarget messagenames all three targets.
Layout
The target follows main's target-ownership rules (
test_modeling_v2_claims.py): everything itruns lives in its own directory. Next to
modeling.pyandweights.pyare:dispatch.py);attention.py,attention_forward.py);layers.py);spec_worker.py);runtime.py);kernels/.The fast paths default on inside the target.
TRTLLM_STAIRCASE_GLM_LL_DISABLE=1turns them alloff, and per-path
TRTLLM_STAIRCASE_GLM_LL_*variables turn off single paths.Executor:
RuntimeHooks(reviewed in #20797)The target takes over parts of the decode step through
RuntimeHooks, a per-model seam thatis off by default:
The target also has a step-input stager and a sampler store writer
(
kernels/deferred_input_prep). They are not served: #19989 onmainreworked the engine pathsthey hooked, so #20797 does not carry those two hooks. The code stays for a follow-up that ports
them onto the new native paths.
The target's side is
runtime.py. The engine side, its tests and the note about #20306 are in#20797.
Kernels and ops
The target's CUDA kernels and their torch ops live in its own
csrc/, under the per-familynative build of #20896: one
<stem>Op.cppper<stem>.cu, declared inmodels/glm_moe_dsa/csrc/CMakeLists.txtand built intolibth_modeling_v2_glm_moe_dsa.so.They register in
torch.ops.mv2_glm_moe_dsa_glm_5_2_nvfp4_sm_103_tp8; the target loads them asnative_opswhen routing picks it, andfake_ops.pyfakes every op. Nothing undercpp/changes.They cover:
Test Coverage
tests/integration/defs/accuracy/modeling_v2/test_glm_moe_dsa.py::TestModelingV2Glm52Nvfp4Sm103Tp8::test_gsm8k_tp8.modeling_v2="require", so a deployment that misses the target, or astep a fast path covers but cannot run, fails instead of silently falling back to the
built-in model.
l0_gb300_multi_nodes_node2_gpu8.the latest merges of main; CI re-runs it on the current head.
tests/unittest/_torch/executor/test_runtime_hooks.pyand the touchedexecutor and attention tests cover the hooks both off (the default) and installed.
[−0.30, +0.99].
PR Checklist
llm_argschange.GitHub Bot Help
To see a list of available CI bot commands, please comment
/bot help.