[None][feat] MiniMax-M3 NVFP4 low-latency decode megakernel for modeling_v2 for Blackwell by peihu-nv · Pull Request #20306 · NVIDIA/TensorRT-LLM

peihu-nv · 2026-10-14T17:23:16Z

Description

At small decode batches, a speculative MiniMax-M3 NVFP4 step on GB300 at TP4/TP8 is bound by many small kernels, the boundaries between them, and host work between CUDA graph replays. This PR adds two modeling_v2 targets for that configuration, minimax_m3_nvfp4__sm_103__tp4 and minimax_m3_nvfp4__sm_103__tp8, plus the executor seams they need. It has two main parts: the TP4 target with the runtime hooks, and the TP8 target. The router picks the target for the matching checkpoint and hardware, modeling_v2: require makes it mandatory, and every kernel it adds is a catalog entry with a torch reference and unit tests.

Step-loop runtime hooks

The hooks are RuntimeHooks in tensorrt_llm/_torch/pyexecutor/runtime_hooks.py; a model declares them through runtime_hooks_of(model). Every hook defaults to off and no in-tree model declares one, so existing models are unchanged. There are no llm_args changes.

Packed step inputs (step_inputs.py, with the trtllm::step_inputs_scatter op): copies made between replays are staged and issued as one upload plus one scatter. The scatter applies host copies first, then device copies, both in recording order.
Skipped unchanged uploads (h2d_cache.py): metadata copies whose values the destination already holds are skipped.
Both are owned per engine and are active only inside that engine's forward. On a failed step, the pack is discarded and the skip cache is cleared.
Private metadata pool: an engine that skips unchanged uploads gives its CUDA graph metadata a buffer pool of its own. Otherwise another engine could overwrite a buffer this one still treats as unchanged.
Executor and sampler:
- the overlap loop can run on the execution stream;
- autotuned tactics can be merged across ranks;
- the greedy RNG upload can be skipped;
- a speculative step-fold protocol lets the sampler read its rows from the model's buffers;
- the sampler takes greedy tokens straight from the model when it produced them.
MSA policy:
- an identity block table for short sequences;
- a hook for the KV update;
- skipping of unchanged block offsets;
- reusable cache-pool views.
- The short/long CUDA-graph threshold is now resolved from the model.
Small op additions: mxfp8_quantize_autotuned, flashinfer_mxfp8_gemm_cute_dsl, and a split_k argument on the CuTe DSL index scorer.

Greedy-only outputs (greedy tokens without logits) are kept only when nothing downstream reads the logits. An executor-side guided decoder turns them off.

The TP4 decode target

A pure generation step of up to 48 tokens runs a fused decode path. Prefill, mixed steps and wider steps take the generic layer path.

Attention: megakernels with long-context block selection and the KV update; a pre-wait-staged NVFP4 QKV/indexer producer; an index scorer that prefetches older index-K pages before its grid wait; long-context sparse decode staged before its grid wait; the dense decode's KV split sized for one request.
Boundaries: multicast all-reduce + residual + RMSNorm, with the routed experts' finalize folded in (batched top-k loads). The Eagle3 capture layers stay on the fused boundary.
GEMMs: tcgen05 MXFP8 GEMVs at low row counts, and stock's autotuned MXFP8 GEMM elsewhere with tactics agreed across ranks.
MoE: CuTe DSL engines for 1–13 rows, with rank-agreed prebuilds.
Speculation: a vocab-parallel greedy verify, an exact sampled verify, an early drafter fc fork, and a fused draft loop. The draft loop's boundaries keep stock Eagle3's structure.
Heads: the target and draft heads stream a lossless 12-bit coding of their BF16 weights. Each coding is checked bit for bit against the weights before use; if the check fails, the BF16 path runs.

Explicit dispatch. dispatch.py lists the row range each fast path serves. Every covered row count is a certified cell of the entry that serves it. Under modeling_v2: require, a covered shape whose precondition is missing raises; it never falls back silently. Every item is on by default and has one constant that turns it off. The target reads no environment variables.

Kernels. The kernels are compiled ops (cpp/tensorrt_llm/kernels/minimaxM3, trtllm::minimax_m3_*) and CuTe DSL modules. Nothing compiles at runtime from source strings, and every compile happens before warmup.

