[Wanli-Jiang]

Description

Completes the modeling_v2 target qwen4_exp/qwen3_8_flash_next_nvfp4__sm_103__tp1 (GB300 sm_103, TP1), on top of #20143 (merged), rebased on main af3e5eb257:

• it closes the target's feature gaps relative to the built-in Qwen3.8-Flash-Next model (TRTLLM-17033);
• it adds the target's GB300 fast paths (TRTLLM-16460).

Every new feature is off unless selected through engine configuration.

The commits are grouped by area for review; the commits between the feature head (77f394148c) and 59b0c08fe8 add ops that 59b0c08fe8 wires in, the next six address the review comments and a CI failure, 2cef6d6cc4 follows main's move of the CuTe DSL helpers (#19275), 53bdd2281e compiles the fused block top-k once for every row count, the next three address the second round of review comments, the next five cut the target's tests down (26 accuracy cases to 6; the modeling_v2 unit run from 3,444 to 1,213 passed cases), and the last one keeps the admission tests passing next to main's KV cache retention check (#20086).

commit	content
109b161cd4 feat	engine, QSA backend and modeling_v2 catalog support (shared code, listed below)
5f4a2ac156 feat	the target's features (table below)
77f394148c test	accuracy gates and CI test lists
24c1ab6201 perf	shared pieces: • catalog entries for a low-M bf16 GEMM and FlashInfer's chunked gated delta rule • PDL in the QSA index-score kernel • QSA launch heuristics
d284f4a871 perf	attention ops
fef084ee3a perf	Gated-DeltaNet ops and their gated norm
c499e28702 perf	hyper-connection ops
eb0808f923 perf	MoE, GEMM, PLE and elementwise ops
59b0c08fe8 perf	modeling.py / weights.py run the fast paths
abd5b89871 chore	cleanup: unreachable paths and duplicated helpers removed, including ops and one catalog entry the forward no longer calls
b1de32022f test	CI cost: four qwen4_exp accuracy gates pre-merge, nine post-merge, the rest on demand
1141bb4781 fix	CI: the beam-search admission test's executor stub gets the model engine the model hook reads
64c1e28466 fix	MTP verify steps run the composed attention prologue (the fused one assumes one row per request)
981b09a526 fix	the gated norm's multi-row prefill launch checks that z covers x
f19ae0bcae chore	routing: each family's route gets the ModelConfig and LLM args next to main's kv_cache_quant_algo; qwen4_exp derives its own engine choices
26c282e864 chore	low_m_bf16_mm and flashinfer_chunk_gated_delta_rule move from catalog/ to the target's ops/
0186621e62 test	review fixes in the target's tests
2cef6d6cc4 chore	the target's CuTe DSL ops import make_ptr and the pipeline helpers from tensorrt_llm/_torch/kernels/blackwell/, where #19275 moved them
53bdd2281e perf	qsa_topk_expand_decode compiles one kernel for every row count: the row count is a runtime argument, so an eager step's context rows no longer pay a JIT compile for each new count
456f198a1e fix	thop attention: FP8 context FMHA for a non-MLA NVFP4 KV cache only on the SM100 family, the only FP8 context FMHA kernels that read NVFP4 KV
9aa0985bd7 perf	the second expert copy is kept only when max_num_tokens can reach a 4096-row launch; TRTLLM_MODELING_V2_MOE_PREFILL_COPY=0 drops it
dcffd603d2 fix	QSA cache manager: the side-plane packing no longer depends on the KV heads a rank holds, so TP1, TP2 and attention-DP peers register the same transfer roles
cf90243a00, 18f2de7752, ed32fa5db5 test	accuracy: 26 cases to 6; each engine runs every check that fits in it (below)
375b5f90e4, 19521ac16d test	unit: each op certifies the shapes and layouts its forward reaches in at most three cells (four where it reaches four compiled variants), with one property test and one single-case guard; wrong-variant, composition and determinism tests are dropped, and each op's note lists what stays certified and what does not
13c21f460b fix	the model hook is looked up through getattr(self, "model_engine", None), as _kv_connector_moves_kv reads its connector, so #20086's object.__new__(PyExecutor) executor stub has no hook to call; this PR's admission tests' stub carries the _is_kv_manager_v2 #20086 reads

