Repository navigation
[None][fix] Kimi K3 / KDA / MNNVL / DFlash: fixes for defects on main - #19988
Merged
Merged
Conversation
…tputs maybe_execute_in_parallel runs fn1 on aux_stream and returns its outputs to the caller, which uses them on its own stream after the event join. Those tensors are allocated in aux_stream order. When the caller drops them, the caching allocator returns their blocks to aux_stream's pool at once and cudaMallocAsync frees them in aux_stream order, so later aux_stream work can take and overwrite the memory while the calling stream's reads are still queued. compute-sanitizer memcheck with cudaMallocAsync reports this as use-after-free on the outputs Kimi K3's KDA mixer and MoE shared expert return this way. The CUDA tensors fn1 returns, directly or in nested tuples, lists and dict values, are now recorded on the calling stream after the join. cudaMallocAsync raises in record_stream on memory it did not allocate (a buffer from outside PyTorch, wrapped as a tensor); only that error is caught, and the tensor is skipped, as the native allocator does. Nothing is recorded under CUDA graph capture or torch.compile tracing. In a capture the caching allocator would hold every recorded block freed during the capture until it ends, so each call's outputs would take new graph pool memory, and under cudaMallocAsync freeing a recorded output inside the capture fails it with "capturing stream has unjoined work". A captured graph relies instead on later aux_stream work waiting on the calling stream, as the next call's event0.wait() does. Dynamo traces neither the capture query nor record_stream, so a compiled caller keeps the previous behavior and still traces fullgraph. test_multi_stream_utils.py reads fn1's output on the calling stream behind a sleep, drops it, and fills a new aux_stream allocation of the same size at once; the read must still see fn1's values. Before, it saw the fill under both allocators. The file also covers a buffer from cudaMalloc, a capture whose outputs outlive it, outputs dropped inside a capture, and a fullgraph torch.compile of a caller. It reruns itself under PYTORCH_CUDA_ALLOC_CONF=backend:cudaMallocAsync and joins the l0_h100 and l0_b200 pre-merge module tests. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…ive for captured graphs Growing the MNNVL all-reduce workspace replaced it and freed the previous buffers, pointer tables and flags while CUDA graphs captured before the growth still launched on them. The replaced workspaces now stay alive with the process, and creating or growing a workspace during graph capture raises instead of failing inside the driver. test_mnnvl_workspace_growth_keeps_captured_graphs captures a one-token all-reduce, grows the workspace with an eager two-shot call, replays the graph and checks growth under capture is refused. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…lti-GPU list test_mnnvl_workspace_growth_keeps_captured_graphs needs fabric-backed MNNVL on two GB200 GPUs, like the checkpoint graph tests next to it. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
The workspace-growth test came from a branch whose yapf settings differ from main's; format it as main's pre-commit does. No behaviour change. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…re guard Creating or growing the MNNVL all-reduce workspace under CUDA-graph capture raises before the collective allocation. Its only test was the last block of the 2-rank growth test, which runs on GB200 multi-GPU stages only. test_mnnvl_workspace_creation_refuses_graph_capture creates the workspace for a one-rank group inside a capture and expects the guard's RuntimeError. The guard runs before any MNNVL call, so one GPU suffices; without it the call enters the workspace construction under capture and fails inside it with another error. The test joins l0_b200's single-GPU PyTorch block. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
… kernels oneshotAllreduceFusionKernel and rmsNormLamport used LamportFlags' cluster arrival, in which every thread arrives on the cluster barrier but only cluster rank 0's first warp waits, and then arrived on the cluster barrier again in the RMSNorm epilogue's cluster reduction. Each thread must arrive and wait once per phase: a second arrival counted toward the first phase lets the reduction read a peer CTA's partial sum before it is written, and counts the cluster in before a late CTA has read the buffer flags. Both kernels now count arrivals per CTA, as the attn_res kernel already does; their cluster barrier is left to the reduction. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…luster reduction writes The RESIDUAL_RMS_NORM epilogues of oneshotAllreduceFusionKernel and rmsNormLamport write each CTA's partial sum into the other CTAs' shared memory through the cluster mapping before any cluster barrier wait, so a target CTA may not have started yet; racecheck reports every such write once the kernels' arrival counting no longer stalls it. A cluster.sync() before the writes makes every CTA of the cluster start first, as the attn_res kernel's early arrive / wait does. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…ads before the next chunk's writes attn_res_fwd_online_v2_kernel reuses one ws_stats buffer every chunk of candidates: lane 0 of each consumer warp writes its row, a named barrier orders the writes before every consumer's cross-warp reads, but nothing ordered those reads before the next chunk's writes. With more than one chunk (N > 4 at H 7168) a warp that finished reading could overwrite its row while a slower warp still read it, mixing the next chunk's statistics into this chunk's logits. Under compute-sanitizer racecheck, which serializes warps, every such call returned a wrong mixture (relative error 0.6-2.1 of max |ref|); natively it was not observed. Add the second barrier the persistent fork already has. Also order every lane's reads of a chunk's slots before lane 0 releases them to the producer (__syncwarp before the bar_consumed arrive), in online_v2 and in the persistent fork. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…eases a TMA-filled slot The online_v2, N = 1 tile and persistent fused kernels read their V (and delta) slots with generic-proxy shared loads and release them with an mbarrier arrive; the producer then refills them with cp.async.bulk (async proxy). Accesses to one location through two proxies need a cross-proxy fence: fence.proxy.async.shared::cta in every reading thread before the release. In the current SASS each release already issues after an instruction that consumes the last load, so this makes that ordering independent of instruction scheduling. Outputs are bit-identical. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…h slot release test_attn_res_proxy_fences.py reads attnResFwd.cu and attnResFwdPersistentFused.cu and checks every consumer release of a cp.async.bulk-filled slot (an arrive on bar_consumed). Each must follow fence.proxy.async.shared::cta, with only the warp sync and the lane-0 guard between the two. The fence orders the consumers' generic-proxy reads of the slot before the producer's async-proxy refill. Without it the order rests on instruction scheduling, which no test of values can see. The check fails on the sources from before the fence was added (four releases) and passes on the current ones. CPU only. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
CUDA-graph padding rows and warmup dummies share the drafter's dummy slot. Their accepted tokens grow its context length like any request's, but nothing reset it, while their page-table rows are the padding request's pages with every other entry mapped to page 0, another request's page. Once the dummy length passed the padding request's own pages, the padding rows' context K / V (k3_ctx_kv, or the torch path) landed on that page: a live request's drafter context overwritten on every padded step, costing acceptance. prepare() now writes 0 to the dummy slot every step, with the evicted slots. test_dflash_dummy_slot.py: padded steps keep the dummy slot at 0 and the real slots untouched; an evicted request and the dummy reset in one step. It fails without the reset. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
… mxfp8_quantize and noaux_tc_op invokeFP4Quantization and invokeMxFP8Quantization launch one block per row (up to an occupancy cap), invokeNoAuxTc one block per token. With no rows the grid has zero blocks: the launch fails with cudaErrorInvalidValue, nothing checks it, and the error stays pending until the next launch that reads the last error raises it, usually an unrelated PyTorch kernel. The three launchers now return before the launch when there are no rows (there is nothing to write; the swizzled scale-factor layout pads zero rows to zero), and their cudaLaunchKernelEx calls are checked. noaux_tc_op still rejects an unsupported configuration at zero tokens. That pending error is what failed the three modeling_v2 catalog tests that run right after a zero-row call in their files; their waivers are removed. Tests: zero-row fp4_quantize (bf16 and fp8 input, both scale-factor layouts) and mxfp8_quantize (linear and swizzled), and zero-token noaux_tc_op (single and multi group), each followed by a normal call that must match the one before it. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
trtllm::kda_prefill returned a view of a buffer it reuses: the output was one of the intermediate tensors _get_buffers caches per input shape, so every call with the same shapes wrote the same storage, and a caller's output changed under it at the next such call. The op's fake declares a new tensor (v.new_empty(v.shape)), so nothing orders the readers of one call's output before the next call. Allocate the output on every call, with the varlen path's one chunk of slack past the last token like the cached buffers, and wrap it for K4 there; the cache no longer holds it. Test: a second call with the same shapes and other inputs must leave the first call's output unchanged (varlen, and one padded equal-length row). Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
The KDA layers (tensorrt_llm/_torch/modules/kimi_kda) import flash-linear-attention's ShortConvolution when they load and run its chunked KDA for prefill batches below four 64-token chunks, but flash-linear-attention is listed only in requirements-dev.txt. CI installs that file, so every test passes, while a plain pip install of the wheel, and the release container, which installs the wheel with its runtime requirements only, cannot load either model: No module named 'fla'. Require fla-core, the distribution that provides fla.modules and fla.ops, at the version requirements-dev.txt already installs through flash-linear-attention 0.5.2. Its only base dependency, einops, is a runtime requirement already. A CPU test checks that every package the KDA layers import when they load is a runtime requirement. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…ture used it Growth retired every replaced workspace for the life of the process, including ones no CUDA graph had captured. Growth rounds to 8 MiB, so a run of rising eager sizes could keep many dead workspaces. The getter now marks a workspace it hands out during capture, and a growth keeps the previous workspace only when it is marked. An unmarked one is released at the replacement, after every rank's device synchronize in the protocol reset. The growth test checks that a second growth, past a workspace no capture used, leaves only the captured one in the retired list and that the graph still replays. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…ken embedding GQADSparkForCausalLM takes the target's embedding in load_weights_from_target_model, so the DFlash worker built every masked block slot from the target's row for the mask token, which the target never trained. A drafter checkpoint that ships its own embedding has trained that row. load_weights now keeps it as mask_token_embedding, and the worker's noise block (dflash_noise_block_embedding, the same computation moved into a function) uses it for every masked slot when the drafter has one for the mask id in use; slot 0 stays the bonus token's shared lookup. Every rank still embeds the mask id through embed_tokens.forward, as before. Drafts change; under greedy verification the target's tokens do not. Plain DFlash and MLA DSpark drafters set no mask_token_embedding and are unchanged. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…y state MambaHybridCacheManagerV2 keeps the KDA replay state (prev_num_accepted_tokens and the six replay caches) per SSM slot, outside the V2 pools, and remembers each request's last slot until free_resources. A V2 suspend (the scheduler's eviction and self-eviction, the pool rebalance) only suspends the pages. Under pressure the host tier takes them, the slot goes to another request, and the suspended request resumes in whichever slot is free. Its next _setup_state_indices then relocated the old slot's contents, by then the other request's replay state, or, back in its old slot, kept them: the next verify replayed another request's drafts. suspend_request now copies the slot's replay state to host memory and forgets the slot; _setup_state_indices writes it into the slot the request resumes in. update_resources gives the host drafter's acceptance of a request suspended after its batch was scheduled to the saved state instead of the old slot. Relocation and the copy share one list of the per-slot buffers. The new tests resume a request in a new slot and in its old one and record a late acceptance on a manager stub, and on the V2 runtime let new caches take a suspended request's slot through the host tier before it resumes. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…KDA replay caches The KDA replay caches and prev_num_accepted_tokens are allocated once, one entry per slot the SSM pool keeps at construction (get_page_index_upper_bound). With enable_kv_pool_rebalance, adjust() resizes the pool groups to the auto-tuner's target ratio, which gives the SSM pool one slot per typical request, so a pool built smaller than that grows past the buffers. A request in a slot past them makes the prev_num_accepted_tokens[slot] = 0 reset fail, and the KDA verify kernels index the replay entries by slot without a bound. Nothing refused the combination: _util.py refuses a rebalance only with block reuse, and the executor's rebalance gate skips only two-model drafters, so one-model speculative decoding on KDA layers reaches it. _build_cache_config now raises ValueError for the KDA replay path with enable_kv_pool_rebalance, as KvCacheConfig does for fp8_ds_mla's fixed pool views. test_v2_kda_replay_refuses_kv_pool_rebalance covers the refusal and the control without the replay caches. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
_get_buffers keeps the prefill scratch for at most _BUF_CACHE_MAX_ENTRIES token-count shapes and evicts the oldest entry on a new shape. The K123 and akk_inv launches wrapped that scratch with _ct_cached, whose id-keyed _input_wrap_cache holds each wrapper until the wrapped tensor is garbage collected. The wrapper's DLPack capsule keeps the tensor alive, so that never happens: every evicted entry's scratch stayed allocated. A server whose prefill batches keep changing token count grows by one entry's scratch per eviction (GSM8K on Kimi K3: about 500 MiB per 200 requests on every rank). Cache the scratch wrappers in the scratch's own _buf_cache entry instead (_scratch_ct), so evicting the entry releases the scratch and its wrappers together; the per-call from_dlpack savings stay. _ct_cached and _input_wrap_cache have no other users and are removed. Test: test_evicted_prefill_scratch_is_released runs cyclic passes over more token-count shapes than the cache holds. After every later pass the cache holds at most _BUF_CACHE_MAX_ENTRIES entries, the allocated bytes are back at their value after the first pass, and the tensors of an entry cached at the end of the first pass, which the next pass evicts, are freed. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…_parallel forward_with_base recorded the calling stream on each LoRA output of the overlapped path. maybe_execute_in_parallel now records fn1's outputs itself, outside CUDA graph capture and torch.compile tracing, so there the LoRA record only repeated it. The overlapped path runs only inside the CUDA graph runner, which turns multi-stream on for its eager warm-up steps and for the capture. During the capture the LoRA record still ran, which is what the helper avoids: the caching allocator holds a recorded block freed in the capture until the capture ends, so each overlapped LoRA call took new graph pool memory, and under cudaMallocAsync a recorded output freed inside the capture breaks it. maybe_execute_in_parallel is now the only place the rule is applied. lora_aux_stream is used only through it, so in a capture each overlapped call's fn1 waits on the calling stream (event0) before it allocates, after the previous call's outputs were read. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…s source Two notes in maybe_execute_in_parallel, no behavior change: - The docstring says that callers should not record fn1's outputs themselves, and that a caller torch.compile traces with disable_on_compile=False runs fn1 on aux_stream without the record, outside CUDA graphs too, so it must order the reuse of fn1's outputs itself. - The comment on the caught cudaMallocAsync error names where its text comes from: the TORCH_INTERNAL_ASSERT in CudaMallocAsyncAllocator::recordStream (c10/cuda/CUDAMallocAsyncAllocator.cpp, torch 2.14.0a0+4fdf77b940.nv26.08), to be re-checked on torch upgrades. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
- test_aux_outputs_dropped_inside_graph_capture asserted that every output in the capture reused one block. That is the native allocator's choice today, not this code's contract. It now asserts fewer blocks than outputs: recording in the capture would hold every output's block until the capture ends. - test_aux_output_over_memory_torch_did_not_allocate imports cuda.bindings instead of importorskip-ing it. cuda-python is a TensorRT-LLM requirement (requirements.txt), so the check of the cudaMallocAsync error text always runs. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…ts instead of a second barrier attn_res_fwd_online_v2_kernel ordered each chunk's cross-warp ws_stats reads before the next chunk's writes with a second NamedBarrier per chunk. Instead, double-buffer ws_stats by chunk parity: chunk gci writes and reads ws_stats[gci & 1]. A warp writes chunk gci + 2 into the same buffer only after it passes chunk gci + 1's barrier, which every warp reaches only after reading chunk gci, so the reads need no barrier of their own. gci keeps counting across token blocks, so the parity holds there too. This removes one bar.sync per chunk for CONSUMER_WARPS * NC * 8 bytes of shared memory. The N == 1 tile kernel shares FwdSmemPlan; it uses buffer 0 and keeps its own second barrier. The __syncwarp() before the lane-0 bar_consumed arrive and the cross-proxy fences stay. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…next chunk fused_kernel123's K1 warps read their scanned prefix row and the chunk total from sPartialLast after the scan barrier, and the next chunk's partial sums overwrite sPartialLast before any K1 barrier (racecheck: write at fused_k123.py:1268 against reads at :1293 / :1299). A K1 warp that reached its reads late used the next chunk's sums: with one K1 warp stalled 20 us after the scan barrier, every case of test_fused_prefill_beta_sigmoid_matches_unfused_kernel changed (output max |d| up to 7e14). One K1 barrier after the reads orders them; outputs are bit-identical to before. test_fused_prefill_does_not_depend_on_k1_warp_timing runs a copy of the kernel with that stall and requires the unpatched kernel's bits; it fails without the barrier. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
The dummy slot's context length grows only where the drafting step adds the step's accepted tokens. Emptying it right after that add, with a zero_() on the fixed slot that a captured step replays, keeps it at 0 between steps without host work. prepare() again writes slot lengths only for evicted requests, instead of two pinned tensors, two H2D copies and an index_copy_ on every step. Dummy rows also attend no context in the step itself, rather than that step's accepted tokens. test_dflash_dummy_slot.py drives the drafting-step update (_advance_ctx_len), including a captured CUDA graph replayed over a dirtied dummy slot, and checks that prepare() writes only evicted slots. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
The fused KDA verify leaves the state after the last golden token in the SSM pool, the drafts the sampler accepted pending in per-slot replay caches, and the conv window only in those caches. commit_kda_pending_drafts moves the drafts and the window into the live pools of one layer, so that the pools hold exactly what the next verify would start from: - kda_decode_mtp_kernel gains a REPLAY_ONLY compile-time mode that runs only the replay steps of a verify and writes the state after them. The replay is the verify's own code, so the result is bit-identical to the state the next verify reaches before its golden token. The mode reads no token inputs or weights; kda_mtp_commit_pending_drafts launches it. - With per-token verify states (kda_state_tok), the accepted drafts are replayed from the per-draft records the Kimi K3 verify kernels keep there (row innovations, beta * k and the decay), with the kernels' arithmetic (replay_kda_token_records). - store_kda_replay_conv_window stores the W - 1 conv columns starting at the accepted count into the conv pool. They are raw bf16 inputs, so the bf16 pool holds them exactly. The default kernel variant is unchanged: the new code is behind const_expr gates and the new parameter is the last one. The new tests check that committing the drafts, storing the window in a bf16 pool and seeding the replay caches from it before a verify gives that verify's outputs, state, caches and windows bit for bit, and the replay of the per-token records. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…lay state into V2 A suspend now leaves the request's whole KDA state in its V2 SSM and conv states instead of a host copy of the replay caches. Before the pages are released, suspend_request commits the slot's pending accepted drafts and its conv window into them (commit_kda_pending_drafts) and clears the draft count; the replay caches hold nothing the request needs any more, and no extra memory is held while it is suspended. The request's next generation step seeds the replay caches of whichever slot it resumes in from the V2 conv state, as a disaggregated generation request's are. Both run on the KV caches' stream: the commit after the current stream's work and before the event the suspend records there, which later page migrations wait for; the seeding after the pages a resume brings back. With a host drafter, update_resources records a batch's acceptance, which the commit needs. Executors record it before they can suspend the request: the host drafter runs without the overlap scheduler, and the PP loop does not evict requests in flight. A suspend between a batch's prepare_resources and its update_resources raises instead of committing a stale count. The host copies and the late acceptance hand-off are removed. The tests resume a request in a new slot and in its old one, check the host drafter guard, and on the V2 runtime with the fused verify kernel compare every verify after a suspend and resume with the run that is never suspended, bit for bit, with the slot taken by another cache meanwhile and without. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
Since both RMSNorm-fused kernels count Lamport arrivals per CTA, nothing instantiates LamportFlags<..., true>. Remove the UseCGA template parameter, the cluster branches of ctaArrive and waitAndUpdate, and the cooperative_groups include and alias that only they used. Also delete the commented-out cluster.sync() placeholders that the added cluster.sync() calls replace. No functional change: every kernel already took the per-CTA path. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…ast it once Follow-ups on the trained mask-token row this branch keeps: - The row is checked at load, before any weight is copied: the shipped embedding must be [rows, hidden_size], and the drafter's mask_token_id must be one of its rows. A ValueError names the problem. Before, an id past the rows kept an empty row, which failed later in the worker's noise block. - The row is cast to the drafter's dtype and placed on the device of the drafter's own fc weight once, at load. The noise block uses it as given, with no per-step cast and no reliance on the current CUDA device. - The docstring says why only the mask row is kept: in the Kimi K3 drafter (RedHatAI/Kimi-K3-speculator.dspark) it is the only one of the 163,840 embedding rows that differs from the target's, and the shipped lm_head is bit-exact with the target's. Tests: test_gqa_dspark_mask_row_takes_the_drafter_dtype_and_device_at_load and test_gqa_dspark_rejects_a_mask_row_the_shipped_embedding_does_not_have (out-of-range, negative and wrong-width cases). Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
… mask id eda9db2 raised for any mask id outside the shipped embedding, including DFlash's default (vocab_size). The drafter takes that default when its config names no mask token in a key it reads. The speculative config can still name the mask token then (a user value, or a key only it reads, such as dspark_noise_token_id), and such a drafter loaded before. It now keeps no row, as a checkpoint without an embedding does. Any other id outside the shipped rows still raises. Test: test_gqa_dspark_keeps_no_mask_row_when_its_config_names_no_mask_token. The out-of-range case of the rejection test now uses vocab_size + 1. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
|
Automatically added "ci: full pre-merge approved" because this PR has satisfied the required GitHub review approvals. Unresolved review conversations and other required checks remain independent merge requirements. |
Collaborator
|
PR_Github #77652 [ run ] completed with state
|
Collaborator
Author
|
/bot run --disable-fail-fast |
Collaborator
|
PR_Github #77681 [ run ] triggered by Bot. Commit: |
Collaborator
|
PR_Github #77681 [ run ] completed with state |
Collaborator
|
I am merging this one since most related developers/codeowners had approved it. |
vsabavat
added a commit
to vsabavat/TensorRT-LLM
that referenced
this pull request
Oct 14, 2026
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
reasonsolo
added a commit
to reasonsolo/TensorRT-LLM
that referenced
this pull request
Oct 15, 2026
…erify kernel NVIDIA#19988 commits a suspended request's pending KDA drafts into its V2 states through a replay-only variant of the fused verify kernel (kda_mtp_commit_pending_drafts), which read the joint q/k/gate replay cache. The tensor-core verify kernel replaced that kernel and split the cache into kda_k_cache (bf16) and kda_g_cache (fp32), with bf16 v / beta / conv caches, so the suspend path no longer had a kernel or a cache to read. - Port the replay-only mode to the tcgen05 chunk kernel (REPLAY_ONLY). The chunk is the request's accepted drafts, the committed row is the last of them, and only the state is stored. With no draft pending the commit row is -1: e^{G} = 1 and Kbar is empty, so the state is stored unchanged. The token inputs, conv weights, gate constants, conv caches and outputs are not touched, and the precompute that writes the caches is compiled out. - kda_mtp_decode_impl takes replay_only (shifted metadata only, no hints, output norm or quantization); kda_mtp_commit_pending_drafts and commit_kda_pending_drafts take k_cache / g_cache, and the V2 manager's _commit_kda_pending_drafts passes the split caches. Every other part of NVIDIA#19988's suspend / resume bookkeeping is unchanged. One property is not kept: the verify after a resume starts a new chunk at the golden token, and the WY-form chunk rounds differently from a verify that replays the drafts in the same chunk, so the outputs and state after a suspend agree to the kernel's accuracy rather than bit for bit. The conv windows and replay caches stay bit-exact. Tests: the commit-then-verify parity test and the V2 suspend/resume test compare the outputs and state with tolerances (the verifies up to the suspend still bit for bit) and keep the caches and windows exact; a new test checks the replay-only commit against an fp32 CPU replay of the accepted drafts, on both the cluster-pair and the one-CTA launch, and that unselected slots, empty slots and all caches are left alone. The cache-manager stubs follow the split caches. Signed-off-by: Lizhi Zhou <1432185+reasonsolo@users.noreply.github.com>
reasonsolo
added a commit
to reasonsolo/TensorRT-LLM
that referenced
this pull request
Oct 15, 2026
CUDA-graph padding and attention-DP dummy requests all map to the DFlash worker's dummy slot, which no request ever frees. If the generation step advanced that slot's context length like a real slot's, the dummies' drafter context would grow every step up to the context cap. Main zeroes it after each update (DFlashWorker._advance_ctx_len, NVIDIA#19988). Add a regression test that drives prepare_1st_drafter_inputs over a real slot and a padding row and checks that only the real slot advances. Signed-off-by: Lizhi Zhou <1432185+reasonsolo@users.noreply.github.com>
This was referenced Oct 15, 2026
reasonsolo
added a commit
to reasonsolo/TensorRT-LLM
that referenced
this pull request
Oct 15, 2026
…erify kernel NVIDIA#19988 commits a suspended request's pending KDA drafts into its V2 states through a replay-only variant of the fused verify kernel (kda_mtp_commit_pending_drafts), which read the joint q/k/gate replay cache. The tensor-core verify kernel replaced that kernel and split the cache into kda_k_cache (bf16) and kda_g_cache (fp32), with bf16 v / beta / conv caches, so the suspend path no longer had a kernel or a cache to read. - Port the replay-only mode to the tcgen05 chunk kernel (REPLAY_ONLY). The chunk is the request's accepted drafts, the committed row is the last of them, and only the state is stored. With no draft pending the commit row is -1: e^{G} = 1 and Kbar is empty, so the state is stored unchanged. The token inputs, conv weights, gate constants, conv caches and outputs are not touched, and the precompute that writes the caches is compiled out. - kda_mtp_decode_impl takes replay_only (shifted metadata only, no hints, output norm or quantization); kda_mtp_commit_pending_drafts and commit_kda_pending_drafts take k_cache / g_cache, and the V2 manager's _commit_kda_pending_drafts passes the split caches. Every other part of NVIDIA#19988's suspend / resume bookkeeping is unchanged. One property is not kept: the verify after a resume starts a new chunk at the golden token, and the WY-form chunk rounds differently from a verify that replays the drafts in the same chunk, so the outputs and state after a suspend agree to the kernel's accuracy rather than bit for bit. The conv windows and replay caches stay bit-exact. Tests: the commit-then-verify parity test and the V2 suspend/resume test compare the outputs and state with tolerances (the verifies up to the suspend still bit for bit) and keep the caches and windows exact; a new test checks the replay-only commit against an fp32 CPU replay of the accepted drafts, on both the cluster-pair and the one-CTA launch, and that unselected slots, empty slots and all caches are left alone. The cache-manager stubs follow the split caches. Signed-off-by: Lizhi Zhou <1432185+reasonsolo@users.noreply.github.com>
reasonsolo
added a commit
to reasonsolo/TensorRT-LLM
that referenced
this pull request
Oct 15, 2026
CUDA-graph padding and attention-DP dummy requests all map to the DFlash worker's dummy slot, which no request ever frees. If the generation step advanced that slot's context length like a real slot's, the dummies' drafter context would grow every step up to the context cap. Main zeroes it after each update (DFlashWorker._advance_ctx_len, NVIDIA#19988). Add a regression test that drives prepare_1st_drafter_inputs over a real slot and a padding row and checks that only the real slot advances. Signed-off-by: Lizhi Zhou <1432185+reasonsolo@users.noreply.github.com>
reasonsolo
added a commit
to reasonsolo/TensorRT-LLM
that referenced
this pull request
Oct 16, 2026
CUDA-graph padding and attention-DP dummy requests all map to the DFlash worker's dummy slot, which no request ever frees. If the generation step advanced that slot's context length like a real slot's, the dummies' drafter context would grow every step up to the context cap. Main zeroes it after each update (DFlashWorker._advance_ctx_len, NVIDIA#19988). Add a regression test that drives prepare_1st_drafter_inputs over a real slot and a padding row and checks that only the real slot advances. Signed-off-by: Lizhi Zhou <1432185+reasonsolo@users.noreply.github.com>
reasonsolo
added a commit
to reasonsolo/TensorRT-LLM
that referenced
this pull request
Oct 16, 2026
…erify kernel NVIDIA#19988 commits a suspended request's pending KDA drafts into its V2 states through a replay-only variant of the fused verify kernel (kda_mtp_commit_pending_drafts), which read the joint q/k/gate replay cache. The tensor-core verify kernel replaced that kernel and split the cache into kda_k_cache (bf16) and kda_g_cache (fp32), with bf16 v / beta / conv caches, so the suspend path no longer had a kernel or a cache to read. - Port the replay-only mode to the tcgen05 chunk kernel (REPLAY_ONLY). The chunk is the request's accepted drafts, the committed row is the last of them, and only the state is stored. With no draft pending the commit row is -1: e^{G} = 1 and Kbar is empty, so the state is stored unchanged. The token inputs, conv weights, gate constants, conv caches and outputs are not touched, and the precompute that writes the caches is compiled out. - kda_mtp_decode_impl takes replay_only (shifted metadata only, no hints, output norm or quantization); kda_mtp_commit_pending_drafts and commit_kda_pending_drafts take k_cache / g_cache, and the V2 manager's _commit_kda_pending_drafts passes the split caches. Every other part of NVIDIA#19988's suspend / resume bookkeeping is unchanged. One property is not kept: the verify after a resume starts a new chunk at the golden token, and the WY-form chunk rounds differently from a verify that replays the drafts in the same chunk, so the outputs and state after a suspend agree to the kernel's accuracy rather than bit for bit. The conv windows and replay caches stay bit-exact. Tests: the commit-then-verify parity test and the V2 suspend/resume test compare the outputs and state with tolerances (the verifies up to the suspend still bit for bit) and keep the caches and windows exact; a new test checks the replay-only commit against an fp32 CPU replay of the accepted drafts, on both the cluster-pair and the one-CTA launch, and that unselected slots, empty slots and all caches are left alone. The cache-manager stubs follow the split caches. Signed-off-by: Lizhi Zhou <1432185+reasonsolo@users.noreply.github.com>
reasonsolo
added a commit
to reasonsolo/TensorRT-LLM
that referenced
this pull request
Oct 16, 2026
CUDA-graph padding and attention-DP dummy requests all map to the DFlash worker's dummy slot, which no request ever frees. If the generation step advanced that slot's context length like a real slot's, the dummies' drafter context would grow every step up to the context cap. Main zeroes it after each update (DFlashWorker._advance_ctx_len, NVIDIA#19988). Add a regression test that drives prepare_1st_drafter_inputs over a real slot and a padding row and checks that only the real slot advances. Signed-off-by: Lizhi Zhou <1432185+reasonsolo@users.noreply.github.com>
VALLIS-NERIA
pushed a commit
to VALLIS-NERIA/TensorRT-LLM
that referenced
this pull request
Oct 19, 2026
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
VALLIS-NERIA
pushed a commit
to VALLIS-NERIA/TensorRT-LLM
that referenced
this pull request
Oct 19, 2026
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
reasonsolo
added a commit
to reasonsolo/TensorRT-LLM
that referenced
this pull request
Oct 23, 2026
CUDA-graph padding and attention-DP dummy requests all map to the DFlash worker's dummy slot, which no request ever frees. If the generation step advanced that slot's context length like a real slot's, the dummies' drafter context would grow every step up to the context cap. Main zeroes it after each update (DFlashWorker._advance_ctx_len, NVIDIA#19988). Add a regression test that drives prepare_1st_drafter_inputs over a real slot and a padding row and checks that only the real slot advances. Signed-off-by: Lizhi Zhou <1432185+reasonsolo@users.noreply.github.com>
reasonsolo
added a commit
to reasonsolo/TensorRT-LLM
that referenced
this pull request
Oct 23, 2026
CUDA-graph padding and attention-DP dummy requests all map to the DFlash worker's dummy slot, which no request ever frees. If the generation step advanced that slot's context length like a real slot's, the dummies' drafter context would grow every step up to the context cap. Main zeroes it after each update (DFlashWorker._advance_ctx_len, NVIDIA#19988). Add a regression test that drives prepare_1st_drafter_inputs over a real slot and a padding row and checks that only the real slot advances. Signed-off-by: Lizhi Zhou <1432185+reasonsolo@users.noreply.github.com>
reasonsolo
added a commit
to reasonsolo/TensorRT-LLM
that referenced
this pull request
Oct 24, 2026
CUDA-graph padding and attention-DP dummy requests all map to the DFlash worker's dummy slot, which no request ever frees. If the generation step advanced that slot's context length like a real slot's, the dummies' drafter context would grow every step up to the context cap. Main zeroes it after each update (DFlashWorker._advance_ctx_len, NVIDIA#19988). Add a regression test that drives prepare_1st_drafter_inputs over a real slot and a padding row and checks that only the real slot advances. Signed-off-by: Lizhi Zhou <1432185+reasonsolo@users.noreply.github.com>
1 task done
zheyuf
added a commit
to zheyuf/TensorRT-LLM
that referenced
this pull request
Oct 26, 2026
Port of main's DFlashWorker._advance_ctx_len (NVIDIA#19988, c98b976), which the unified KV cache stack relies on but agentx lacks. CUDA-graph padding, warmup and attention-DP dummy rows share one dummy slot that no request owns, and each step added the dummies' accepted tokens to its length. On a unified KV cache the drafter's allocation check covers those rows, whose pages are the padding request's, so the growing length tripped the device assert during CUDA-graph capture at server start. Empty the slot right after the add, inside the captured step; with the windowed arena its base too. Signed-off-by: Zheyu Fu <zheyuf@nvidia.com>
6 of 12 tasks
feih-nv
added a commit
to feih-nv/flashinfer
that referenced
this pull request
Nov 11, 2026
…heck the launch Port of TensorRT-LLM 16df7eaf2 (NVIDIA/TensorRT-LLM#19727, the noAuxTcKernels.cu hunk only) and c98b9765d (NVIDIA/TensorRT-LLM#19988). With no tokens the launcher issued a zero-sized grid. cudaLaunchKernelEx rejected it, nothing checked the return value, and the error surfaced on the next unrelated CUDA call. The launcher now returns before the launch when num_tokens == 0 and fails the call itself when the launch does not succeed, as the other FlashInfer launchers do. The regression test routes zero tokens on the grouped, 256-expert and 128-expert single-group paths, with and without routing_replay_out, and checks that the following CUDA call is clean (6 of 6 cases failed before).
feih-nv
added a commit
to feih-nv/flashinfer
that referenced
this pull request
Nov 11, 2026
…heck the launch Port of TensorRT-LLM 16df7eaf2 (NVIDIA/TensorRT-LLM#19727, the noAuxTcKernels.cu hunk only) and c98b9765d (NVIDIA/TensorRT-LLM#19988). With no tokens the launcher issued a zero-sized grid. cudaLaunchKernelEx rejected it, nothing checked the return value, and the error surfaced on the next unrelated CUDA call. The launcher now returns before the launch when num_tokens == 0 and fails the call itself when the launch does not succeed, as the other FlashInfer launchers do. The regression test routes zero tokens on the grouped, 256-expert and 128-expert single-group paths, with and without routing_replay_out, and checks that the following CUDA call is clean (6 of 6 cases failed before).
feih-nv
added a commit
to flashinfer-ai/flashinfer
that referenced
this pull request
Nov 12, 2026
…heck the launch (#9345) <!-- .github/pull_request_template.md --> ## 📌 Description Port of TensorRT-LLM [16df7eaf2](NVIDIA/TensorRT-LLM@16df7ea) (NVIDIA/TensorRT-LLM#19727; only its `noAuxTcKernels.cu` hunk, the rest of that commit is a DeepSeek-V4.1 feature) and [c98b9765d](NVIDIA/TensorRT-LLM@c98b976) (NVIDIA/TensorRT-LLM#19988, which folds in NVIDIA/TensorRT-LLM#19840: "skip zero-row launches in fp4_quantize, mxfp8_quantize, noaux_tc_op" with tests that failed on their main) to the DeepSeek-V3 fused routing launcher (`invokeNoAuxTc`, `csrc/fused_moe/noAuxTcKernels.cu`). - **Bug.** `fused_topk_deepseek` with zero tokens set `gridDim = 0`. `cudaLaunchKernelEx` rejects that with `cudaErrorInvalidValue`; the return value was not checked, so the call returned normally and the error surfaced on the next unrelated CUDA call (`AcceleratorError: CUDA error: invalid argument` from, say, `torch.ones(4).sum()`), on every path and with or without `routing_replay_out`. - **Fix.** Return before the launch when `num_tokens == 0`; otherwise check the launch with `TLLM_CUDA_CHECK`, as upstream does (it throws with the CUDA error string). No change for non-empty inputs. - **Test.** `test_zero_tokens_launch_nothing`: grouped, 256-expert and 128-expert single-group configurations, with and without a replay buffer; asserts the output shapes and that the following CUDA call is clean. 6 of 6 cases failed before the fix. ## 🔍 Related Issues NVIDIA/TensorRT-LLM#19727, NVIDIA/TensorRT-LLM#19988 ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.). - [x] This PR requires **full-scope** CI, not targeted tests (expensive; justify below). Reason: <!-- e.g. cross-cutting change to shared infrastructure --> Suggested CI test targets (optional; paths under `tests/`, whitespace-separated — a starting point for `@flashinfer-bot run <targets>` / `/bot run <targets>`, reviewers may adjust): `tests/model_optimizations/test_dsv3_fused_routing.py` Run on B200 (SM100) in the CUDA 13.0 CI container at this branch's head: ``` pytest tests/model_optimizations/test_dsv3_fused_routing.py \ -k "zero_tokens_launch_nothing or (backend_correctness and default) or routing_replay_out_extended" 84 passed, 30 skipped ``` `test_zero_tokens_launch_nothing` on the pre-fix kernel: 6 failed. ## 🔬 Experimental Track <!-- Only for PRs submitted under the experimental policy (CONTRIBUTING.md → "Experimental APIs and Backends"). Leave this section untouched for normal PRs. --> - [ ] This PR is **experimental**: it adds or changes code under `flashinfer/experimental/` and/or an `@flashinfer_experimental_api`. Tracking issue: # - [ ] The tracking issue names an owner, the reason for the experimental path, and a graduation plan with a target release. - [ ] Core changes are limited to a thin entry point (signature, shared validation, feature-gate check, backend selection, handoff). - [ ] Tests live in `tests/experimental/` and were validated on the intended hardware; a runnable example is included. - [ ] Nothing is registered in `flashinfer/aot.py`, and no experimental backend is reachable from `backend="auto"` without `FLASHINFER_ALLOW_EXPERIMENTAL_AUTO_BACKENDS=1`. (Calling an `@flashinfer_experimental_api` or naming a backend explicitly is itself the opt-in and needs no environment variable.) - [ ] **Test scope declared below.** The experimental CI lane runs exactly these targets, so keep them as narrow as the change allows. <!-- Required for experimental PRs. Replace the commented lines below with your targets. Do not delete the fence or change its `experimental-tests` tag — the experimental-track watcher reads it verbatim to decide which targets to ask CI for. --> ```experimental-tests # One target per line: a directory or a file. (A pytest ::selector is not # supported -- the sharding runner cannot consume one.) Must be under # tests/experimental/ and must exist. Delete these comment lines and add yours, e.g. # # tests/experimental/test_my_backend.py # tests/experimental/my_backend/ # # Declaring the whole tree (tests/experimental/) is allowed but means every # experimental PR pays for every other feature's tests, in every matrix cell. ``` ## Reviewer Notes The Cake backend already rejects zero tokens in its contract; this is the default backend only. The new test sits with the other input-contract test (`test_dsv3_fused_routing_rejects_strided_views`). #9344 adds its test elsewhere in the same file; whichever merges second gets rebased. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> <!-- coderabbit-ai-generated:v1 --> ## Summary by CodeRabbit * **Bug Fixes** * Zero-token routing operations now complete without launching unnecessary GPU work across grouped and single-group configurations. Empty outputs retain their expected shapes, and subsequent GPU operations can proceed normally. * GPU kernel launch failures are now checked and reported before stream error synchronization, while successful launches continue through the existing error checks. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
5 of 11 tasks
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Description
Twelve independent fixes for defects on main. Each one comes with a test that fails on main and passes here. These were
reviewed as separate PRs and are combined at the maintainers' request; every reviewed commit is kept.
modules/multi_stream_utils.py,peft/lora/layer.pydistributed/ops.pycommon/lamportUtils.cuh,mnnvlAllreduceKernels.cusPartialLastbefore the next chunkkimi_k3_kda/fused_k123.pycustom_ops/cute_dsl_kimi_k3_custom_ops.pycustom_ops/cute_dsl_kimi_k3_custom_ops.pykv_cache/mamba_cache_manager.py, KDA MTP opsspeculative/dflash.pymodels/modeling_dspark.py,speculative/dflash.pyfp4_quantize,mxfp8_quantize,noaux_tc_op; unwaive the three tests of nvbugs/6852473kernels/quantization.cu,kernels/noAuxTcKernels.cufla-coreis a runtime dependency of the KDA layersrequirements.txtthop_attentioncatalog test: drop a probe whose context length exceeds the call's tokens, which wrote past its output buffertests/.../test_modeling_v2_thop_attention.pyThe linked PRs have the full failure analyses.
Test Coverage
GB200. Each fix's new test was run on main's code and on this branch:
modules/test_multi_stream_utils.pymulti_gpu/test_mnnvl_allreduce.py::test_mnnvl_workspace_growth_keeps_captured_graphs(4 GPUs)multi_gpu/test_mnnvl_allreduce.pyunder compute-sanitizer racechecktest_kda_prefill_op.py::test_fused_prefill_does_not_depend_on_k1_warp_timingtest_kda_prefill_op.py::test_kda_prefill_op_output_survives_the_next_calltest_kda_cache_soundness.py::test_evicted_prefill_scratch_is_releasedtest_mamba_cache_manager.py(suspend / rebalance),test_kda_mtp_decode_cute_parity.pyspeculative/hw_agnostic/test_dflash_dummy_slot.pytest_kimi_k3_dspark_semantics.py(GQA mask-row tests)test_noaux_tc.py,test_fp8_quantize.py,test_fp4_gemm_quantize.py -k zero_rows, the 3 unwaived testsothers/test_kimi_kda_runtime_dependencies.pymodeling_v2/activation+test_modeling_v2_thop_attention.pyunder memcheckTested on main f891b3d. Wider groups were run against main: kv_cache, speculative, kimi_kda and mamba show no
failure that main lacks.
Limitations:
[None][fix] DFlash: start the dummy slot at context length 0 every step #19827's GB300 PerfSanity timeout is unresolved; it ran DeepSeek-V4 with MTP, which does not reach the DFlash
path this PR changes.
PR Checklist
Dev Engineer Review
torch.compile. The LoRA parallel path no longer callsrecord_stream; verify output lifetime under its stream-ordering conditions.fla-core==0.5.2is added for KDA runtime imports.QA Engineer Review
thop_attentioncontext-length behavior.thop_attentiontest removes a context-length probe that can write past the output buffer and documents the validcontext_lengthscondition.Per-File QA Perspective
cpp/tensorrt_llm/common/lamportUtils.cuh: Verify CTA arrival and coordinator behavior on architectures below and above SM 7.0. TheUseCGAparameter and cluster-coordination path are removed.cpp/tensorrt_llm/kernels/communicationKernels/mnnvlAllreduceKernels.cu: Verify one-shot and two-shot reductions and cluster RMSNorm synchronization. The kernels use CTA-counted Lamport arrivals and synchronize before peer shared-memory writes.cpp/tensorrt_llm/kernels/noAuxTcKernels.cu: Verify zero-token behavior and launch-error propagation throughTLLM_CUDA_CHECK.cpp/tensorrt_llm/kernels/quantization.cu: Verify zero-row FP4 and MXFP8 behavior and launch-error propagation throughTLLM_CUDA_CHECK.requirements.txt: Verify runtime installation includesfla-core==0.5.2for KDA imports.tensorrt_llm/_torch/custom_ops/cute_dsl_kimi_k3_custom_ops.py: Verify scratch-wrapper lifetime and fresh output ownership for padded and variable-length calls.tensorrt_llm/_torch/custom_ops/cute_dsl_kimi_k3_kda_mtp_ops.py: Verify replay-only validation, compile-cache separation, and accepted-draft commits.tensorrt_llm/_torch/cute_dsl_kernels/blackwell/kimi_k3_kda/fused_k123.py: Verify K1 warps synchronize after loading scan values and before Pass 2a.tensorrt_llm/_torch/cute_dsl_kernels/blackwell/kimi_k3_kda/kda_mtp_decode.py: Verify replay-only execution updates recurrent state without normal decode work.tensorrt_llm/_torch/distributed/ops.py: Verify workspace growth fails during capture and captured workspaces remain available to existing graphs.tensorrt_llm/_torch/models/modeling_dspark.py: Verify speculative mask-token selection, checkpoint embedding validation, and shared-embedding fallback.tensorrt_llm/_torch/modules/kimi_kda/_kda_kernels.py: Verify accepted-draft replay and convolution-window writes for selected state slots.tensorrt_llm/_torch/modules/multi_stream_utils.py: Verify recursive output recording outside capture and compilation, including nested containers and allocator error handling. Its tests are listed in B200 and H100 test-db lists.tensorrt_llm/_torch/peft/lora/layer.py: Verify parallel LoRA output lifetime after removal of therecord_streamcalls.tensorrt_llm/_torch/pyexecutor/kv_cache/mamba_cache_manager.py: Verify acceptance tracking, state commits on suspension, restoration on resume, slot relocation, and rebalance rejection when KDA replay is active.tensorrt_llm/_torch/speculative/dflash.py: Verify noise-block embedding selection and dummy-slot context-length reset.tensorrt_llm/_torch/pyexecutor/config_utils.py: Verify speculators-format DSpark configs translate to the expected Qwen3 config and disable the sliding window.tensorrt_llm/llmapi/llm_args.py: Verify DSpark config detection, translation, validation, target-layer ID conversion, and DFlash config resolution.tests/integration/test_lists/test-db/l0_b200.yml: Adds the multi-stream utility test to the B200 pre-merge list.tests/integration/test_lists/test-db/l0_gb200_multi_gpus.yml: Adds the MNNVL workspace-growth graph test to the GB200 multi-GPU list.tests/integration/test_lists/test-db/l0_h100.yml: Adds the multi-stream utility test to the H100 pre-merge list.tests/integration/test_lists/test-db/l0_dgx_b200.yml: Adds the MNNVL workspace-creation capture-guard test to the B200 pre-merge list.tests/integration/test_lists/waives.txt: Removes three modeling-v2 MoE and quantization waivers.tests/unittest/_torch/executor/kv_cache/test_mamba_cache_manager.py: Covers KDA replay rebalance restrictions, host-drafter acceptance, and suspend/resume. No changed test-db entry is identified.tests/unittest/_torch/modeling_v2/attention/test_modeling_v2_thop_attention.py: Removes an out-of-range context-length probe and clarifies valid input behavior. No changed test-db entry is identified.tests/unittest/_torch/modules/kimi_kda/test_kda_cache_soundness.py: Covers bounded scratch-cache size, memory growth, and eviction release. No changed test-db entry is identified.tests/unittest/_torch/modules/kimi_kda/test_kda_mtp_decode_cute_parity.py: Covers committed accepted drafts, recurrent-state parity, and convolution-window restoration. No changed test-db entry is identified.tests/unittest/_torch/modules/kimi_kda/test_kda_prefill_op.py: Covers K1 timing interference and output ownership for padded and variable-length prefill. No changed test-db entry is identified.tests/unittest/_torch/modules/test_multi_stream_utils.py: Covers output lifetime across streams, graph capture, and compilation. Listed in B200 and H100 test-db lists.tests/unittest/_torch/multi_gpu/test_mnnvl_allreduce.py: Covers workspace growth with captured graphs and workspace creation during capture. The graph-growth test is listed in the GB200 multi-GPU test-db list; the capture-guard test is listed in the B200 test-db list.tests/unittest/_torch/speculative/hw_agnostic/test_dflash_dummy_slot.py: Covers dummy-slot reset during preparation and graph replay. No changed test-db entry is identified.tests/unittest/_torch/speculative/hw_agnostic/test_kimi_k3_dspark_semantics.py: Covers trained mask embeddings, mask-ID validation, and speculative-config resolution. No changed test-db entry is identified.tests/unittest/_torch/speculative/hw_agnostic/test_dspark_speculators_checkpoint.py: Covers published DSpark config translation, equivalence with a hand-converted config, draft-length validation, and unchanged handling of other configs. No changed test-db entry is identified.tests/unittest/_torch/thop/parallel/test_fp4_gemm_quantize.py: Covers zero-row FP4 output shapes and a subsequent nonempty call. No changed test-db entry is identified.tests/unittest/_torch/thop/parallel/test_fp8_quantize.py: Covers zero-row MXFP8 output shapes and a subsequent nonempty call on supported devices. No changed test-db entry is identified.tests/unittest/_torch/thop/parallel_hw_agnostic/test_noaux_tc.py: Covers zero-token output shapes and a subsequent nonempty call. No changed test-db entry is identified.tests/unittest/others/test_kimi_kda_runtime_dependencies.py: Checks that KDA import providers are declared in runtime requirements. No changed test-db entry is identified.Current main-conflict repair validation
Published
72bb8ca0b8. Native builds and affected CPU gates passed; remaining collective, 16-rank functional and model-serving comparisons are pending. Local qualification continues in parallel with CI; no pending gate is claimed passed.