Repository navigation
[None][feat] Rubin merge back: MLA, DSpark, DWDP, Helix - #21067
mingyangHao wants to merge 28 commits into
Conversation
|
PR_Github #81701 [ run ] completed with state
|
|
/bot run --disable-fail-fast |
|
PR_Github #81827 [ run ] triggered by Bot. Commit: |
|
PR_Github #81827 [ run ] completed with state
|
JunyiXu-nv
left a comment
JunyiXu-nv
left a comment
There was a problem hiding this comment.
Runtime part LGTM (autotuner short-list sharding, the MNNVL fusion-op entry and the padding-slot guard): each change is gated to PARALLEL distributed tuning, a fusion op nothing requests yet, or rejection sampling with padding slots, so default behavior and the per-iteration path are unchanged. I did not review the kernel/model parts. Note that functional.py now conflicts with main's AllReduceParams.fp8_payload; keep both that and the extended assert.
Resolve the MNNVL AllReduce overlaps by retaining the FP8 payload dispatch from main and the FP8 block-scale scale-factor leading dimension from the port. Preserve both the payload option and residual-free block-scale fusion in AllReduceParams. Signed-off-by: Mingyang Hao <mingyangh@nvidia.com>
|
/bot run --disable-fail-fast |
|
PR_Github #81920 [ run ] triggered by Bot. Commit: |
|
|
||
| @functools.cache | ||
| def _compile_dspark_rmsnorm_rope_draft_block(block_size: int, eps: float): | ||
| def _compile_dspark_rmsnorm_rope_draft_block(block_size: int, eps: float, device: int): |
There was a problem hiding this comment.
Why this change is needed? I assume TRTLLM has a one-process-one-gpu map and other kernels also don't have this.
| main_qkv, qkv = qkv_rows.split((main_rows.shape[0], block_rows.shape[0]), dim=0) | ||
| main_qkv = main_qkv.view(*main_x.shape[:-1], qkv_rows.shape[-1]) | ||
| qkv = qkv.view(*x.shape[:-1], qkv_rows.shape[-1]) | ||
| q_latent, kv = qkv.split((q_lora_rank, qkv.shape[-1] - q_lora_rank), dim=-1) |
There was a problem hiding this comment.
(By Codex)
This split leaves KV with a row stride of q_lora_rank + 512. is_fused_dspark_attention_preparation_supported() accepts this regular strided layout, but cute_dsl_dspark_rmsnorm_rope_draft_block() subsequently requires x.is_contiguous() through _validate_dspark_preparation(). The input therefore passes the support check but fails preparation, even for batch size 1 with a 5/6-token block. Please align the producer’s layout with the preparation kernel’s requirements.
There was a problem hiding this comment.
The new main_kv_input.contiguous() and block_kv_input.contiguous() calls add GPU copies before preparation when the QKV slices are non-contiguous. Please confirm this is intended.
| ) | ||
| return (h, ffn_pre) if pre_mix is not None else h | ||
|
|
||
| def forward( |
There was a problem hiding this comment.
(By Codex)
Several paths appear to lack production integration: the managed-page forward(kv_pages=...) / write_context_pages() path, fused_dsv4_dspark_attention_fp8() whose model-side caller is in that managed-page path, and configure_generation_autotuner(). The worker still uses forward_batched() with rolling windows and does not call the autotuner configuration entry point. Are these intentionally staged for a follow-up integration, missing wiring in this PR, or leftovers from an earlier implementation? Please clarify the intended call sites and whether these paths should remain in this change.
|
PR_Github #81920 [ run ] completed with state
|
sunnyqgg
left a comment
sunnyqgg
left a comment
There was a problem hiding this comment.
Reviewed at 1034f82 (33 files, +7185/-1808). Findings by level: 7 blocking, 12 possibly blocking (a measurement or a closer look decides) and 10 required; nice-to-improve items are left out here. The inline comments cover points with a single anchor in the diff. The items below have no such anchor, or they extend an existing thread.
-
[Blocker] +1 to pengbowang-nv's thread on tensorrt_llm/_torch/models/modeling_dspark.py:4855 (strided KV split vs
_validate_dspark_preparation). It is wider than the KV input, and it hits the default configuration:- With
TRTLLM_DSPARK_LINEAR_BACKENDunset and TP=1 or attention DP,post_load_weights(weights)bindsqkv_proj. - On SM100/103 the strided
q_latentfrom_qkv_projectionsgoes through_q_projection→_rmsnorm_rope_batched→cute_dsl_dspark_rmsnorm_rope. - That op's gate (
_has_regular_row_stride, dspark_rmsnorm_rope_custom_op.py:55-72) accepts the strided view, but its kernel is compiled withmake_fake_compact_tensor((rows, hidden), stride_order=(1, 0))under--enable-tvm-ffi(:200-239). The TVM-FFI stride check rejects it for G·block ≥ 5. main_kvreaches the compact-compiled cache-write op the same way.- So DeepSeek-V4 + DSpark on B200/B300/GB200/GB300 fails at the first draft forward. The non-waived pre-merge
TestDeepSeekV4FlashDSpark::test_gsm8k_1p1d_dep4(l0_dgx_b200.yml:211,enable_attention_dp=True) takes this path.
A KV-only fix is not enough. Make
_qkv_projectionsreturn compact tensors, or have every DSpark fused-op gate requireis_contiguous(). Then add a B200 test that runs with the real bound FP8 projections and the fused ops enabled. - With
-
[Blocker] +1 to coderabbitai's thread on tensorrt_llm/_torch/models/modeling_dflash.py:1988 and pengbowang-nv's on :1976. The unvalidated folded ratio crashes supported configurations; it is not only a limit question:
- For ratios above 64, trtllm-gen picks
tileSizeQ=128and then requiresnumHeadsQ % min(ratio, 128) == 0(flashinfer fmhaKernels.cuh:353-357). - So any TRTLLM-backend drafter with
block_size × groupabove 128 and not a multiple of 128 now fails the first draft forward with "The numHeadsQ/numHeadsKv is not supported". Examples are group 12 × block 16, and the CI's Nemotron-3.5-Lightning-30B-A3B DFlash with group 16 × block 9–15. - Init only checks
group <= 32(dflash_attention.py:176-182). At HEAD these layers ran on the context kernel.
Please validate
group * block_sizeagainst the folded decode at init, or split the fold. Report the folded lowering invalidate_block_attention_windows, and drop the unreachable context-kernel branch. - For ratios above 64, trtllm-gen picks
-
[Blocker] CI: no pre-merge run has passed at any revision:
- The head run (L0_MergeRequest_PR #67568 at 1034f82) failed on infrastructure. Three DGX_B200 stages could not download TensorRT-LLM.tar.gz, DGX_H100-PyTorch-4 lost its Slurm monitor, and SBSA setup aborted.
- #67341/#67362/#67480 at 72e6cc7 were UNSTABLE, with the multi-GPU stages blocked on the label.
This PR touches allreduceOp.cpp, communicationKernels and _torch/distributed, so the multi-GPU stages matter. None of these has run:
unittest/_torch/multi_gpu(l0_dgx_h100),test_autotuner_distributed_strategy,test_dspark_drafter_tp_fusion,TestDeepSeekV4FlashDSpark::test_gsm8k_1p1d_dep4(l0_dgx_b200), andTestDeepSeekV41::test_gsm8k_dsparkon GB300. Given the strided-view crash above, the non-waived Flash dep4 case should fail as the code stands. Please re-run now thatci: full pre-merge approvedis set, and link a run where the B200, SBSA single-GPU and multi-GPU stages all passed. -
[Possibly blocking] +1 to pengbowang-nv's thread on tensorrt_llm/_torch/models/modeling_dspark.py:2317 about
configure_generation_autotunerhaving no production caller. It also costs performance on Rubin:- It is the only place that installs DSpark's bucketed generation tuning specs.
- The default
_DSPARK_GENERATION_M_SPEC = DynamicTensorSpec(0, 0)(dspark_attention_custom_op.py:45) keeps M exact, and the autotune warmup runs a single generation request. - So on SM107 with the default FP8 binding, the FP8 draft projections and the OA BMM are tuned for one batch size only. Every other CUDA-graph batch size misses the cache and bakes in the fallback tactic (128×128, 1×1 cluster, no swap_ab).
Please install bucketed specs at bind time from
max_batch_size, or call it from the worker before the autotune warmup. Measurement that decides it: on SM107, DeepSeek-V4 + DSpark with ADP and CUDA graphs for batch 1..128. Count the AutoTuner "fallback tactic, due to cache miss" logs for the DSpark ops, and compare draft-step latency at G ∈ {2, 8, 32, 128}, tuned vs fallback. A regression above about 2% vs HEAD makes it a blocker. -
[Possibly blocking] Also on pengbowang-nv's :2317 thread, about whether the managed-page path belongs in this PR. Besides the
SymIntcompile error infused_dsv4_dspark_attention_fp8(inline on dspark_attention_custom_op.py:171), the first caller of the paged stack would inherit these defects:- The
q_b_norm_enabled=Falsebranch (:725) applies the single-request_rope_last_dimsto[G, block, rd/2]freqs and raises for G ≥ 2. V4.1 sets the flag to False. - The CuTe branch (:727) is chosen by arch and block only, so strided inputs or non-BF16 pools raise instead of falling back. The fallback's Triton writer (attention/backends/dspark.py) casts BF16 into an FP8 pool at scale 1 and then fails in
torch.cat. - The kernel (attention_kernel.py:1146-1175, :1196-1197) and the Triton writer bound page-table columns only by capacity, not by table width, and do not mask rows beyond capacity or with page -1.
- The writer makes the token count and strides
tl.constexpr, so every new shape triggers a JIT compile.
I'd drop the paged stack from this PR until the worker uses it.
- The
-
[Possibly blocking] +1 to coderabbitai's thread on cpp/tensorrt_llm/kernels/communicationKernels/mnnvlAllreduceKernels.cu:561. The misalignment is real:
- With
accessGroupSize=1for this pattern (:833-836, :1630-1631),adjustGridConfigpicks for example 72-thread CTAs for hidden 2304. 1152/2816/3840/4608/5376/6656/7680 also misalign at small token counts. - 16-lane groups then straddle 128-element blocks in both one-shot and two-shot, and some elements are dequantized with a scale off by 2×, with no error.
- K3's 3584/7168 are fine, and the out-of-bounds early-return part does not apply when
hidden % 128 == 0.
It is latent until something calls the fusion. Please pass an access-group size of 16 for this pattern and use a 16-lane shuffle mask.
- With
-
[Required] +1 to coderabbitai's thread on tests/unittest/_torch/speculative/hw_agnostic/test_kimi_k3_dspark_semantics.py:887. Confirmed:
- DGX_H100-PyTorch-5 skipped both parametrizations ("needs an SM100-family GPU").
- l0_cpu does not collect the file, since it has no cpu_only marker.
- No B200/B300 list selects it.
It is the only GPU test of the new modeling_dflash.py code (folded decode, FP8 page writes, workspace and counter sizing). Please list it in l0_b200.yml next to test_dspark_noncausal_window.py.
-
[Required] PR description and title. The description is only the CodeRabbit summary, and it describes the PR before it was trimmed: the cache-domain API, the standalone draft cache and its disagg transfer (removed in 36ae25e after dc3671's request), helixAllToAll.cu, mla.py, kv_cache_manager_v2.py and other files that are not in the current 33-file diff. The template's Description and Test Coverage sections are empty. Nothing in it covers what actually changes defaults here:
- the DSpark FP8 linear backend;
- the default-on FP4 MLA wave plan;
- the three new env vars (TRTLLM_DSPARK_LINEAR_BACKEND, TRTLLM_FP4_MLA_PACKED_WAVE_PLAN, TRTLLM_HELIX_V_UP_BEFORE_A2A);
- the MNNVL FP8 block-scale fusion.
It also gives no test results, and "DWDP" in the title matches nothing in the diff. Please rewrite the description for the current diff (motivation for each piece, default changes and env vars, tests and the hardware they were validated on) and fix the title.
-
[Required] Scope. AGENTS.md: "One concern per PR — avoid scope creep. If a PR touches unrelated areas, split it." This PR bundles at least eight unrelated changes:
- the MNNVL FP8 block-scale fusion;
- autotuner short-list sharding;
- the FC12 MoE
x_ready_event; - the rejection-sampling padding-slot check;
- the FP4 MLA packed-decode kernel rewrite and wave plan;
- DSpark paged attention and FP8 projections;
- DFlash FP8 pages and folded decode;
- K3 Helix V-up binding.
Three of them have no production caller: the MNNVL fusion, the paged DSpark stack, and the FC12 event join (
ctx.x_ready_eventis only set when the backend hasquantize_input_async, moe_scheduler.py:319, and FC12 does not define it). Please split the PR by area so each owner can review their part, and drop the unwired pieces until the PRs that wire and test them.
| _arrive_smem_p4_mapped_leader_mbar(leader_pair_p_source_ready_mbar) | ||
| elif prims.elect_sync(): | ||
| _arrive_smem_p4_mapped_leader_mbar(leader_pair_p_source_ready_mbar) | ||
| # P_SOURCE_READY is the only P-side signal the MMA warp needs: it implies |
There was a problem hiding this comment.
[Blocker] This tail no longer orders the compute warps' next row-max publish after the correction warps' read of the current one. Before this change, compute warps (0-3) waited on direct_p_ready_mbar, which all 8 producer warps arrive on only after reading the row max and storing P, plus the SLOT0/SLOT1 named barriers, before moving to the next tile. Now they only advance() here, and the remaining wait below (pair_rowmax_read_done_mbar) orders the opposite direction: correction warps wait for compute warps.
Since QK now runs two tiles ahead, qk_full(t+1) is usually already complete. A compute warp can therefore enter tile t+1 and _publish_p4_n256_atomic_rowmax (:3776, inside _load_p4_n256_score_half_from_tmem, before p4_smem_slot0_producer.acquire() and the DUAL_HALF barrier) into the single-buffered sAtomicRunningRowMax while a correction warp of the same rows still has to do its tile-t read after the DUAL_HALF barrier (:5119/:5134). That warp then reads max(M_t, m_{t+1}). On tile 0 the anchor is taken directly, and later tiles can rebase differently, so the two 128-column halves of a row are exponentiated against different anchors and O/LSE come out silently wrong. running_row_anchor is per-warp, so the mismatch lasts to the end of the KV range.
Every task that takes the per-tile loop on SM107 reaches this in the default packed FP4 MLA decode: 2-tile splits (pipelined_tiles < 2), masked tail tiles of packed MTP/Helix decode, and the folded path. The pipelined path is fine because sPipelinedRowMax is double-buffered by tile parity. The race is timing-dependent, so no current test catches it.
Fix: double-buffer the per-tile row max by tile parity like sPipelinedRowMax. Alternatively, add a "row-max read done" mbarrier that all 8 producer warps arrive on after reading new_row_max, and make compute warps wait on it before publishing tile t+1. A stress test that delays the correction warps (for example a debug-only nanosleep) and compares against the reference would pin it.
| page = cutlass.Int32(-1) | ||
| if position >= max(0, end - length) and position < capacity: | ||
| page = common_params.mPT_win[position // self.history_page_size, batch] | ||
| common_params.sPT[row, load_pt_producer_state.index] = page |
There was a problem hiding this comment.
[Blocker] The window page ids are now written with plain st.shared and published through load_pt_pipeline.producer_commit, but load_pt_pipeline is a PipelineCpAsync (:2590). Its producer commit lowers to cp.async.mbarrier.arrive.noinc (in cutlass-dsl 4.8.0, AsyncLoad maps to cute.arch.cp_async_mbarrier_arrive_noinc). That instruction only tracks the thread's earlier cp.async operations, so it does not order generic stores, and there is no fence before it. The TMA warp then reads sPT[context_row, stage] in load_history (:1410/:1419) and uses it as the TMA page coordinate. Under the PTX memory model it can therefore see a stale or uninitialized page id and load the wrong K/V, giving silently wrong draft attention.
At HEAD the window page id was loaded with cpasync.CopyG2SOp, which that arrive does cover; only the single draft-block slot used .fill. This PR moves the whole window stage (128 ids per stage) to plain stores, and the production rolling-window DSpark path reaches it through the ring adapter on SM100/103/107. I don't expect it has been observed, but the fix costs nothing: use PipelineAsync (its mbarrier.arrive has release semantics) for the page-table stage, or add a shared-memory fence and a plain mbarrier.arrive after the stores.
| # Helix MLA generation applies the V up-projection of all heads before the CP | ||
| # all-to-all (see ``MLA._helix_project_v_partials``), which needs every head's | ||
| # BF16 V weight on each rank. | ||
| _K3_HELIX_V_UP_BEFORE_A2A = os.environ.get("TRTLLM_HELIX_V_UP_BEFORE_A2A", "0") == "1" |
There was a problem hiding this comment.
[Blocker] TRTLLM_HELIX_V_UP_BEFORE_A2A does not control what it gates. Here is why:
KimiK3MLAAttention(MLA)(modules/kimi_k3_mla/kimi_k3_mla_attention.py:120) does not overridetransform_weights, and it is built withdtype=torch.bfloat16(:2102).- The post-load walk (model_loader.py:1987-1992) visits the MLA before its
kv_b_projchild. - So
MLA.transform_weights→_build_helix_v_up_weight()(mla.py:2511 → :1823-1865) binds the all-head BF16 V weight under Helix. It does so for bothkv_b_projlayouts the K3 loader leaves behind: BF16[H*n | H*v], and FP8 codes with FP32 128×128 scales, which the_Fp8BlockScaleWeightReadLinearswap keeps untransformed. - That sets
_helix_v_up_before_exchange = Truewhatever this variable says.
With the default "0", K3 under Helix CP still runs the V up-projection before the all-to-all. That also holds for anyone setting 0 as an A/B or escape hatch, like _K3_DISABLE_MIN_LATENCY_LATENT_PROJ right below. Setting "1" only adds an earlier bind at :3040-3044 / :3362-3363. That bind allocates a per-layer BF16 [num_heads_tp, 128, 512] copy plus FP32 temporaries before the walk rebinds the weight. No test sets the variable.
Fix: drop the env var and both K3 bind sites, since MLA already owns the binding. If an opt-out is really needed, put one documented flag inside MLA._build_helix_v_up_weight. Add a K3 Helix load test that asserts _helix_v_up_before_exchange with the flag unset and set.
|
|
||
| def _use_dspark_bf16_linear(mapping) -> bool: | ||
| """Select BF16 for an explicit override or sharded attention projections.""" | ||
| backend = os.environ.get(_DSPARK_LINEAR_BACKEND_ENV, "fp8").lower() |
There was a problem hiding this comment.
[Blocker] This new default changes DeepSeek-V4 DSpark draft numerics on every GPU, not just Rubin.
- With
TRTLLM_DSPARK_LINEAR_BACKENDunset and TP=1 or attention DP,_use_dspark_bf16_linearreturns False. post_load_weights(weights)→_bind_attn_projections()(:1601-1604, :2592) then binds the FP8 MLA linears for q_a/wkv, q_b, o_a and o_b.- There is no SM check, so on H100/H200/B200/B300 the draft attention projections switch from HEAD's BF16 dequantized weights (
_cache_attn_weights, HEAD :1273) to FP8 activation quantization. - On B200/B300 this default is also what currently crashes the first draft step (see the summary bullet on :4855). The pre-merge
TestDeepSeekV4FlashDSpark::test_gsm8k_1p1d_dep4runs with it.
Two more problems with the switch on this line:
- An explicit
TRTLLM_DSPARK_LINEAR_BACKEND=fp8is silently turned into BF16 when TP>1 without attention DP. - The variable is not documented anywhere.
The PR has no acceptance-length or accuracy data for the BF16 → FP8 switch on non-Rubin GPUs. Please:
- keep BF16 as the default outside SM107 until there is an A/B of acceptance length (BF16 vs FP8 draft projections on B200/GB300 and H200), or add that data here;
- reject an explicit
fp8that the mapping cannot honor, or at least warn; - document the variable.
| ) | ||
| output_scale_fake = cute.runtime.make_fake_compact_tensor( | ||
| cutlass.Float32, | ||
| (output_groups, group_dim // 128, (rows + 3) // 4 * 4), |
There was a problem hiding this comment.
[Possibly blocking] With the pinned nvidia-cutlass-dsl[cu13]==4.8.0 (requirements.txt:96), this line raises TypeError at compile time:
batch = cute.sym_int()(:152), sorows = batch * block_size(:162) is aSymInt(through__mul__).- In 4.8.0,
class SymIntincutlass/cute/typing.pydefines only__mod__/__rmod__/__mul__/__rmul__/__floordiv__. The only__add__in that file belongs toPointer. - So the first call of
trtllm::fused_dsv4_dspark_attention_fp8, the only op that passesoutput_groups, fails. Failed compiles are not cached (:304-309), so every later call fails again.
Today the only model caller is the paged dspark_attention_forward (modeling_dspark.py:751), which production does not call yet. But the PR's own SM107 test test_fused_dsv4_dspark_attention_fp8_matches_separate_quant_and_replays_graph would fail on Rubin, and nothing in CI runs it, because the file is not in l0_gr100.yml. Also, the FP8 variant's shared storage (about 233,472 B) looks larger than the 232,448 B available on SM100/103, yet _run_paged_dspark_attention accepts output_groups there.
Fix:
- Use an independent symbolic dimension for the padded scale rows (e.g.
cute.sym_int(divisibility=4)), or compute the padded rows on the host and pass them in. - Reject
output_groupsoutside SM107. - Run the SM107 test on a Rubin machine before merging the op.
| static_assert(kELTS_PER_THREAD == 8, "FP8 block-scale quant expects eight elements per 16-byte access."); | ||
| constexpr int kQuantBlock = 128; | ||
| constexpr int kThreadsPerQuantBlock = kQuantBlock / kELTS_PER_THREAD; | ||
| float amax = 0.f; |
There was a problem hiding this comment.
[Required] Two written-rule issues in the new C++:
- AGENTS.md: "NVIDIA copyright header on ALL new files (update year on modified files)". allReduceFusionKernels.h and customAllReduceKernels.h are meaningfully modified here (new enum values and traits), but line 2 of both still says
Copyright (c) 2022-2024. Every other modified file is already at 2026. - CODING_GUIDELINES.md: "Literal suffixes should be upper case." This epilogue uses
0.f,0xFFFFFFFFu,1e-10f,448.0f,127.f,1.0fand0u(:551-611), while the rest of the file uses forms like448.F.
Please bump both headers to 2022-2026 and use upper-case suffixes (0.F, 1e-10F, 0xFFFFFFFFU, ...).
| self.lm_head = target_model.lm_head | ||
| self.dspark_model.lm_head = target_model.lm_head | ||
|
|
||
| def configure_generation_autotuner(self, max_batch): |
There was a problem hiding this comment.
[Required] CODING_GUIDELINES.md (Static Typing #4): "Always annotate functions. Make the return type None if the function does not return anything." These new functions have no annotations:
- the public wrappers
configure_generation_autotuner(self, max_batch)here andwrite_context_pages(self, main_hidden, positions, mask, pools, block_tables, capacities)at :2863, although the methods they delegate to (:2668, :2710) are fully annotated; - the new mufu16 helpers
_issue_pipelined_score_loads,_finish_pipelined_score_loadsand_prepack_pipelined_tile(fp4_mla_cutedsl_mufu16.py:5550, :5580, :5625); - the new tests (test_autotuner.py:1848, test_moe_backend.py:6106, test_fp4_mla_packed.py:385).
Please copy the signatures from DSv4DSparkDraftModel and add -> None where nothing is returned.
| # plain tensor parallelism (attention_dp off, tp_size > 1) the expert-sharded | ||
| # MoE output must be all-reduced across ranks -- exactly what the target MoE | ||
| # does (enable_allreduce = not (POST_MOE_FUSION or tp_size == 1)). Previously | ||
| # this was hard-coded to False, which dropped the reduction on the non-ADP |
There was a problem hiding this comment.
[Required] This new _forward_stage is a byte-for-byte copy of _forward_stage_legacy (:1985-2132) apart from the attention call. The copy covers the mHC prelude and the 64-line MoE/allreduce tail, and it brings this history narration into new code. AGENTS.md: "Put a change's rationale and history (why it was made, what it replaces, ticket/PR references) in the PR description, not in code comments. Code comments should explain the code as it stands for a future reader."
Beyond the rule, production runs _forward_stage_legacy while the new paged tests run this copy. The next fix to the lagged-pre / fused-HC / MoE-allreduce logic, the logic whose earlier bug silently halved acceptance, can then land in only one of them. Please factor the shared pre/post stage logic into one helper that takes the attention call, and drop the "Previously ..." sentence.
| mutates_args=("output", "packed_scale"), | ||
| device_types="cuda", | ||
| ) | ||
| def cute_dsl_dspark_fp8_bmm_quantize_rubin_out( |
There was a problem hiding this comment.
[Required] This file copies machinery that already exists:
- This op, its fake and the two BMM tuning
ConstraintSpecs are essentially copies oftrtllm::cute_dsl_fp8_bmm_quantize_rubin_outandCuteDSLFp8RubinBmmQuantRunner.tuning_config(cute_dsl_custom_ops.py:13655-13804). Only the tuning plumbing differs. _get_cute_dsl_mxfp8_rubin_alpha(:700) re-implements_get_mxfp8_gemm_alpha(cute_dsl_custom_ops.py:103-126) but drops its CUDA-graph-capture guard.- The R128c4 packed-scale size is derived again in four places (here at :869/:940, dspark_rmsnorm_rope_custom_op.py:772, modeling_dspark.py:4967-4968) instead of using
mxfp8_scale_infer_shape(_torch/utils.py:483).
If the parent layout changes, the DSpark op keeps validating and tuning with the old sizes. Please reuse the existing op, or derive the configs with dataclasses.replace(CuteDSLFp8RubinBmmQuantRunner.tuning_config, ...). Use _get_mxfp8_gemm_alpha, and compute scale sizes in one helper.
| numPackedSfK = ((hiddenDim / 128) + 3) / 4; | ||
| int64_t const tPadded = (numTokens + 3) / 4 * 4; | ||
| constexpr int64_t kTmaAlignedUint32Elems = 4; | ||
| scaleLeadingDim = (tPadded + kTmaAlignedUint32Elems - 1) / kTmaAlignedUint32Elems * kTmaAlignedUint32Elems; |
There was a problem hiding this comment.
[Required] This block re-implements existing quantization pieces instead of reusing them:
- The allocation and the
from_blobMN-major view copy the legacy (use_r128c4_layout=False) branch offp8_quantize_1x128_packed_ue8m0(fp8Quantize.cpp:209-237). - The kernel epilogue copies the 1x128 UE8M0 math already in quantization.cuh (
reciprocal_approximate_ftz,exp2f_rcp,FP8_E4M3_MAX).
Other problems in the same block:
- The comment above says the consumer GEMM "cannot tell this from a standalone-quantized activation". That op now defaults to R128c4 (fp8Quantize.cpp:348), so a consumer written against the default would misread these scales.
- This rounding is a no-op, because
tPaddedis already a multiple of 4. - The launchers duplicate the
tokenDim % 128check (mnnvlAllreduceKernels.cu:822-826, :1535-1539) but never check what the kernel actually needs:sfLeadingDim >= numTokensand% 4 == 0.
There is no caller yet, but please:
- share one allocate/view helper with fp8Quantize;
- call the shared quant helpers from the epilogue;
- say "deep_gemm MN-major (not R128c4)" in the comment;
- check the full
sfLeadingDimprecondition once, with the values in the message.
Fix DFlash precision and workspace sizing, DSpark projection layouts and shared ring metadata, FP4 synchronization and planner overhead, and short autotuner tactic sharding. Reuse existing Linear and quantization dispatch, remove unused draft integration and MNNVL fusion, and add regression and Rubin CI coverage. Keep optional wave simulation disabled by default. Signed-off-by: Mingyang Hao <mingyangh@nvidia.com>
Signed-off-by: Mingyang Hao <mingyangh@nvidia.com>
Shared scratch is overwritten for each output column tile. Mark the inline PTX loads as memory-dependent side effects so the compiler cannot reuse a previous tile's values. The existing FP8 numerical and graph regression now passes with its original tolerance. Validated with 62 passing Rubin GPU tests, including production ring inputs, page-boundary handling, FP8 output, and CUDA graph replay. Signed-off-by: Mingyang Hao <mingyangh@nvidia.com>
Keep folded ratios through128 in a single launch and split larger blocks without changing the attention mask. Size workspace for the largest chunk, validate initialization, and test BF16 success plus FP8 rejection before cache mutation. Signed-off-by: Mingyang Hao <mingyangh@nvidia.com>
Keep production rolling-window attention and its paged-kernel adapter. Remove the unused standalone page interface and exclusive writers, retain numerical and graph coverage, and allow architecture-conditional op monkeypatches. Signed-off-by: Mingyang Hao <mingyangh@nvidia.com>
Signed-off-by: Mingyang Hao <mingyangh@nvidia.com>
| # Matching Q/KV dtypes and q_len=1 select a Q tile up to 128 in TRTLLM-Gen. | ||
| # Keep each folded launch in that range; larger model blocks are split. | ||
| # This is a selected launch range, not a native or model block-size limit. | ||
| DFLASH_FOLDED_DECODE_MAX_RATIO_PER_LAUNCH = 128 |
There was a problem hiding this comment.
🟡 A folded launch runs with group × chunk Q heads per KV head, up to 128 here, while FlashInferTrtllmGenFmha.MAX_HEADS_RATIO_GENERATION is 32 (flashinfer_trtllm_gen.py:353, raised from 16 for Nemotron-H in #12620) and validate_dflash_trtllm_gen_runtime only checks the unfolded group. The added tests mock the workspace sizes on CPU; nothing launches a folded shape above 32 on GPU.
Could you add a B200 case with a folded ratio above 32 (e.g. group 4 × block 16 = 64, Qwen3-8B DFlash b16) compared against the context kernel or a reference? A missing cubin would fail loudly, but the numerics of the 64/128 tiles on this path are not covered yet.
|
/bot run --disable-fail-fast |
|
PR_Github #82461 [ run ] triggered by Bot. Commit: |
Brings in the Rubin merge-back review fixes (cf8309c), the DSpark FP8 scratch-load ordering fix (ba5a32d) and main up to 5865809. Conflicts: - MNNVL all-reduce fusion: drop RMS_NORM_QUANT_FP8_BLOCKSCALE as the PR does (nothing calls it) and keep our RMS_NORM_QUANT_MXFP8, renumbered into the freed slot (AllReduceFusionOp 10, kARRMSNormMXFP8 7, kMXFP8 3) in customAllReduceKernels.h, allReduceFusionKernels.h, functional.py, mnnvlAllreduceKernels.cu, allreduceOp.cpp and distributed/ops.py. - modeling_dspark.py: keep _helix_post_process; write_dspark_context now comes from attention/kernels/dspark.py (moved by the PR). - dflash.py: keep HelixDraftShard; DFLASH_CONTEXT_LOWERING is gone. - test_fp4_mla_packed.py: keep both sides' plan cases and tests. Signed-off-by: Zhenhuan Chen <zhenhuanc@nvidia.com>
Merge origin/main 745c1ce. The 50 main commits up to 5865809 came in through the NVIDIA#21400 and NVIDIA#21067 updates; the last 8 merge without conflicts. Signed-off-by: Zhenhuan Chen <zhenhuanc@nvidia.com>
|
PR_Github #82461 [ run ] completed with state
|
sunnyqgg
left a comment
sunnyqgg
left a comment
There was a problem hiding this comment.
Follow-up to my review of 2026-11-03: 16 of the 21 lower-priority points I left out of it still apply at c8e67f3. The other 5 are fixed or no longer apply. They are marked [Nice to have] and none blocks the merge on its own.
| def dequant(name: str) -> torch.Tensor: | ||
| weight = weights[f"{prefix}{name}.weight"].to(device) | ||
| scale = weights[f"{prefix}{name}.scale"].to(device) | ||
| return self._block_dequant(weight, scale) |
There was a problem hiding this comment.
[Nice to have] _bind_bf16_attn_projections re-implements _cache_attn_weights (same per-stage loop and dequant) but calls self._block_dequant(weight, scale) with the default 128 block (modeling_dspark.py:2158), dropping the block=self.checkpoint_fp8_block_size that _cache_attn_weights passes (:1427). That is right for V4 today, but V4.1 (checkpoint_fp8_block_size = 32) would decode with the wrong scales if it moved to post_load_weights(weights). Keep one BF16 binder (reuse _cache_attn_weights and reset the FP8 keys), and update the stale cache_attn_weights_from_state_dict docstring (:1466-1467) and the finished TODO(step 3) (:1447).
| attn = stage._dspark_attn | ||
| if attn.get("qkv_proj") is not None: | ||
| head_dim = int(self._attn_params["head_dim"]) | ||
| return attn["qkv_proj"](main_x)[..., -head_dim:] |
There was a problem hiding this comment.
[Nice to have] The context KV written to the rolling window is sliced differently from the main_kv that attention uses: _qkv_projections takes main_qkv[..., q_lora_rank:] (modeling_dspark.py:4347), while the window write takes [..., -head_dim:] (:2277), forward_slice(q_lora_rank, q_lora_rank + head_dim) (:2223-2224) or the trailing head_dim weight rows of the stacked GEMM (:4518). They agree only while the fused projection is exactly q_lora_rank + head_dim wide; otherwise the draft silently attends to the wrong context. Derive all of them from one slice, or assert that width once in _bind_attn_projections.
| _DSPARK_NUM_HEADS = 128 | ||
| _DSPARK_HEAD_DIM = 512 | ||
| _DSPARK_WINDOW_SIZE = 128 | ||
| _DSPARK_PAGE_SIZE = 8 |
There was a problem hiding this comment.
[Nice to have] The six _DSPARK_* constants added at dspark_attention_custom_op.py:29-34 are unused (the op reads the _DSV4_DSPARK_* set at :54-59), and _DSPARK_PAGE_SIZE = 8 contradicts the 16–256-token pages the op accepts (:190), so code importing it gets a value the op rejects. The support-gate message also says 'q must have 128 heads' (:506) although 64 heads are accepted. Delete the unused constants and change the message to '64 or 128 heads'.
| _DSPARK_LINEAR_BACKEND_ENV = "TRTLLM_DSPARK_LINEAR_BACKEND" | ||
|
|
||
| if IS_CUTLASS_DSL_AVAILABLE: | ||
| from ..custom_ops import dspark_attention_custom_op # noqa: F401 |
There was a problem hiding this comment.
[Nice to have] The two back-to-back if IS_CUTLASS_DSL_AVAILABLE: blocks (modeling_dspark.py:140-148 and :151-159) import from the same two modules, and from ..custom_ops import dspark_attention_custom_op # noqa: F401 (:141) is redundant because :152 already imports that module, which registers its ops. Merge them into one guarded block and drop the side-effect import together with its noqa.
| 2, | ||
| ): | ||
| raise ValueError("DSpark inverse RoPE frequencies must be FP32 [batch, block, 32, 2]") | ||
| if kv_pages.stride(1) != dim or kv_pages.stride(2) != 1 or kv_pages.stride(0) % 8: |
There was a problem hiding this comment.
[Nice to have] _run_paged_dspark_attention lets kv_pages.stride(0) == 0 through (dspark_attention_custom_op.py:219 only checks stride(0) % 8), so an expand()ed or as_strided pool passed directly to trtllm::fused_dsv4_dspark_attention makes every page alias page 0 and returns wrong attention instead of raising. The model path is safe only because is_fused_dsv4_dspark_attention_supported rejects stride(0) <= 0 (:516); the public op should reject it too. Add kv_pages.stride(0) <= 0 to this check.
| # TMEM allocation is published to the producers and the PV issue warp. | ||
| TMEM_BAR_THREADS = 32 * (2 + SMEM_P4_TOTAL_PRODUCER_WARPS) | ||
| O_STORE_BAR_ID = 2 | ||
| SMEM_P4_P4_SLOT0_BAR_ID = 6 |
There was a problem hiding this comment.
[Nice to have] Parts of the removed score-slot/row-sum protocol remain: SMEM_P4_CONTROL_MBAR_COUNT budgets SMEM_P4_TMEM_SCORE_PIPELINE_STAGES * 2 barriers for tmem_score_slot0_mbars, which is no longer allocated (fp4_mla_cutedsl_mufu16.py:615); the P4_READY and HALF1_ROWSUM_READY barriers are initialized (:6583-6594) but used only by the dead _runtime_t336_producer_tile (:5193); and that dead function alone keeps _store_smem_p4_n256_rowsum_part (:3573) and SMEM_P4_P4_SLOT{0,1}_BAR_ID (:926-927) alive. Delete the dead function with its helpers and barrier ids and shrink the barrier budget, so readers can tell which protocol is live.
| # QK and PV B scale factors use disjoint TMEM columns: a ring holding one QK data stage of K scales | ||
| # (4 kblocks) and one PV N tile of V scales (4 kblocks). Each issuer reuses its | ||
| # columns in issue order. | ||
| PV_ISSUE_WARP_ID = ROWMETA_WARP_ID |
There was a problem hiding this comment.
[Nice to have] Several identifiers no longer match what they hold: PV_ISSUE_WARP_ID = ROWMETA_WARP_ID (fp4_mla_cutedsl_mufu16.py:527); SMEM_P4_PAIR_SCORE_CONSUMED_MBAR_OFFSET is a third alias of the ROWSUM_DONE/DIRECT_P_READY slots (:448-451); sScoreSlotRowMax stores per-warp row-sum parts (:7379-7391); _compute_grid(o, ...) receives q_latent (attention_kernel.py:537), and page_size_win is the window length next to the new history_page_size; _DSV4_DSPARK_NUM_HEADS is a tuple (dspark_attention_custom_op.py:54). Rename them after their contents (e.g. sRowSumParts, _SUPPORTED_NUM_HEADS, window_len) or drop the aliases.
| context window and the draft block's own rows); see | ||
| ``attention.py`` for the supported DSpark contract and | ||
| ``tests/unittest/_torch/speculative/test_dspark_cute_dsl_attention.py`` for | ||
| ``tests/unittest/_torch/speculative/test_dspark_paged_attention.py`` for |
There was a problem hiding this comment.
[Nice to have] The kernel docstring points to tests/unittest/_torch/speculative/test_dspark_paged_attention.py for validation coverage (attention_kernel.py:61), but that file doesn't exist; the coverage is in test_dspark_cute_dsl_attention.py, which the docstring named before. Also, the comment at test_fp4_mla_packed.py:155 ('as before the wave simulation') describes history rather than the code; say what the default plan does (rule-based task length, request order, no wave simulation).
| ) -> None: | ||
| tiles_per_task, begins = _plan_mod().plan_packed_tasks(tiles, capacity, wave_tasks) | ||
| tiles_per_task, begins = _plan_mod().plan_packed_tasks( | ||
| tiles, capacity, wave_tasks, wave_plan=True |
There was a problem hiding this comment.
[Nice to have] With the wave plan opt-in, wave_plan=False is the production default, yet test_plan_packed_tasks_balances_by_kv_length only checks wave_plan=True (test_fp4_mla_packed.py:97-99), and the default plan is pinned for just two batches (:156-160). The fixed results this test had for [128] * 15 + [200], [128, 128, 128, 1024] and [8] are no longer asserted anywhere. Parametrize the balances test over wave_plan with both expected tuples so a change to the default plan fails a test.
| ) | ||
| SMEM_P4_V_PIPELINE_STAGES = 3 | ||
| # Poll period of the K/V load warps while their ring slot is still in use. | ||
| SMEM_P4_LOAD_WAIT_SLEEP_NS = 128 |
There was a problem hiding this comment.
[Nice to have] The wave planner is opt-in, but the pipelined leading-tile producer, the QK issue two tiles ahead and the new 128 ns load-wait poll (SMEM_P4_LOAD_WAIT_SLEEP_NS, fp4_mla_cutedsl_mufu16.py:460) change every packed FP4 MLA decode by default, and I couldn't find latency numbers for them. Please post SM107 decode latency before/after for representative batch and KV-length shapes (the PR description is fine) and note on the constant where it was tuned.
Dev Engineer Review
QA Engineer Review
unittest/disaggregated/test_extractor.pyin bothl0_a10.ymlandl0_h100.yml. It found no exact matching entries for the other searched changed tests.l0_dgx_b200.ymllists a different autotuner test,test_autotuner_distributed_strategy. No matching manual-QA entries were found.Per-File QA Perspective
Source files
cpp/tensorrt_llm/batch_manager/kv_cache_manager_v2/common.h— VerifyCacheDomainvalues and the documented storage order.cpp/tensorrt_llm/batch_manager/kv_cache_manager_v2/config.h— Verify thatAttentionLayerConfigdefaults to the target cache domain.cpp/tensorrt_llm/batch_manager/kv_cache_manager_v2/lifeCycleRegistry.cpp— Verify lifecycle construction forwards the configured cache domain.cpp/tensorrt_llm/batch_manager/kv_cache_manager_v2/lifeCycleRegistry.h— Verify domain-aware comparison and scratch-range behavior.cpp/tensorrt_llm/batch_manager/kv_cache_manager_v2/storage/config.cpp— Verify pool grouping separates cache domains while preserving intended reuse.cpp/tensorrt_llm/kernels/communicationKernels/allReduceFusionKernels.h— Verify the new fusion and quantization enum values match dispatch consumers.cpp/tensorrt_llm/kernels/communicationKernels/mnnvlAllreduceKernels.cu— Verify FP8 block quantization, scale packing, padding, and launch validation.cpp/tensorrt_llm/kernels/communicationKernels/mnnvlAllreduceKernels.h— Verify callers setsfLeadingDimcorrectly.cpp/tensorrt_llm/kernels/customAllReduceKernels.h— Verify the fusion operation name and value match kernel and Python dispatch.cpp/tensorrt_llm/kernels/helixAllToAll.cu— Verify the asynchronous shared-memory fence protects FIFO data before reuse.cpp/tensorrt_llm/nanobind/batch_manager/kvCacheManagerV2.cpp— Verify binding defaults and properties match the C++ API.cpp/tensorrt_llm/thop/allreduceOp.cpp— Verify output padding and the returned scale-view shape.tensorrt_llm/_torch/attention/backends/dspark.py— Verify captured-context writes respect masks, capacities, and page-table bounds.tensorrt_llm/_torch/attention/backends/fp4_mla/cache_manager.py— Verify draft-layer dtype and cache-size accounting across Helix and non-Helix configurations.tensorrt_llm/_torch/attention/backends/fp4_mla/config.py— Verify the wave-plan environment setting and default.tensorrt_llm/_torch/attention/backends/fp4_mla/fp4_mla_cutedsl_mufu16.py— Verify packed-task bounds, split-output merging, and pipeline synchronization.tensorrt_llm/_torch/attention/backends/fp4_mla/fp4_mla_packed_reduce.py— Verify candidate limits, makespan selection, tie-breaking, and disabled-planner behavior.tensorrt_llm/_torch/attention/backends/fp4_mla/packed_plan.py— Verify environment settings reach the planner launch.tensorrt_llm/_torch/attention/backends/sparse/deepseek_v4/cache_manager.py— Verify draft-cache quota and fixed-reservation calculations.tensorrt_llm/_torch/attention/backends/sparse/dsa/cache_manager.py— Verify local-layer pool sizing and draft-cache accounting.tensorrt_llm/_torch/attention/kernels/blackwell/dspark/attention.py— Verify paged-attention arguments and supported geometry match callers.tensorrt_llm/_torch/attention/kernels/blackwell/dspark/attention_kernel.py— Verify history-page masking, capacity handling, and FP8 output scales.tensorrt_llm/_torch/attention/kernels/blackwell/dspark_rmsnorm_rope.py— Verify quantization scales and bounds checks on managed page writes.tensorrt_llm/_torch/attention/mla.py— Verify chunked-prefill dispatch and Helix V-projection behavior against fallback paths.tensorrt_llm/_torch/autotuner.py— Verify the full tactic list is used when it is shorter than the tuning group.tensorrt_llm/_torch/custom_ops/dspark_attention_custom_op.py— Verify custom-op validation, kernel caching, and the removed cache/slot interface.tensorrt_llm/_torch/custom_ops/dspark_rmsnorm_rope_custom_op.py— Verify device-specific compilation, page-write validation, and FP8 precompile gating.tensorrt_llm/_torch/disaggregation/native/auxiliary.py— Verify draft-history serialization and recycled-slot clearing.tensorrt_llm/_torch/disaggregation/native/bounce/impl.py— Verify byte accounting forcp_as_tppaged groups.tensorrt_llm/_torch/disaggregation/native/mixers/attention/peer.py— Verify peer checks reject incompatible overlapping attention groups.tensorrt_llm/_torch/disaggregation/native/peer.py— Verify draft-layer compatibility and rank mapping forcp_as_tp.tensorrt_llm/_torch/disaggregation/native/rank_info.py— Verify the draft-layer count reflects the active cache manager.tensorrt_llm/_torch/disaggregation/native/transfer.py— Verify auxiliary-slot requirements and draft-history transfer behavior.tensorrt_llm/_torch/disaggregation/resource/cache_reuse.py— Verify block-ordinal conversion for sharedcp_as_tppages.tensorrt_llm/_torch/disaggregation/resource/kv_extractor.py— Verify draft-page expansion and pool separation in extracted page tables.tensorrt_llm/_torch/disaggregation/resource/page.py— Verify deserialization remains compatible when the new group fields are absent.tensorrt_llm/_torch/disaggregation/transceiver.py— Verify draft-history validation and restoration.tensorrt_llm/_torch/distributed/ops.py— Verify the new fusion operation reaches supported MNNVL dispatch.tensorrt_llm/_torch/models/modeling_dflash.py— Verify FP8 page writes and block-query layout conversion.tensorrt_llm/_torch/models/modeling_dspark.py— Verify managed-page and legacy paths, projection fallback, and required paged metadata.tensorrt_llm/_torch/models/modeling_kimi_linear.py— Verify the Helix V-projection weight path when enabled.tensorrt_llm/_torch/moe/fused_moe/cutedsl/fused_fc12.py— Verify the input-ready event is joined before fused execution.tensorrt_llm/_torch/moe/fused_moe/fused_moe_cute_dsl_b12x.py— Verify event ordering before either execution path.tensorrt_llm/_torch/pyexecutor/_util.py— Verify unified draft-cache selection and rejection of unsupported manager combinations.tensorrt_llm/_torch/pyexecutor/engine/runners/decoder/runner.py— Verify draft-cache preparation and history snapshots occur only in applicable decode steps.tensorrt_llm/_torch/pyexecutor/kv_cache/kv_cache_manager_v2.py— Verify draft allocation, sizing, history lifecycle, and target/draft pool ratios.tensorrt_llm/_torch/pyexecutor/kv_cache/mamba_cache_manager/common.py— Verify draft storage is included in hybrid-cache cost estimates.tensorrt_llm/_torch/pyexecutor/kv_cache/mamba_cache_manager/mamba_cache_manager_v2.py— Verify draft pools use base-manager roles and free-block counts.tensorrt_llm/_torch/pyexecutor/py_executor.py— Verify pending history updates publish after sampler updates and handle absent state.tensorrt_llm/_torch/pyexecutor/sampler/sampler.py— Verify the optional update field remains compatible with sample-state consumers.tensorrt_llm/_torch/pyexecutor/standalone_draft_cache.py— Verify history validation and stale-cache protection during publication.tensorrt_llm/_torch/speculative/dflash.py— Verify unified-cache preparation, page-table fallback, and overflow handling.tensorrt_llm/_torch/speculative/dspark.py— Verify managed-page allocation, context seeding, and progress-slot handling.tensorrt_llm/_torch/speculative/interface.py— Verify managed-cache hooks and padding-slot validation across worker implementations.tensorrt_llm/functional.py— Verify residual-input validation permits only the intended new fusion.tensorrt_llm/runtime/kv_cache_manager_v2/__init__.py— Verify the publicCacheDomainexport.tensorrt_llm/runtime/kv_cache_manager_v2/__init__.pyi— Verify stub signatures and defaults match the bound API.Test files
cpp/tests/unit_tests/batch_manager/kvCacheManagerV2ConfigTest.cpp— Covers cache-domain defaults, lifecycle comparison, and pool grouping. No test-list entry was established.tests/unittest/_torch/attention/multi_gpu/test_mla_helix.py— Covers V-up projection placement and Helix all-to-all behavior in CUDA graphs. No matching test-list entry was found.tests/unittest/_torch/attention/test_fp4_mla_packed.py— Covers planner wave simulation, ordering, and disabled-planning behavior. No matching test-list entry was found.tests/unittest/_torch/executor/kv_cache/test_kv_cache_estimation.py— Covers standalone draft-cache sizing and dtype selection. No matching test-list entry was found.tests/unittest/_torch/misc/test_autotuner.py— Covers short tactic lists across Helix CP ranks. The list search found a different autotuner test, not this test by name.tests/unittest/_torch/speculative/hw_agnostic/test_dflash_worker.py— Covers managed-cache block-table decoding and FP8 context-page writes. No matching test-list entry was found.tests/unittest/_torch/speculative/hw_agnostic/test_dspark_attention.py— Covers projection paths, backend selection, and page writes. No matching test-list entry was found.tests/unittest/_torch/speculative/hw_agnostic/test_dspark_cuda_graph.py— The reported change removes comments only; test behavior is unchanged. No matching test-list entry was found.tests/unittest/_torch/speculative/hw_agnostic/test_dspark_worker.py— Covers managed draft pages, request slots, context seeding, and generation progress. No matching test-list entry was found.tests/unittest/_torch/speculative/hw_agnostic/test_kimi_k3_dspark_semantics.py— Covers paged decode parity for BF16 and FP8 histories. No matching test-list entry was found.tests/unittest/disaggregated/region/test_aux.py— Covers draft-history serialization round trips. No matching test-list entry was found.tests/unittest/disaggregated/test_extractor.py— Updates the fake-manager fixture. Listed inl0_a10.ymlandl0_h100.yml.Description
Test Coverage
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-compatibleorapi-breaking. Forapi-breaking, includeBREAKINGin 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.