[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

litaotju · 2026-10-19T09:55:27Z

+
+} // namespace
+
+void invokeStepInputsScatter(uint8_t const* stage, cudaStream_t stream)


shall these 2 also be put into minimaxM3 folder?

Done in dcdbd8c: stepInputsScatter.{cu,h} are now in kernels/minimaxM3/.

litaotju · 2026-10-19T09:56:19Z

+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0


I am wondering shall we put these m3 specific torch ops into thop/minimaxM3?

Easy to manage the review rules and changed based testing.

Done in dcdbd8c: all MiniMax-M3 torch ops, including stepInputsScatterOp.cpp, are now in thop/minimaxM3/, which adds them to th_common_attention from its own CMakeLists.txt (like thop/moe/).

litaotju · 2026-10-19T09:59:43Z



+@dataclass(frozen=True)
+class RuntimeFacts:


@WeiHaocheng could you help to take a look at this one?

We can directly use the 'def route(pretrained_config, mapping, compute_capability, spec_config, kv_cache_quant_algo)' function interface here.

litaotju · 2026-10-19T10:02:05Z

@@ -418,6 +418,60 @@ def _(
        return scores.new_empty((scores.shape[2], scores.shape[0], topk),
                                dtype=torch.int32)

+    @torch.library.register_fake("trtllm::step_inputs_scatter")


I think we should register these under the modeling_v2 m3 speficif folder since its only used there.

Done in dcdbd8c: the fake impls moved to modeling_v2/models/minimax_m3/fake_ops.py, imported by the M3 package, pyexecutor/step_inputs.py and the fake-registration test.

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

tensorrt-cicd · 2026-10-20T03:15:51Z

PR_Github #79048 [ run ] completed with state FAILURE. Commit: 87b9d8d
/LLM/main/L0_MergeRequest_PR pipeline #65117 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

…ares is_spec_decode DecoderRunner reads DecoderRunnerConfig.is_spec_decode at construction (the spec step fold binds the engine buffers). The compiled-body validation test builds the runner from a SimpleNamespace config; give it the field the real config always has. Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

peihu-nv · 2026-10-20T03:21:29Z

/bot run --disable-fail-fast

Resolve the modeling_v2 router against NVIDIA#20267: route() now takes the resolver's KV-cache algorithm as its fifth argument for every routing module, and modules that set ROUTES_ON_RUNTIME (MiniMax-M3) still receive runtime=RuntimeFacts; the MiniMax-M3 route() accepts the new argument. Keep both sides' new helpers (RuntimeFacts / resolved_mode and tp_reaches_across_nodes) and both test-list additions. Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

peihu-nv · 2026-10-20T03:29:51Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-10-20T03:36:31Z

PR_Github #79124 [ run ] triggered by Bot. Commit: 5a9c484 Link to invocation

tensorrt-cicd · 2026-10-20T09:12:46Z

PR_Github #79124 [ run ] completed with state FAILURE. Commit: 5a9c484
/LLM/main/L0_MergeRequest_PR pipeline #65180 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

WeiHaocheng · 2026-10-20T09:56:47Z

For the new Minimax M3 kernels, please move them to the Minimax directory under modeling_v2 (you can add a csrc/ folder for the C++/CUDA sources). Shared kernels for TP4/TP8 belong in models/minimax_m3, while target-specific kernels should reside in their respective target folders. In short, new ops should be scoped to the level where they take effect rather than dumped into generic modules.

Resolve the MiniMax-M3 index-score kernels against NVIDIA#19275 (c146094), which split cute_dsl_kernels/ into the domain kernel trees: the TP4 and TP8 m3_index_score imports now take cute_ptx_utils and utils from tensorrt_llm._torch.kernels.blackwell, and the docstring points at attention/kernels/blackwell/minimax_m3_index_decode_score.py. Both moved modules are unchanged (pure renames). Signed-off-by: peihengh <259410613+peihu-nv@users.noreply.github.com>

peihu-nv · 2026-10-20T12:47:09Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-10-20T12:54:18Z

PR_Github #79261 [ run ] triggered by Bot. Commit: 2343e84 Link to invocation

WeiHaocheng · 2026-10-20T13:56:00Z

+        alpha: Optional[float],
+        beta: Optional[float],
+    ) -> torch.Tensor:
+        return torch.ops.trtllm.silu_and_mul(


For model-specific ops implemented in C++/CUDA, a simple pybind is sufficient; they don't necessarily need to be bound as 'torch.ops.trtllm.x'.

WeiHaocheng · 2026-10-20T14:00:30Z

+"""Fixed M3 H6144 gather/raw residual/RMS, reusing pinned K3 norm kernel.
+
+128 threads x 8 BF16 x 6 iterations covers the row. Unlike the original
+K3 wrapper this target does not promise FlashInfer's bit-identical reduction.


Please move the code under 'tensorrt_llm/_torch/_experimental/modeling_v2/models/minimax_m3/minimax_m3_nvfp4__sm_103__tp4/kernels/' to the 'ops' directory at the same level.

WeiHaocheng · 2026-10-20T14:05:51Z

+    captured = getattr(spec_config, "eagle3_layers_to_capture", None)
+    checks = (
+        (
+            cfg.get("num_hidden_layers") == 1,


Adding so many model architecture checks in the routing logic is unnecessary. The fact that this interface is invoked already guarantees that the model architecture is Minimax M3. These checks are essentially guarding against scenarios that will never happen.

WeiHaocheng · 2026-10-20T14:06:44Z

  checkpoint's config, the parallel mapping, the GPU's compute capability and
-  the speculative-decoding config. It returns the target class, or raises
+  the speculative-decoding config (and, for a module that sets
+  ``ROUTES_ON_RUNTIME``, the engine arguments in :class:`RuntimeFacts`). It


This line is also can be deleted.

tensorrt-cicd · 2026-10-20T17:34:42Z

PR_Github #79261 [ run ] completed with state SUCCESS. Commit: 2343e84
/LLM/main/L0_MergeRequest_PR pipeline #65288 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-20T20:05:35Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-10-20T20:12:36Z

PR_Github #79375 [ run ] triggered by Bot. Commit: 2343e84 Link to invocation

tensorrt-cicd · 2026-10-20T22:53:22Z

PR_Github #79375 [ run ] completed with state SUCCESS. Commit: 2343e84
/LLM/main/L0_MergeRequest_PR pipeline #65374 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-21T11:15:49Z

/bot run --disable-fail-fast

tensorrt-cicd · 2026-10-21T11:24:02Z

PR_Github #79618 [ run ] triggered by Bot. Commit: 2343e84 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

Open

1 task done

peihu-nv requested review from SimengLiu-nv, bo-nv, chuangz0, eopXD, heyuhhh, joyang-nv, lfr-0531, longlee0622 and rosong11 October 19, 2026 23:04

WeiHaocheng reviewed Oct 20, 2026

View reviewed changes


		} // namespace

		void invokeStepInputsScatter(uint8_t const* stage, cudaStream_t stream)



		@dataclass(frozen=True)
		class RuntimeFacts:

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!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

tensorrt-cicd commented Oct 19, 2026

Uh oh!

tensorrt-cicd commented Oct 20, 2026

Uh oh!

peihu-nv commented Oct 20, 2026

Uh oh!

peihu-nv commented Oct 20, 2026

Uh oh!

tensorrt-cicd commented Oct 20, 2026

Uh oh!

tensorrt-cicd commented Oct 20, 2026

Uh oh!

WeiHaocheng commented Oct 20, 2026

Uh oh!

peihu-nv commented Oct 20, 2026

Uh oh!

tensorrt-cicd commented Oct 20, 2026

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

peihu-nv commented Oct 14, 2026 •

edited

Loading