Numerics. Weights and activations keep stock M3's precisions; the drafter and both LM heads run in BF16 over the full vocabulary. Results differ from stock only by floating-point summation order, so at most a few BF16 units in the last place. The greedy verify reproduces stock's acceptance bit for bit, and the MoE-finalize adds the shared expert in stock's order.

One known difference: the long-context block selector breaks score ties toward the smaller block, as its contract says. The in-tree selector instead takes tied blocks in atomic arrival order when more than 1024 blocks tie exactly.

The one change outside the target and the executor seams is a small hook in fmha/msa_decode.py. It lets a layer cap the dense decode's KV-split length, and it is a no-op unless a layer sets the cap.

The TP8 decode target

minimax_m3_nvfp4__sm_103__tp8 serves the same checkpoint on 8 GB300s across two trays (attention and experts both tensor-parallel 8). It runs the TP4 target's exact decode path at the TP8 geometry, plus:

MoE: router, routed experts and shared expert in one CuTe DSL launch for 1–8 tokens, with the sum pushed into the boundary all-reduce.
Attention: an 8-head long-context sparse decode (minimaxM3Tp8/sparseDecodeLc8), and index scoring fused with block selection for up to 2 requests (minimaxM3Tp8/fusedScoreSelect).
Drafter and heads: fixed-tactic low-M BF16 GEMMs, and BF16 target and draft heads with no coded copies.

The two targets are isolated. The TP8 target imports nothing from the TP4 target, owns copies of the MoE engine and MX GEMV it changes, and a drift test pins the files it copies unchanged. Outside the target, the TP8 part touches only routing.py (the tp8 branch), the new minimaxM3Tp8 kernels and their thop ops, and a two-line no-op hook in fmha/flashinfer_trtllm_gen.py that applies a layer's KV-split cap.

Test Coverage

Unit tests under tests/unittest/_torch/modeling_v2/models/minimax_m3/minimax_m3_nvfp4__sm_103__tp4/: a torch-reference test for every catalog entry; dispatch tables vs certified cells; the require-mode raise; and exact-numerics tests of the verify and drafter against stock Eagle3. All 1156 pass on GB300.
Runtime hook tests under tests/unittest/_torch/ (executor/, sampler/, speculative/, attention/sparse/msa/): all 548 pass on GB300.
Accuracy vs stock (same build, modeling_v2 off):
- GSM8K (1319 questions), McNemar test: greedy 97.35% vs 97.04% (p = 0.52); sampled at temperature 1.0, top-p 0.95, 96.82% vs 96.97% (p = 0.80);
- Eagle3 natural acceptance length over 264 requests: 2.561 ± 0.029 vs 2.536 ± 0.029.
TP8 unit tests: 1019 pass on GB300. 2 are skipped because they need a world size of 8 across two trays; those comm cells were run at 2 and 4 ranks instead.
TP8 accuracy vs stock:
- GSM8K greedy 97.12% vs 97.27% (p = 0.82); sampled 96.74% vs 97.04% (p = 0.57);
- natural acceptance length 2.536 vs 2.553.
End to end on GB300: about 1.55x interactivity for TP4 and ~1.84x for TP8 over the baseline configuration.

PR Checklist

Please review the following before submitting your PR:

PR description clearly explains what and why. If using CodeRabbit's summary, please make sure it makes sense.
PR Follows TRT-LLM CODING GUIDELINES to the best of your knowledge.
Test cases are provided for new code paths (see test instructions)
If PR introduces API changes, an appropriate PR label is added - either api-compatible or api-breaking. For api-breaking, include BREAKING in the PR title.
Any new dependencies have been scanned for license and vulnerabilities
CODEOWNERS updated if ownership changes
Documentation updated as needed
Update tava architecture diagram if there is a significant design change in PR.
The reviewers assigned automatically/manually are appropriate for the PR.
Please check this after reviewing the above items as appropriate for this PR.

GitHub Bot Help

To see a list of available CI bot commands, please comment /bot help.

🤖 Generated with Claude Code

peihu-nv · 2026-10-14T17:23:31Z

/bot run

tensorrt-cicd · 2026-10-14T17:33:28Z

PR_Github #77887 [ run ] triggered by Bot. Commit: 044b375 Link to invocation

tensorrt-cicd · 2026-10-14T18:00:21Z

PR_Github #77887 [ run ] completed with state FAILURE. Commit: 044b375
/LLM/main/L0_MergeRequest_PR pipeline #64205 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

peihu-nv · 2026-10-14T19:50:17Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-10-14T20:20:07Z