Features

feature	how it is selected
CUTEDSL MoE backend for the trunk's NVFP4 routed experts	moe_config.backend: CUTEDSL (TRTLLM/AUTO keep trtllm-gen)
MTP (draft length 1..7): • GDN cached replay • QSA sparse attention in the draft layer • IndexShare • speculative rollback of the PLE and QSA side state	MTPDecodingConfig (the legacy TRTLLM_USE_GDN_REPLAY=0 path stays at draft <= 3)
Video input; the vision tower runs in groups of whole frames, with admission-time validation	automatic
FP8 KV cache (unit scales; a lossy substitution of the checkpoint's bf16 K/V)	kv_cache_config.dtype: fp8
NVFP4 KV cache: • uint8 e2m1 pages + e4m3 block scales • QSA side planes packed so the data and scale pools keep one page-index scale	kv_cache_config.dtype: nvfp4
Gated-DeltaNet state in bf16 (halves the recurrent state: 113 MB -> 57 MB per slot)	kv_cache_config.mamba_ssm_cache_dtype: bfloat16 (the engine's auto resolves to bf16 for this checkpoint)
Canonical QSA selection: the selection row is a pure function of the scores, the window and k (every block strictly above the k-th score, ties to the highest index), not of the radix kernel's launch order	always on for the sparse path
Breakable prefill CUDA graph	prefill_cuda_graph_backend: breakable

Routing declines what the target's constructor would refuse, from main's kv_cache_quant_algo (#20267) and the engine choices qwen4_exp derives from the ModelConfig and LLM args (EngineChoices):

• an unsupported KV algorithm, model dtype, MoE backend or recurrent-state dtype;
• a draft length that changes at runtime;
• CUTEDSL, or the legacy GDN verify path, with draft > 3;
• the piecewise prefill graph.

So modeling_v2="auto" falls back to the built-in model instead of failing mid-build.

Performance

Setup:

• ISL 8192 / OSL 1024, 1x GB300, modeling_v2: require.
• Storage and backends: bf16 KV, fp32 GDN state, trtllm-gen MoE.
• Engine: max_batch_size 128, max_num_tokens 16384, chunked prefill, CUDA graphs at batch 1..128, free_gpu_memory_fraction 0.6.

	con1 tok/s/user	con128 tok/s/GPU
feature commits only	90.4	2,122
with the fast paths	~370 *	4,787 - 4,827 **

* Measured at an intermediate step of the perf work; the later steps targeted the mixed and con128 decode steps.

** 3 same-host rounds, TPOT 24.8 ms, TTFT 1.87 s. Measured before the last four perf steps, which were measured per kernel in place: 0.4 - 1.4 ms per mixed step each, and 55 us per decode step for the logits cast.

What runs on the fast paths:

• Decode steps run as chains of programmatically dependent launches:
  • attention: one launch from the input projection's output to the index scores (QK norm + rope, index store and compression) when a step has one row per request (MTP verify steps keep the composed launches); the block top-k fused with its expansion to tokens; a CuTe DSL split-K sparse GQA;
  • Gated-DeltaNet: the input projection with the conv step as its epilogue; the recurrence and the gated norm in one launch (small batches) or one streaming kernel over the state pool (larger batches);
  • each hyper-connection site: an inject + norm kernel, a tcgen05 first projection and an up projection with the sigmoid gate and the stream mean in its epilogue (two GEMV launches at one row); the MoE site's mixer writes the routed experts' NVFP4 input directly.
• Context and mixed steps:
  • a row-tile sparse GQA on tcgen05 (K/V tiles read once per group of rows of one request);
  • one Gated-DeltaNet prefill conv pass;
  • the shared expert's gate logit inside its gate|up GEMM;
  • the CuTe DSL MoE for launches of >= 4096 rows;
  • the PLE n-gram embedding in one launch;
  • a persistent CuTe QK norm + rope.
• The host-offloaded PLE table is mapped from registered pageable memory; on the pinned allocator's pages, the device's random row gathers were up to ~14x slower.
• The logits are widened bf16 -> fp32 in one vectorized kernel.

The fast paths cover bf16 KV and an fp32 GDN state, the checkpoint's own storage. FP8 / NVFP4 KV and a bf16 state keep the feature commits' paths. At con128 they are slower: FP8 KV -20% and bf16 state -11.5% against bf16 KV / fp32 state, measured at an intermediate step.

Shared code touched (please review)

• tensorrt_llm/_torch/pyexecutor/py_executor.py: _validate_request calls an optional validate_request(request) on the model (found through the executor's model engine, if it has one), mirroring sampler.validate_request; no model in the tree defined one before. An oversize or inconsistent multimodal request now fails alone instead of ending the executor. The executor stub of test_beam_search.py's admission test gains a model_engine.
• tensorrt_llm/_torch/pyexecutor/engine/runners/decoder/inputs.py: the mRoPE position delta is carried through one-model MTP verify groups (DecoderInputPreparer._decode_mrope_position_delta; _pack_extend_requests and _pack_generation_requests share HostInputs.has_gen_mrope_delta). This also changes verify-row rotary for built-in mRoPE models with a one-model drafter.
• tensorrt_llm/_torch/pyexecutor/model_loader.py, _experimental/modeling_v2/_router_index.py, tensorrt_llm/_torch/model_config.py:
  • effective_mamba_ssm_cache_dtype is factored out of validate_and_set_mamba_ssm_cache_dtype;
  • modeling_v2_resolve passes model_config and llm_args to every family's route(), after main's kv_cache_quant_algo: qwen4_exp derives its EngineChoices from them, gpt_oss and deepseek_v3 accept and ignore them;
  • moe_backend_requested keeps the requested MoE backend.
• tensorrt_llm/_torch/attention/backends/sparse/qsa/:
  • cache_manager.py packs the index and position planes into one side buffer when their page size could equal a K/V or block-scale page at any split of the KV heads over ranks (FP8 and NVFP4 KV at this geometry, never bf16). The decision is the same on every TP / attention-DP layout, so disaggregated peers agree on the transfer roles. This fixes a trunk exact_div assert and a silent draft-layer misaddressing; the built-in path could not construct its manager with NVFP4 KV before.
  • metadata.py: the draft block table and prepare_for_draft_forward.
  • kernels.py:
    • an additive displace_oldest flag (default off);
    • the index-score kernel takes use_pdl / prefetch (both default off);
    • the launch heuristics change for every caller: context-row tiles are 128 dot columns wide (was 16); batches past 4 rows use 64-block tiles (was 8 below 257 rows); a program of packed context rows walks 8 tiles (was 32). They were tuned on this target's geometry (24 q / 2 KV heads, head_dim 256) on GB300 and are not measured on other QSA models.
• cpp/tensorrt_llm/thop/attentionOp.cpp (in 109b161cd4 and 456f198a1e, needs reviewer attention):
  • mFP8ContextFMHA is also set for an FP4 KV cache when MLA is not enabled, on the SM100 family only. Only those FP8 context FMHA kernels read FP4 KV, and a bf16-output attention layer aborted in initialize() before; other architectures still fail AttentionOp's FP4 KV check.
  • The MLA / DSA / DeepSeek-V4 / GLM FP4-KV paths use mUseNvfp4MlaKvCache / mFP8ContextMLA. Reading the code, the expression is unchanged for them and for every non-FP4 call, but this is not measured on those models.
• modeling_v2 catalog:
  • attention/thop_attention.py: additive cells and constants (qwen4_exp FP8 / NVFP4 / MTP verify widths); the gpt_oss / deepseek cells and reference branches are unchanged;
  • graph/eager_on_graph.py: the breakable prefill graph's break;
  • attention/fused_qk_norm_rope.py: cells for this target's geometry.

Decisions worth knowing

• GVR Top-K (enable_heuristic_topk) is not included. Measured on GB300:
  • it is neutral at max_seq_len 8192 and 1-5% slower per decode step at 32K;
  • its selection kernel is 2.5-3.4x slower than radix on decode rows (it only ties at 2080 columns);
  • it adds a bind-time warm-up and a new capture-time failure mode.
• Canonical selection: radix selection depends on launch order. Over 50 launches of 256 rows, every row came back in a different order and up to 252 rows selected more than one set; canonical rows were identical in all cases.
  • As a composition, canonical selection costs decode +2..4% per step and prefill +1% (8K) to +6% (32K / 131K tables).
  • The fast paths produce the same rows from one fused launch: 16K context rows 287 vs 1,721 us; 128 decode rows 11.4 vs 110 us.
  • The composition remains for score tables wider than 4096 columns and in the MTP draft layer.
• Second expert copy under trtllm-gen MoE (the default backend for this target): the trunk keeps a CuteDSL-layout copy of the routed experts so that context launches of >= 4096 rows run the CuteDSL MoE.
  • GB300 per layer, CuteDSL vs trtllm-gen: 8192 rows 451 vs 555 us; 16384 rows 656 vs 882 us.
  • Cost: about 1.3 GiB per layer, about 63 GiB over the 48 layers, taken from the KV cache budget.
  • It is kept only when max_num_tokens >= 4096, the smallest launch it serves; below that every launch stays on trtllm-gen and the memory stays in the KV pool. TRTLLM_MODELING_V2_MOE_PREFILL_COPY=0 drops it for memory-constrained deployments (gsm8k 95.98 without it, 96.17 with it).
• Breakable prefill graph only removes launch overhead:
  • TTFT (median of generate(max_tokens=1)), eager -> BCG: 128 tokens 158 -> 30 ms, 512 tokens 104 -> 33 ms, 2048 tokens 104 -> 69 ms, 6000 tokens 204 -> 205 ms (MTP3: 123 -> 43, 124 -> 42, 124 -> 73, 212 -> 213).
  • Costs: +5 s boot, about 0.4 s per bucket, +3.4 GiB graph pool for buckets up to 3072 (+6.5 GiB up to 6144).
  • A prompt that fills its bucket is bitwise identical to eager; a padded prompt differs as a different eager batch size does (top-token logprob up to ~0.25 nats).
  • At con128 it is neutral: the mixed steps are GPU-bound.
• bf16 recurrent state: MTP with the engine default mamba_ssm_cache_dtype: auto is now claimed by this target (it fell back to the built-in model before); fp16 is refused.

Known limitations / not verified

• NVFP4 KV works but is not fast yet. Its sparse GQA kernels spill registers at head_dim 256: about 948 ms per layer on a 16K-row mixed step, 134 tok/s/GPU at con32. A fix is prototyped and will follow separately: the context rows read pages dequantized into a bf16 scratch through the row-tile pass, and the other rows use 16-token split-K tiles.
• Memory headroom with the second expert copy: in an earlier measurement, free_gpu_memory_fraction: 0.9 did not boot; the perf configuration runs at 0.6. TRTLLM_MODELING_V2_MOE_PREFILL_COPY=0, or max_num_tokens < 4096, gives the copy's memory back to the KV pool.
• Speculative rollback has no end-to-end output oracle: all-rejected outputs differ from no-speculation through GEMM / attention kernel selection by row count. It is guarded by a bitwise unit test that replays all-1 and all-8 acceptance; forced-acceptance runs measured acceptance length exactly 1.0 / 8.0 at draft 7 during bring-up and are not kept as a gate.
• NVFP4 KV coverage:
  • the thop cells cover P <= 4; P = 5..8, G = 32 and keys past 2048 are not covered;
  • NVFP4 requires chunked prefill, block reuse or speculation (a bind-time assert; route() cannot see those).
• Not verified:
  • disaggregated serving with MTP on the generation worker only;
  • tokens_per_block != 32;
  • bf16 state at >= 128 rows with the FlashInfer-compatible mapping;
  • BCG combined with bf16 state, CUTEDSL or MTP5-7;
  • BCG buckets above 6144 tokens;
  • BCG with video in CI.
• Lossy storage: the bf16 state and the NVFP4 / FP8 KV substitutions are lossy storage choices, and the references have no discriminator for the SSM dtype.
• Lossy KV and deep drafts: at draft 7 the KV dtype sets the acceptance length on the long pass keys: bf16 7.857, FP8 7.500, NVFP4 6.87 - 6.97 (IndexShare and the bf16 state each move it by about 0.1). At draft 3 the same lossy KV costs about 0.03. Every pass key is still stated back; the built-in model at draft 7 over NVFP4 KV was not measured.
• scripts/cbts: these tests were not run in the feature validation; one unix-socket test hangs in the container, unrelated to this diff.

Test Coverage

GB300 (sm_103), TP1, modeling_v2: require, full gsm8k (1319 samples).

Final head (13c21f460b, rebased on main af3e5eb257). The rebase applied every commit without conflicts; the sections below it were measured on the pre-rebase hashes.

check	result
test_kv_connector_initialization.py, test_py_executor.py, test_beam_search.py -k vbws_rejects, disaggregated/test_chunked_transfer.py, test_trtllm_serve_config_schemas.py	431 passed; 1 failed: test_checked_in_schemas_are_current, which main fails too since 93b17dcdd9 set the version to 1.4.0rc2 without regenerating the snapshots (waived)
unittest/_torch/modeling_v2 (not comm)	1213 passed, 43 skipped
l0_cpu.yml's modeling_v2 files, -m cpu_only	407 passed, 43 skipped
QSA unit tests	101 passed
mypy (scripts/run_mypy.sh)	clean
test_gsm8k	96.25 (ref 96.36)
test_mtp1_gsm8k_mmmu_and_video	gsm8k 95.98 (ref 95.57), acceptance length 1.971; MMMU 57.89 (ref 59.22); videos pass
test_mtp3_long_context_passkey	pass, acceptance length 4.000

Test slimming (f29b8c6f3b, before the rebase):

check	result
test_gsm8k (unchanged)	not rerun
test_mtp1_gsm8k_mmmu_and_video	gsm8k 95.53 (ref 95.57), acceptance length 1.969; MMMU 58.33 (ref 59.22); videos pass
test_mtp3_long_context_passkey	pass, acceptance length 4.000
test_mtp3_cutedsl_fp8_kv_bf16_state	gsm8k 95.94 (ref 95.94), acceptance length 3.509; pass keys, acceptance length 3.962
test_mtp7_long_context_passkey (NVFP4 KV, bf16 state)	pass, acceptance length 6.867 / 6.967 (two runs)
test_breakable_prefill_matches_eager (unchanged)	not rerun
unittest/_torch/modeling_v2 (not comm)	1213 passed, 43 skipped
l0_cpu.yml's modeling_v2 files, -m cpu_only	407 passed, 43 skipped
mypy (scripts/run_mypy.sh), pre-commit over the changed files	clean

Review fixes (6b039cb3c6):

check	result
QSA unit tests, including the TP1 / TP2 / ADP2 transfer-view regression test for bf16, FP8 and NVFP4	805 passed
gsm8k, with and without the second expert copy	96.17 / 95.98 (ref 96.36)
MTP1 gsm8k over FP8 KV / acceptance length	96.02 / 1.963
long-context passkey, MTP3 over FP8 / NVFP4 KV / with IndexShare	pass, acceptance length 3.954 / 3.949 / 4.000
unittest/_torch/modeling_v2 (not comm)	3444 passed, 43 skipped
CuteDSL MoE unit tests, this branch and main (test_cute_dsl_moe.py)	904 passed each

71b2b3cf66 (the top-k compile fix on d0b2d5dfa6):

check	result
qsa_topk_expand_decode and the target's other attention ops' unit tests, including one compiled kernel across 1 to 300 rows	658 passed
a call at a row count the process has not run yet (10 counts, 1 to 16256 rows)	0.1 ms (360 - 410 ms before, nearly all of it the compile)
gsm8k	96.06 (the gate took 345 s, against 411 s on d0b2d5dfa6)
long-context passkey, MTP3 with IndexShare	pass, acceptance length 4.000

d0b2d5dfa6 (on main c14609482b):

check	result
gsm8k	95.75 (ref 96.36)
MTP1 gsm8k / acceptance length	96.36 / 1.969
MMMU, MTP1	57.67 (ref 59.22)
long-context passkey, MTP3 with IndexShare	pass, acceptance length 4.000
long-context passkey, MTP7 (3k-12k tokens)	pass, acceptance length 7.839
unittest/_torch/modeling_v2 (not comm)	3436 passed, 43 skipped
l0_cpu.yml's modeling_v2 files, -m cpu_only	437 passed, 43 skipped
test_beam_search.py -k vbws_rejects	5 passed
test_pytorch_model_engine.py, test_py_executor.py, disaggregated/test_chunked_transfer.py, QSA, test_model_config.py	559 passed
pre-commit over the PR's files (mypy through scripts/run_mypy.sh in the container)	clean

Before the rebase (dcff1146d0, on main 72cd8b57a3):

check	result
gsm8k	96.171
MTP1 gsm8k / acceptance length	95.603 / 1.971
long-context passkey, MTP3 with IndexShare	pass, acceptance length 3.998
breakable prefill graph against eager, with the decode overlap asserted	pass
test_beam_search.py -k vbws_rejects (failed in CI on f22959cf9a)	5 passed
pre-commit over the changed files	clean

Before the review fixes (the same eleven patches as f22959cf9a, on main a3766d68f9):

check	result
gsm8k	95.982
MTP1 gsm8k / acceptance length	95.641 / 1.969
MMMU, MTP1	59.667 (ref 59.22)
video, MTP1	pass
unittest/_torch/executor	GPU: 4048 passed; -m cpu_only with no GPU visible: 2663 passed
QSA, qwen4_exp support and runtime wiring, test_model_config	205 passed

In unittest/_torch/executor the only failures were:

• tests main waives (nvbugs/6818710): 35 on GPU, 11 of them under -m cpu_only;
• the MpiPoolSession tests, which cannot start workers in the test container (one timed out, two were deselected);
• on GPU, the five test_overlap_scheduler.py tests, whose Qwen3-0.6B checkpoint this cluster does not have.

All 26 qwen4_exp gates passed on the same eleven patches before the rebase onto #20134 (02e2d7a7ea on main 3163e66b2c):

gate	gsm8k	acceptance length
no MTP	96.247
CUTEDSL experts	96.096
bf16 GDN state	95.868
breakable prefill graph	96.171
MTP1	95.830	1.969
MTP1, FP8 KV	95.982	1.963
MTP1, NVFP4 KV	96.361	1.959
MTP3	95.754	3.693
MTP3, CUTEDSL experts	96.323	3.701
MTP3, bf16 GDN state	96.133	3.696
MTP3, breakable prefill graph	95.944	3.697
MTP5	95.754	5.065
MTP7	96.058	6.048

The other gates:

• MMMU at MTP1: 56.333 (ref 59.22); 59.111 with the breakable prefill graph.
• Long-context passkeys: MTP3 with IndexShare, FP8 KV, NVFP4 KV and the breakable prefill graph; MTP7 with the fp32 and the bf16 state.
• Forced acceptance at draft 7, all rejected and all accepted.
• MTP1 video; MTP1 on the legacy GDN verify path; the breakable prefill graph against eager.

The acceptance lengths were read from the worker's iteration statistics. The test cluster cannot spawn MPI workers, so the LLM API worker runs in-process, and there LLM.get_stats() returns nothing. The engine keeps the last 1000 iterations, so the values cover the end of each run. CI runs the gates unmodified.

Earlier validation:

• Feature commits, before the perf work:
  • the gates passed with gsm8k 95.75 - 96.36;
  • acceptance length 1.971 / 3.704 / 4.428 / 5.063 / 5.609 / 6.052 at draft 1 / 3 / 4 / 5 / 6 / 7;
  • KV bytes per token, bf16 : FP8 : NVFP4 = 29976 : 16664 : 10840 (1 : 0.56 : 0.36).
• Disaggregated serving (ctx TP1 -> gen TP1, NIXL), gsm8k:
  • no speculation 95.830; MTP1 on both workers 95.982 (AL 1.969); FP8 KV 95.982; NVFP4 KV 96.740; CUTEDSL 95.906;
  • greedy tokens match the aggregated engine on prompts of 32 to 6000 tokens (including 2048, 2049, 2050, 4096 and 4097), with bf16 and NVFP4 KV, on the sequential passes.
• Fast paths:
  • every new or changed op has a GPU unit test parametrized on its certified cells;
  • gsm8k on the performance configuration after the port and after each step that changes numerics: 95.94 - 96.02; with MTP3, 96.17 (the deployment guide's stock NVFP4 + MTP3 number is 95.57);
  • the remaining steps are bitwise equal to what they replace, per their unit tests.

CI entries (l0_b300.yml) are chosen for cost: every gate builds the full engine, and a gsm8k gate takes roughly 8-13 min on one GB300. test_qwen4_exp.py holds six cases, each engine running every check that fits in it, and all six are listed:

• Pre-merge (3):
  • test_gsm8k;
  • test_mtp1_gsm8k_mmmu_and_video: one MTP1 engine for gsm8k with the acceptance-length floor (the only gate that sees a miscomputed draft layer), MMMU and synthetic videos. It replaces main's test_mtp1_gsm8k_and_acceptance and test_mtp1_mmmu;
  • test_mtp3_long_context_passkey: IndexShare over bf16 KV, the only prompts long enough to run the sparse attention path.
• Post-merge (3):
  • test_mtp3_cutedsl_fp8_kv_bf16_state: CuteDSL experts, FP8 KV and a bf16 state in one engine, gsm8k and then the pass keys;
  • test_mtp7_long_context_passkey: seven drafts (the verify step's eight rows per request) over NVFP4 KV and a bf16 state;
  • test_breakable_prefill_matches_eager.

l0_cpu.yml adds the target's CPU tests: the multimodal splice and admission check on a stub tower, the KV pool contract, the KV storage derivation and the attention-call arguments.

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.

Dev Engineer Review

The inspected checkout shows formatting-only import changes in three source files. It shows no behavior, API, or default changes in those files. The supplied change inventory describes a much broader Qwen4-Exp implementation, but that scope is not present in the inspected diff and cannot be verified here.

No current review findings were supplied. Test results and completion status are unavailable.

QA Engineer Review

One test file changed, but only its imports were reformatted. No test functions, test IDs, selectors, or waivers changed. No integration test-list files changed. The modified unit test is not expected in the test-db/ or qa/ integration lists. Coverage verdict: needs follow-up to reconcile the supplied change inventory with the inspected diff.

Per-File QA Perspective

• tensorrt_llm/_torch/model_config.py — Import formatting only; no observable behavior change to verify.
• tensorrt_llm/_torch/pyexecutor/model_loader.py — Import formatting only; no observable behavior change to verify.
• tensorrt_llm/_torch/pyexecutor/py_executor.py — Import formatting only; no observable behavior change to verify.
• tests/unittest/_torch/executor/test_pytorch_model_engine.py — Import formatting only; existing model-engine tests remain, with no new or changed test coverage. No integration-list entry applies.

[Wanli-Jiang]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #78865 [ run ] triggered by Bot. Commit: bfa1ea6 Link to invocation

[Wanli-Jiang]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #79774 [ run ] triggered by Bot. Commit: 257d5f0 Link to invocation

[tensorrt-cicd]

PR_Github #79774 [ run ] completed with state SUCCESS. Commit: 257d5f0
/LLM/main/L0_MergeRequest_PR pipeline #65674 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

[Wanli-Jiang]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #79889 [ run ] triggered by Bot. Commit: 257d5f0 Link to invocation

[tensorrt-cicd]

PR_Github #79889 [ run ] completed with state SUCCESS. Commit: 257d5f0
/LLM/main/L0_MergeRequest_PR pipeline #65766 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

[ZhanruiSunCh]

LGTM for infra part

[Wanli-Jiang]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #80019 [ run ] triggered by Bot. Commit: 257d5f0 Link to invocation

[tensorrt-cicd]

PR_Github #80019 [ run ] completed with state FAILURE. Commit: 257d5f0
/LLM/main/L0_MergeRequest_PR pipeline #65860 completed with status: 'FAILURE'

CI Report

⚠️ Resource Maintenance — Test Coverage Reduced:

Resource maintenance in LLM/main/L0_Test-x86_64-Multi-GPU #2983: B200 multi-GPU stages are paused to relieve the CI queue backlog (109 running / 459 pending on 2026-10-22). Owner: @sharabeshj-nvda. The CI team decides when to lift it.
Affected stages:
DGX_B200-2_GPUs-PyTorch-1
DGX_B200-4_GPUs-PyTorch-1
DGX_B200-4_GPUs-PyTorch-2
DGX_B200-4_GPUs-PyTorch-3
DGX_B200-4_GPUs-PyTorch-Ray-1
DGX_B200-8_GPUs-PyTorch-1
DGX_B200-8_GPUs-PyTorch-2
DGX_B200-8_GPUs-PyTorch-3
DGX_B200-8_GPUs-PyTorch-4
DGX_B200-8_GPUs-PyTorch-Ray-1

⚠️ 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

[Wanli-Jiang]

/bot run --disable-fail-fast

[tensorrt-cicd]

PR_Github #80081 [ run ] triggered by Bot. Commit: 257d5f0 Link to invocation

[tensorrt-cicd]

PR_Github #80081 [ run ] completed with state SUCCESS. Commit: 257d5f0
/LLM/main/L0_MergeRequest_PR pipeline #65903 completed with status: 'SUCCESS'

CI Report

⚠️ Resource Maintenance — Test Coverage Reduced:

Resource maintenance in LLM/main/L0_Test-x86_64-Multi-GPU #3317: B200 multi-GPU stages are paused to relieve the CI queue backlog (109 running / 459 pending on 2026-10-22). Owner: @sharabeshj-nvda. The CI team decides when to lift it.
Affected stages:
DGX_B200-2_GPUs-PyTorch-1
DGX_B200-4_GPUs-PyTorch-1
DGX_B200-4_GPUs-PyTorch-2
DGX_B200-4_GPUs-PyTorch-3
DGX_B200-4_GPUs-PyTorch-Ray-1
DGX_B200-8_GPUs-PyTorch-1
DGX_B200-8_GPUs-PyTorch-2
DGX_B200-8_GPUs-PyTorch-3
DGX_B200-8_GPUs-PyTorch-4
DGX_B200-8_GPUs-PyTorch-Ray-1

Link to invocation