PR_Github #77904 [ run ] triggered by Bot. Commit: 157a99a Link to invocation

tensorrt-cicd · 2026-10-15T02:23:23Z

PR_Github #77904 [ run ] completed with state FAILURE. Commit: 157a99a
/LLM/main/L0_MergeRequest_PR pipeline #64221 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

peihu-nv · 2026-10-15T03:00:57Z

/bot run --disable-fail-fast

… step-loop runtime hooks Adds the modeling_v2 target minimax_m3_nvfp4__sm_103__tp4 (GB300, TP4, Eagle3 one-model speculation) and the executor seams it opts into. - Runtime hooks (off unless a model declares them): packed step inputs, skipped unchanged H2D uploads, a private graph metadata pool, sampler greedy-token and step-fold paths, MSA runtime policy, and rank-merged autotuned tactics. Greedy-only outputs are kept only when nothing reads the logits (no guided decoder, non-empty batch). - TP4 decode target: fused attention with long-context selection, a pre-wait-staged NVFP4 QKV/indexer producer and index scorer, multicast all-reduce + RMSNorm boundaries with the MoE finalize folded in, MXFP8 GEMVs, CuTe DSL MoE engines, an exact greedy and sampled verify, a fused draft loop, and lossless 12-bit coded BF16 target and draft heads. - Explicit dispatch: every covered shape is a certified catalog cell; under modeling_v2: require a covered shape that cannot run raises. Every item is on by default with one constant that turns it off; no environment reads. - Numerics match stock up to summation order; the long-context selector breaks >1024-way score ties by its contract (smaller block). Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

Pick the row source explicitly, raise when neither logits nor greedy tokens are present, and assert logits before the fast greedy kernel. Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

Adds the modeling_v2 target minimax_m3_nvfp4__sm_103__tp8 (GB300, 8 ranks over two trays, tensor-parallel attention and experts, one-model Eagle3 or no speculation) and routes world = tp = moe_tp = 8 to it; the tp4 branch of the router is unchanged. - Decode path: the tp4 target's exact decode path at the TP8 geometry, plus the TP8 items: the 1..8-token MoE engine at expert width 384 (logits replicated 32x at <= 4 tokens), an 8-head one-launch long-context sparse decode staged before its grid wait at <= 8 rows, a fused index score + top-16 select for spans of <= 2 requests, the producer releasing its dependents on entry, the long-row select's early release at <= 8 rows, the TP8 MXFP8 GEMV table, fixed-tactic low-M BF16 GEMMs for the drafter linears, the drafter decode's KV split capped at 16K for one request, BF16 target and draft heads, and the drafter fc forked beside the verify head. - Isolation: the target imports nothing from the tp4 target; the MoE engine and the MX GEMV op it changes are its own copies inside the target. The two new kernels are compiled ops (cpp/tensorrt_llm/kernels/minimaxM3Tp8, trtllm::minimax_m3_tp8_*). A drift test pins the files copied unchanged from the tp4 target. - fmha/flashinfer_trtllm_gen.py applies a layer's KV-split cap (msa_decode.dense_decode_max_seq_len); a no-op unless a layer sets one. - Same contract as the tp4 target: every kernel a catalog entry with CELLS and a torch reference, dispatch.py listing each fast path's rows, a covered shape that cannot run raising under modeling_v2: require, kernels built before warmup, every item on by default with one constant to turn it off. - Numerics match stock up to summation order (BF16 heads and drafter, NVFP4 experts, MXFP8 projections). Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

CUtensorMap already carries its own alignment, which is stricter than 64 bytes in newer CUDA toolkits; clang-hosted nvcc rejects the weaker alignas with "alignment cannot be set to less than the default alignment" (SBSA build). Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

peihu-nv · 2026-10-15T03:18:21Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-10-15T03:26:33Z

PR_Github #77987 [ run ] triggered by Bot. Commit: bb8674c Link to invocation

tensorrt-cicd · 2026-10-15T08:38:53Z

PR_Github #77987 [ run ] completed with state FAILURE. Commit: bb8674c
/LLM/main/L0_MergeRequest_PR pipeline #64288 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

Resolve two additive conflicts with the Qwen4-exp modeling_v2 target by keeping both router entries and both routing tests. Port the sampler's greedy_tokens_usable to main's request_row_params after the TorchSampler strategy refactor (NVIDIA#20188) removed _request_strategy. Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

- Register fake implementations for the new step_inputs_scatter and MiniMax-M3 router, low-M GEMM, producer and allreduce-norm ops. - Give the warmup tests' stand-in engines the default runtime hooks the decoder runner now reads. - Leave the identity table unset in the piecewise-attention metadata mock. - Stub the eager K/V view resolution only while the runtime-policy fixture builds its manager, so the pool-view tests still exercise get_buffers. Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

peihu-nv · 2026-10-15T17:53:19Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-10-15T18:00:09Z

PR_Github #78201 [ run ] triggered by Bot. Commit: a427f56 Link to invocation

tensorrt-cicd · 2026-10-15T22:13:39Z

PR_Github #78201 [ run ] completed with state SUCCESS. Commit: a427f56
/LLM/main/L0_MergeRequest_PR pipeline #64456 completed with status: 'UNSTABLE'

CI Report

⚠️ Multi-GPU Label Required:
Multi-GPU tests require the ci: full pre-merge approved label on this PR. Either:

Wait for the PR to be fully approved — the label is added automatically once approval is complete. Having unresolved open comments is fine, or
If needed, ask a member of NVIDIA/trt-llm-ci-approvers to add the label manually.
Then re-trigger CI with the same bot command (no rebase needed).

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

Link to invocation

Restructure the minimax_m3_nvfp4__sm_103__tp4 and __tp8 targets into the flat layout the modeling_v2 exemplars use, with no change to what they compute. Each target owns its own copy of every file (the targets share no code): * modeling.py: the dispatch tables (FAST_PATHS, row ranges, miss/uncovered policy), the flat core, PrefillTarget / DecodeTarget, the one-model Eagle3 drafter and worker, the step fold, runtime hooks and the shell. The core selects the target on num_contexts alone: every pure generation step is DecodeTarget's (its decode kernels up to 48 tokens, its own generic layers for wider steps or embeddings), every step with context rows PrefillTarget's. The two targets share no body: each owns its layer, attention and MoE compositions, and its step_args is the one projection of the MSA metadata its forward reads (DecodeTarget states num_contexts=0 / num_ctx_tokens=0). Shared module-level code is limited to single-call op adapters that read no step metadata. * weights.py: the language-model and drafter weight tables (host-side TP slicing, consumed checkpoint keys released during the load; at TP8 the 4 KV / index heads are duplicated on rank pairs as the built-in loader does). * ops/: catalog entries with certified cells for every call the model makes; kernel wrappers route through their entries. The MSA and drafter attention entries carry a numeric reference and GPU tests that drive every cell against real caches and prepared metadata. * kernels/: the CuTe DSL GEMV, MoE engine and embedding+first-norm kernels move into each target (the old shared cute_dsl_kernels/minimax_m3 copy is removed); capacity annotations on every constant. * The legacy dispatch/layers/linears/runtime/spec_fold/spec_worker modules are removed; the unit tests are ported to the flat targets. * The unfused QKV / index projection runtime (fuse_qkv_index_projection=False, FP8 indexer, FP8 or NVFP4 main KV) runs as the built-in model does. Routing (models/minimax_m3/routing.py) now declines what the targets do not build or run, so modeling_v2="auto" falls back to the built-in model and "require" names the missed criterion: * Eagle3 drafters other than the targets' one-layer drafter (read from the drafter checkpoint's config.json). * Runtimes other than MSA with an FP8 indexer and an FP8 or NVFP4 main KV cache. The resolver hands routing modules that set ROUTES_ON_RUNTIME a RuntimeFacts (sparse attention config and the KV cache quantization the engine will run with); only MiniMax-M3 sets it. _validate_runtime stays as the backstop. Adds a whole-model accuracy test under modeling_v2="require" (GSM8K without speculation, and with one-model Eagle3 plus a chat-GSM8K acceptance probe), on the 4-GPU GB300 list (TP4) and the 2-node GB300 list (TP8). Bitwise identical to the previous targets on GB300 (tokens and per-step outputs on every rank, no speculation and Eagle3, fused and unfused projections, TP4 and TP8, including pure generation steps wider than the decode kernels' 48 tokens). Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

peihu-nv · 2026-10-19T00:47:51Z

/bot run --disable-fail-fast

Resolve cache_manager.py: keep the decode shortcuts and pooled views, and the HND main KV storage with the strided NHD view from NVIDIA#20159. Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

peihu-nv · 2026-10-19T03:45:45Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-10-19T03:54:00Z

PR_Github #78786 [ run ] triggered by Bot. Commit: 6d9af78 Link to invocation

tensorrt-cicd · 2026-10-19T11:19:16Z

PR_Github #78786 [ run ] completed with state SUCCESS. Commit: 6d9af78
/LLM/main/L0_MergeRequest_PR pipeline #64914 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

Conflicts: - sampler.py: the stable greedy path keeps the forward's greedy tokens (rows_cuda) and takes main's removal of the two-model d2t check and the d2t argument of fast_greedy_sample_kernel. Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

peihu-nv · 2026-10-22T19:53:50Z

/bot run --disable-fail-fast

Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

peihu-nv · 2026-10-26T22:10:44Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-10-26T22:17:20Z

PR_Github #80708 [ run ] triggered by Bot. Commit: eea8a02 Link to invocation

tensorrt-cicd · 2026-10-27T03:21:38Z

PR_Github #80708 [ run ] completed with state SUCCESS. Commit: eea8a02
/LLM/main/L0_MergeRequest_PR pipeline #66456 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

Measured on GB300 with NVIDIA#20306's CI wheel (its MiniMax-M3 TP4 modeling_v2 target) against the built-in model in the same wheel, the gen worker's median device step time per decode iteration was 6.07 vs 10.44 ms at concurrency 1 and 47.48 vs 45.56 ms at concurrency 64. So the con1 lane sets `modeling_v2: auto`, and the con64 lane pins `off`. Main has no MiniMax-M3 route, so `auto` runs the built-in model today (same step time as with the option unset) and the lane picks the target up once it is on main. Signed-off-by: chenfeiz0326 <203214996+chenfeiz0326@users.noreply.github.com>

Mgluhovskoi

LGTM from a telemetry POV. Nit: the PR description still says “There are no llm_args changes,” which is now inaccurate.

github-actions · 2026-10-27T16:42:48Z

Automatically added "ci: full pre-merge approved" because this PR has satisfied the required GitHub review approvals. Unresolved review conversations and other required checks remain independent merge requirements.

Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

…sh executor tests Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

peihu-nv · 2026-10-27T20:15:27Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-10-27T20:21:52Z

PR_Github #80966 [ run ] triggered by Bot. Commit: d1a23f1 Link to invocation

Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

peihu-nv · 2026-10-27T22:26:02Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-10-27T22:33:43Z

PR_Github #80996 [ run ] triggered by Bot. Commit: bf11325 Link to invocation

tensorrt-cicd · 2026-10-27T22:37:13Z

PR_Github #80966 [ run ] completed with state ABORTED. Commit: d1a23f1

Link to invocation

tensorrt-cicd · 2026-10-28T03:33:20Z

PR_Github #80996 [ run ] completed with state FAILURE. Commit: bf11325
/LLM/main/L0_MergeRequest_PR pipeline #66716 completed with status: 'FAILURE'

CI Report

⚠️ Action Required:

Please check the failed tests and fix your PR
If you cannot view the failures, ask the CI triggerer to share details
Once fixed, request an NVIDIA team member to trigger CI again

CI Agent Failure Analysis

Link to invocation

github-actions Bot assigned peihu-nv Oct 14, 2026

peihu-nv added 4 commits October 14, 2026 20:10

[None][fix] Satisfy mypy on the MiniMax-M3 sampler fast-greedy path

c76faa4

Pick the row source explicitly, raise when neither logits nor greedy tokens are present, and assert logits before the fast greedy kernel. Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

peihu-nv force-pushed the peihengh/m3-megakernel branch from 306e8da to bb8674c Compare October 15, 2026 03:17

peihu-nv added 2 commits October 15, 2026 10:27

xwang233 mentioned this pull request Oct 18, 2026

[TRTLLM-17096][feat] modeling_v2: GLM-5.2 NVFP4 low-latency decode target for GB300 #20524

Open

6 tasks done

[None][chore] Merge main into MiniMax-M3 megakernel PR

6d9af78

Resolve cache_manager.py: keep the decode shortcuts and pooled views, and the HND main KV storage with the strided NHD view from NVIDIA#20159. Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

litaotju reviewed Oct 19, 2026

View reviewed changes

litaotju mentioned this pull request Oct 19, 2026

[TRTLLM-17033][feat] Qwen3.8-Flash-Next modeling_v2: CUTEDSL, MTP 1..7, video, FP8/NVFP4 KV, bf16 state, canonical QSA selection, breakable prefill graph, GB300 fast paths #20276

Merged

1 task done

peihu-nv requested a review from BowenFu October 22, 2026 10:24

Wanli-Jiang mentioned this pull request Oct 23, 2026

[None][feat] modeling_v2: build target C++/CUDA as per-family libraries #20896

Open

1 task done

WeiHaocheng approved these changes Oct 25, 2026

View reviewed changes

WeiHaocheng mentioned this pull request Oct 26, 2026

[None][refactor] Route per-step device copies through a model-declared copy plan #20972

Draft

1 task

[None][chore] Merge main 3205aea into the MiniMax-M3 megakernel PR

eea8a02

Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

Shixiaowei02 reviewed Oct 27, 2026

View reviewed changes

Comment thread ...erimental/modeling_v2/models/minimax_m3/minimax_m3_nvfp4__sm_103__tp4/ops/m3_mc_allreduce.py Outdated

litaotju approved these changes Oct 27, 2026

View reviewed changes

liji-nv approved these changes Oct 27, 2026

View reviewed changes

chenfeiz0326 mentioned this pull request Oct 27, 2026

[None][test] Add DeepSeek-V4.1-Flash, GLM-5.3-Flash, Qwen3.8-Flash-Next, Kimi-K3 and MiniMax-M3 perf-sanity cases #21080

Merged

1 task done

Mgluhovskoi approved these changes Oct 27, 2026

View reviewed changes

trtllm-agent added the ci: full pre-merge approved label Oct 27, 2026

peihu-nv added 3 commits October 27, 2026 13:07

[None][chore] Merge main into the MiniMax-M3 megakernel PR

294d255

Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

[None][fix] Use TP process groups for MiniMax-M3 workspaces and refre…

6e4caea

…sh executor tests Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

[None][fix] Keep MiniMax-M3 modeling_v2 opt-in like other models

d1a23f1

Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

[None][fix] Merge main preserving M3 page-table materialization

bf11325

Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

peihu-nv enabled auto-merge (squash) October 27, 2026 22:37

BowenFu approved these changes Oct 28, 2026

View reviewed changes

Conversation

peihu-nv commented Oct 14, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Description

Step-loop runtime hooks

The TP4 decode target

The TP8 decode target

Test Coverage

PR Checklist

GitHub Bot Help

Uh oh!

peihu-nv commented Oct 14, 2026

Uh oh!

tensorrt-cicd commented Oct 14, 2026

Uh oh!

tensorrt-cicd commented Oct 14, 2026

Uh oh!

peihu-nv commented Oct 14, 2026

Uh oh!

tensorrt-cicd commented Oct 14, 2026

Uh oh!

tensorrt-cicd commented Oct 15, 2026

Uh oh!

peihu-nv commented Oct 15, 2026

Uh oh!

peihu-nv commented Oct 15, 2026

Uh oh!

tensorrt-cicd commented Oct 15, 2026

Uh oh!

tensorrt-cicd commented Oct 15, 2026

Uh oh!

peihu-nv commented Oct 15, 2026

Uh oh!

tensorrt-cicd commented Oct 15, 2026

Uh oh!

tensorrt-cicd commented Oct 15, 2026

Uh oh!

peihu-nv commented Oct 19, 2026

Uh oh!

peihu-nv commented Oct 19, 2026

Uh oh!

tensorrt-cicd commented Oct 19, 2026

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

tensorrt-cicd commented Oct 19, 2026

Uh oh!

peihu-nv commented Oct 22, 2026

Uh oh!

peihu-nv commented Oct 26, 2026

Uh oh!

tensorrt-cicd commented Oct 26, 2026

Uh oh!

tensorrt-cicd commented Oct 27, 2026

Uh oh!

Uh oh!

Mgluhovskoi left a comment • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

github-actions Bot commented Oct 27, 2026

Uh oh!

peihu-nv commented Oct 27, 2026

Uh oh!

tensorrt-cicd commented Oct 27, 2026

Uh oh!

peihu-nv commented Oct 27, 2026

Uh oh!

tensorrt-cicd commented Oct 27, 2026

Uh oh!

tensorrt-cicd commented Oct 27, 2026

Uh oh!

tensorrt-cicd commented Oct 28, 2026

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

peihu-nv commented Oct 14, 2026 •

edited

Loading

Mgluhovskoi left a comment •

edited

Loading