Update TensorRT-LLM#4

kaiyux · 2023-10-11T06:26:45Z

No description provided.

Update TensorRT-LLM

#4550) Signed-off-by: Superjomn <328693+Superjomn@users.noreply.github.com>

# This is the 1st commit message: kernel Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> remove prints Signed-off-by: Ubuntu <dafrimi@nvidia.com> test pass Signed-off-by: Ubuntu <dafrimi@nvidia.com> test refactor with more use cases Signed-off-by: Ubuntu <dafrimi@nvidia.com> refacor Signed-off-by: Ubuntu <dafrimi@nvidia.com> refacor_2 Signed-off-by: Ubuntu <dafrimi@nvidia.com> add tuner wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> autotuner works Signed-off-by: Ubuntu <dafrimi@nvidia.com> bfloat16 works. moer changes to the thop file Signed-off-by: Ubuntu <dafrimi@nvidia.com> is tune for autotuner is True --> gets real tactics configs Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> zeros + quant mode is works Signed-off-by: Ubuntu <dafrimi@nvidia.com> act int8 Signed-off-by: Ubuntu <dafrimi@nvidia.com> removed fp8 for now Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> w4a16 linear module Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> changed cutalss for sm==89 Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> test linear work Signed-off-by: Ubuntu <dafrimi@nvidia.com> add license Signed-off-by: Ubuntu <dafrimi@nvidia.com> works! Signed-off-by: Ubuntu <dafrimi@nvidia.com> refactor + linear test pass Signed-off-by: Ubuntu <dafrimi@nvidia.com> preprocess in load weights Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> refactor + rebase Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> Blackwell not supported Signed-off-by: Daniel Afrimi <dafrimi@nvidia.com> wip Signed-off-by: Daniel Afrimi <dafrimi@nvidia.com> skip blackwell Signed-off-by: Daniel Afrimi <dafrimi@nvidia.com> wip Signed-off-by: Daniel Afrimi <dafrimi@nvidia.com> works Signed-off-by: Ubuntu <dafrimi@nvidia.com> # This is the commit message NVIDIA#2: rebased Signed-off-by: Ubuntu <dafrimi@nvidia.com> # This is the commit message NVIDIA#3: align with my pld worked version of linear Signed-off-by: Ubuntu <dafrimi@nvidia.com> # This is the commit message NVIDIA#4: wip Signed-off-by: Ubuntu <dafrimi@nvidia.com> # This is the commit message NVIDIA#5: refactor Signed-off-by: Daniel Afrimi <danielafrimi8@gmail.com> # This is the commit message NVIDIA#6: refactor Signed-off-by: Daniel Afrimi <danielafrimi8@gmail.com> # This is the commit message NVIDIA#7: refactor Signed-off-by: Daniel Afrimi <danielafrimi8@gmail.com> # This is the commit message NVIDIA#8: refactor Signed-off-by: Daniel Afrimi <danielafrimi8@gmail.com> # This is the commit message NVIDIA#9: sys path Signed-off-by: Daniel Afrimi <danielafrimi8@gmail.com> # This is the commit message NVIDIA#10: sys path Signed-off-by: Daniel Afrimi <danielafrimi8@gmail.com>

Update sm100f cubins

Signed-off-by: Yuxian Qiu <142763828+yuxianq@users.noreply.github.com>

Signed-off-by: Yuxian Qiu <142763828+yuxianq@users.noreply.github.com> Signed-off-by: Fanrong Li <23290157+lfr-0531@users.noreply.github.com>

- Revert utils.py docstring changes (NVIDIA#1) - Remove disable_deep_gemm flag, use module-level quant_config=None (NVIDIA#2) - Replace Flux2SwiGLU with shared swiglu from _torch/modules (NVIDIA#3) - Unify Flux2PosEmbed into FluxPosEmbed with parameterized axes_dim (NVIDIA#4) - Replace Flux2FeedForward with shared GatedMLP (NVIDIA#6) - Revert pipeline.py post_load_weights; add create_weights() in FLUX transformer __init__ to match WAN's __post_init__ pattern (NVIDIA#7) - Refactor _get_quant_algo_for_layer to encapsulate hasattr logic (NVIDIA#9) - Remove partial docstrings that just repeat class names (NVIDIA#10) - Simplify bias flags to direct booleans (NVIDIA#11) - Unify apply_rotary_emb: delete FLUX copy, FluxPosEmbed outputs [1,S,1,D] matching WAN format so shared function unchanged (NVIDIA#12) - Remove dynamic=True from maybe_compile (NVIDIA#13) - Add @torch.compiler.disable on FluxPosEmbed.forward - Tighten FLUX.2 PSNR test threshold from 12 to 20 dB Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com> Signed-off-by: Kanghwan Jang <861393+karljang@users.noreply.github.com>

… PRs landed Apply nine consistency fixes against the post-NVIDIA#13671/NVIDIA#13672/NVIDIA#13673/NVIDIA#13674 state of the investigation: 1. Front-matter Status block: replace the "sig NVIDIA#6 root-caused, validation in flight" wording with the post-run8 picture (all 6 TRT-LLM PRs in review; NVIDIA#7 is an out-of-scope NIXL bug; deadline work is the TRT-LLM-side fallback for NVIDIA#7). 2. Front-matter Branches in this worktree: add the four new sig #4 / #5 / NVIDIA#6 branches. 3. Front-matter Related PRs: add NVIDIA#13674 / NVIDIA#13671 / NVIDIA#13672 / NVIDIA#13673 with chained-on-NVIDIA#13640 callout for NVIDIA#13673. 4. "Configurations that did not reproduce": #5 and NVIDIA#6 now do reproduce in single-process unit tests via the new tests added by NVIDIA#13672 and NVIDIA#13673; only #3 and NVIDIA#7 remain field-only. 5. Phase 6 close: the sig #4 regression test is no longer isolated in local/rc11-disagg-repro - it is now in the chained NVIDIA#13674 / NVIDIA#13671 pair. 6. Signature NVIDIA#6 section: drop the "(suspected)" qualifier and the "(most likely)" hedging on Where-it-lives - both are confirmed by run7 and run8. Rename the section header to describe the actual failure shape (recv buffer index leak via !isReady early return wedging assignBufferIndexForRecv) rather than the early control-path-stall hypothesis. Mirror the rename in the Signature - PR Map row. 7. File / Branch Index "New unit tests": add the new sig #5 (test_cancel_queued_gen_request_fulfills_receiver_future) and sig NVIDIA#6 (test_cancelled_after_ready_does_not_leak_recv_buffer_index, NIXL backend) tests. 8. Signature #3 status hypothesis: add a one-paragraph note that NIXL (signature NVIDIA#7) is now also a candidate cause of the half-initialized state, so a future field hit is not misattributed to a fresh TRT-LLM bug. 9. Phase 5 narrative: add a forward link explaining that the underlying terminal driver of the Phase-5 wedge was already NVIDIA#7 (NIXL), but #4 was the visible TRT-LLM-side symptom because the gen event loop was self-blocking before any of the later layers could surface. Signed-off-by: Chien-Chun Hung <2679986+chienchunhung@users.noreply.github.com> Made-with: Cursor

… section Fold two related pieces of analysis into the report as a new section between the Investigation Timeline and "Why the Existing Tests Did Not Catch This": (1) Signature taxonomy refining the naive "burst -> timeout -> cancellation -> bug" framing. Four-of-seven signatures (#1, #3, #5, NVIDIA#6) are direct cancellation-handling bugs; #4 is a structural latent blocking bug that cancellations expose; #2 is an eviction-driven bug that burst traffic exposes via memory pressure; NVIDIA#7 is a NIXL-internal contention bug that the same load shape happens to trigger but which is not strictly a cancellation bug. Includes a refined trigger chain diagram and two precise corrections (burst alone is not the trigger; "cancellation" is one of several entry points to cleanup paths). (2) Cascade map distinguishing two kinds of inter-signature tangling: - Type 1 (a fix produces a new signature): only one case, the #1 fix produces NVIDIA#6 by making the receiver-side !isReady early-return path reachable in production where a latent recv-buffer leak existed. This is why NVIDIA#6 PR (NVIDIA#13673) is explicitly chained on #1 fix PR (NVIDIA#13640). - Type 2 (a fix exposes a pre-existing signature): three cases where #4 fix exposes #5 / NVIDIA#6 and NVIDIA#6 fix exposes NVIDIA#7 because the upstream fix removes the masking effect on the downstream bug. These are not regressions of the fixes; they were latent pre-existing issues. - Subtler third relationship: #4 is structurally a defensive catcher for any upstream bug that produces a never-resolving receiver future. The #4 fix is independently valuable as defence in depth, not just a symptomatic patch. Also includes a fix-to-file mapping showing that the fixes do not overlap in code; the only structural dependency is the NVIDIA#6 -> #1 chain enforced by the PR base. Signed-off-by: Chien-Chun Hung <2679986+chienchunhung@users.noreply.github.com> Made-with: Cursor

…cess retrospective Add a new section between "Why the Existing Tests Did Not Catch This" and "Architectural Reflections" that answers the natural follow-up question after the cascade map: with the e2e view in hand, what would the cleaner approach have been from the start? Two halves: (1) What was not actually possible at T0 — explains why a "design one comprehensive fix" approach (e.g. a TransferSession-like abstraction introduced in PR #0) was strictly impossible given the information state at T0: you cannot design an abstraction to fix bugs you have not found yet, and the field was wedged under P0 urgency that did not allow a multi-thousand-line refactor of a critical path. (2) What we should have done first, in priority order, to change the meta-process rather than the fix: - PR #0: deadline enforcement (Next Steps item 7) as a containment layer. Converts every cleanup-path bug from a silent wedge into a per-request error, gives every subsequent bug an attributable failure point, gives orchestration a real signal so customer deployments self-heal via pod restart. - Write down the seven invariants from the Architectural Reflections section. Would have caught #5 and NVIDIA#6 at #1's PR review (Type 1 cascade prevented at design time) and made #4 visible to any code search. - Add the cancel-during-transfer integration test for the customer load shape. Would have surfaced all six TRT-LLM signatures as CI test failures instead of as a customer field hit. Plus a clarifying "What this section is not arguing" subsection to prevent over-reading: not advocating for a TransferSession rewrite as PR #0, not claiming one PR could fix all seven signatures, and not claiming the sequential discovery was avoidable given the meta-process we actually had at T0. The key insight is that the bottleneck of the investigation was observability and attribution, not fix complexity. The three meta-process changes attack that bottleneck directly. Signed-off-by: Chien-Chun Hung <2679986+chienchunhung@users.noreply.github.com> Made-with: Cursor

NVIDIA#7 confirmed independent of TRT-LLM-side fix strategy Add Phase 11 to the Investigation Timeline documenting the pr13056_run1 experiment: ran the same 1P1D long-prompt burst harness against an independent fix stack (a comprehensive end-to-end shared_ptr<LlmRequest> + BufferIndexHolder RAII + deadline-enforcement refactor) on the same rc11 base. Same outcome as run8: NO RECOVERY after 180s idle, with a gdb-confirmed pthread_mutex_lock frame in CacheSender::Impl::response on the ctx worker, alongside the same NIXL plugin threads in the same process. Coverage map shows the two stacks converge on the same TRT-LLM-side bugs (#1, #4, #5, NVIDIA#6) via different mechanisms (surgical patches vs comprehensive refactor), eliminating two alternative hypotheses: - our chained PRs introduced a regression that masquerades as NVIDIA#7 - comprehensive deadline enforcement alone clears the field reproducer Both refuted by the experiment. The independent stack's defensive diagnostics ([buf] CANCEL: 0, [buf] STILL_WAITING: 0, kNETWORK_ERROR: 0, broken-promise: 0, deadline-driven failures: 0) confirm the wedge is below where any TRT-LLM-side deadline can reach. Update Next Steps item 8 to reference the pr13056_run1 stack dump as a second independent reproducer for the NIXL/UCX bug filing — much stronger evidence than run8 alone because it shows the deadlock is independent of any TRT-LLM fix strategy. Update Phase 10 title to remove the (current) qualifier since Phase 11 is now the latest phase. Signed-off-by: Chien-Chun Hung <2679986+chienchunhung@users.noreply.github.com> Made-with: Cursor

… colored fix components Two improvements to 00-tldr.md based on review feedback: 1. Add a 7-row table after the architecture section that gives a one-line "where it lives" + "symptom" for each signature #1-NVIDIA#7. The TL;DR previously mentioned the seven signatures by number without describing them, forcing the reader to jump to 02-failure-signatures.md to make sense of references like "sig #4" or "sig NVIDIA#7". 2. Color-code the four fix components in the combo diagram: - PR NVIDIA#13056: blue (lifetime + cancel-flag + RAII) - PR NVIDIA#13495: orange (NIXL TransferStatus::release) - eval-order fix: green - Python idempotency guards: purple Each L1-L8 layer node is tinted with the lighter shade of whichever fix closes it, so the "closes" arrows are reinforced by colour-matching. Makes the fix-to-layer mapping visually scannable at a glance. Note added underneath the diagram explaining the colour scheme. Signed-off-by: Chien-Chun Hung <2679986+chienchunhung@users.noreply.github.com>

GLM-5 has hidden_size=6144 with 256 MoE experts, but the dsv3_router_gemm custom kernel only had compile-time instantiations for hidden_dim=7168 (DeepSeek-V3). On Blackwell this caused the routing GEMM to fall through to cublas_mm_out, which dispatched BF16xBF16->FP32 to cutlass_80_simt_sgemm_64x64_8x5_tn_align1 (Ampere SIMT, no tensor cores) — ~5.4% / ~1 ms per decode iter on GLM-5 FP8 MTP=3 BS=1 ISL=1K B200 TP=8. The router_gemm_kernel template is generic over kHiddenDim as long as kHiddenDim is divisible by VPT*kBlockSize = 1024. 6144/1024 = 6 iterations, so adding the 16 num_token instantiations is sufficient. dsv3_router_gemm_op now dispatches between K=7168 and K=6144 paths; unsupported shapes still fall back to cublas_mm_out. Test coverage extended to hidden_size in {7168, 6144}. Refs: nvbug 6108841 — TRT-LLM vs TileRT GLM-5 perf gap, target NVIDIA#4 in COMPARISON_TRTLLM_VS_TILERT.md (cutlass_80_simt_sgemm Ampere fallback). Signed-off-by: Yijing Li <257409031+yijingl-nvidia@users.noreply.github.com>

…view Apply 5 review-driven edits to §16 to address residual concerns: Walk ordering (concern NVIDIA#8): change GMS RO and MX-receiver alias walks from per-module to top-level model.setup_aliases(). Matches the §7 mitigation contract from ai-dynamo/dynamo PR NVIDIA#7053 ("Call model.post_load_weights() (top-level only) before materialize_module_from_gms()"). transform_weights() and cache_derived_state() walks remain per-module since those bodies live on submodules. New "Why setup_aliases() is top-level-only" callout documents the asymmetry. Lifecycle of _weights_transformed (concern #4): new subsection specifying explicit set/reset/orthogonality rules. Includes a 2x2 truth table showing _weights_removed and _weights_transformed track different lifecycles and can take any combination. Reset is the orchestrator's responsibility (e.g., ModelLoader.reload() resets the flag before re-binding tensors); subclasses do not manage reset. Hard preconditions (concern #5): promote MX source-identity completeness from "open question" to "hard precondition P1." Lists transform-affecting parameters that MX identity must cover (attn_backend, quant backend list, FP8/NVFP4 fusion strategy, TP/EP layout, model revision). Specifies an in-tree backend-fingerprint fail-safe as the fallback if upstream MX cannot guarantee completeness. P2 documents that orchestrator owns _weights_transformed reset. Removes redundant open question NVIDIA#6 from the table. Cosmetic fixes (concerns #2, NVIDIA#6): "four stages" -> "three per-module stages plus orchestrator-managed per-process finalization." cache_derived_state description softened to "reserved for data-dependent state where it exists; many existing modules will have empty bodies." Scope clarifications (concerns #1, #3, NVIDIA#7): - Tiny PR scope reframed as "duck-typed helpers, not inheritance" with citations to existing model_loader.py walker pattern. Lists 4 walker helpers (_setup_aliases, _walk_transform, _walk_cache_state, _walk_full_post_load). - Migration callout: when migrating a subclass, the old post_load_weights() override MUST be removed; otherwise the new staged calls silently no-op. - Family PR #2 (Linear/Attention) gains a "quant-method callback decision" note with default = keep quant_method.post_load_weights callback name (no rename). No code changes. Drives Tiny prep PR scope and family-PR migration sequence. References: - TRT-LLM PR NVIDIA#13926 (GMS-only) - TRT-LLM PR NVIDIA#14151 (MX shim refactor) - ai-dynamo/dynamo PR NVIDIA#7053 (upstream GMS prototype) Signed-off-by: Chien-Chun Hung <chienchunh@nvidia.com> Signed-off-by: Chien-Chun Hung <2679986+chienchunhung@users.noreply.github.com>

Bug fixes: - phase4_histogram_snap: add missing cute.arch.barrier() between the smem_hist[0..NW-1] read loop (recomputing block_min/block_max from warp-staged cmax slots) and the histogram-zeroing write loop. Without this, warp-0 thread-0 could finish the unrolled read and start zeroing smem_hist[0] while warp-N was still reading the staged cmax → squashed bmax_r → all candidates land in bin 0 → wrong K-th threshold. Hit-rate-dependent; covered by existing tests once they span realistic preIdx hit rates (NVIDIA#4 from prior review). - Add runtime 32B alignment assert on logits.data_ptr() when the use_256bit_load heuristic fires (LDG.E.256 path). Catches view-with- unaligned-offset misuse with a clear error rather than silent miscompiled addresses / faults. Code hygiene: - Fix stale "unrolled for 64 times" comment in the block_min/max recompute loop; it unrolls num_warps times (16 or 32). - Fix phase1_preidx_stats s_thr docstring: [3] [threshold, val_lo, val_hi]; pmax_saved lives in the separate s_thr_extra allocation. - Document the (lge << 16) | lgt packing's <2^16 per-warp count precondition (currently holds via cand_count ≤ kC ≤ 6144 in GvrParams; future kC bump past 65536 would silently corrupt). Signed-off-by: Mindy Li <11663212+limin2021@users.noreply.github.com>

… cancel-poison design Adds a new investigation doc 19 capturing the external forensic A/B (fengyul/dynamo-disagg exp 4) that decomposes the decode-side wedge into three independent failures: F1 (Broken promise UAF), F2 (engine- loop freeze on unbounded future.get()), and F3 (eager-free poisons UCX progress thread → permanent wedge). Maps F1/F2/F3 onto the existing layer model (L1, L3, L4+L5) and signatures (#1, #4), and onto the design doc's C4 invariant. Strategic conclusions surfaced from the experiment: * PR NVIDIA#14979's shared_ptr lifetime port closes F1 (a co-occurring crash class) but does not recover the field wedge on NIXL by itself. * Bounded polling alone (F2 fix) does not progress a stuck transfer on NIXL because UCX runs its own background progress thread — engine-freeze and transfer-stall are separate concerns. * The load-bearing fix is F3 done safely = quiescence-gated freeing, which is structurally tangled across py_executor.py and the transfer manager API and not cleanly cherry-pickable. The deployable that passes the reproducer is PR NVIDIA#13713 in full. Updates to investigation README: * Adds doc 19 to the 'How to read' index. * Adds a 'Can we ship a small subset of NVIDIA#13713?' suggested reading path. Updates to disagg-inflight-cancel-poison design README: * Adds the exp 4 evidence to 'Empirical evidence motivating this work'. * Adds a 'load-bearing for recovery, not optional polish' note on Phase 2 (deferred un-poison) — the polling-until-quiescence mechanism realizes C4 on the V1 + C++ path and is direct empirical answer to the field wedge, not just operability improvement on top of the existing fail-closed surface. Signed-off-by: Chien-Chun Hung <2679986+chienchunhung@users.noreply.github.com>

#1 FD <-> chunk position misalignment in P2pHandleExporter Old code kept POSIX FDs in a flat mExportedFds and addressed them in removeHandles() by accumulating chunk counts across earlier pools. Pool-extension paths in exportHandles() append chunks to an existing pool but FDs to the END of mExportedFds, breaking that invariant. Removing a registered range could close another pool's FDs and leak the actual ones, also corrupting any in-flight UDS handshake. Fix: store FDs on P2pMemPool.fds in lockstep with chunks. Removal walks only the owning pool's fds - no cross-pool indexing possible. NVIDIA#2 UDS socket exposed POSIX FDs to other local users /tmp/trt_llm_p2p_fd_*.sock was created with default umask, no peer-credential check. Any same-host process that could connect could SCM_RIGHTS-receive the FDs and import the exporter's GPU memory. Fix: tight umask + explicit chmod 0600 around bind, fail closed if either fails; SO_PEERCRED gate at accept() to refuse non-same-uid clients (defense in depth). NVIDIA#3 P2pTransferContextPool never erased per-thread Contexts Long-lived processes whose callers are transient threads leaked ~64MB cubTempStorage + a worker pool per dead thread; OS thread::id reuse could also hand a fresh thread a stale Context built in a different CUDA context. Fix: pool is now held by shared_ptr (factory + private ctor enforce this); contextForCurrentThread installs a thread_local guard that, on thread exit, calls eraseForCurrentThread via weak_ptr (no-op if pool is already gone, so destruction order between agent and caller threads is safe in either direction). NVIDIA#4 Reads of P2pHandleExporter::mLocalInfo were unprotected getLocalAgentDesc did isSupported() then getLocalInfo().serialize() - a TOCTOU pair where a concurrent registerMemory could append a pool between the two reads, producing a half-written wire blob. Fix: single mInfoMutex protects mLocalInfo / mDetectedHandleType / mUdsPath end-to-end. New serializeIfSupported() does both checks under the lock and replaces the pair at the only caller. The previous mPoolsMutex (added for FDs) is the same mutex, renamed. NVIDIA#8 Document acquire-then-record contract on CudaEventPool Recycled events carry their prior recording; calling cudaEventQuery before record() returns stale state. All current callers already follow acquire->record->query, but the contract is now spelled out on the API doc so future callers don't diverge. Tests: p2pTransferAgentTest +2 regressions - SerializeIfSupportedEmptyOnFreshAgent: locked accessor returns empty on a freshly constructed agent. - TransientCallerThreadsReleaseContextOnExit: 32 transient caller threads each take a Context from the same agent; after join, pool->numContexts() must be exactly 0 (proves thread-exit guard fires). Local validation: p2pMemInfoTest 11/11 transferAgentTest 18/18 (AgentDesc + VmmDescSplitter) p2pTransferAgentTest 36/36 nixlP2pE2ETest 7/7 (+2 env-gated, validated separately) Signed-off-by: Chuang Zhu <111838961+chuangz0@users.noreply.github.com>

…che (P0) Implement WriteThroughPolicy.ON_COMMIT (DESIGN §4.3): committed blocks are immutable, so their D2H write-out can happen the moment they commit, making the common free path copy-free. - _commit_block (arena mode, ON_COMMIT, host tier present) copies the just-committed block's pages to fresh host slots. The sources are still locked, so their ready events do not cover in-flight writes: the copy is ordered after an event recorded on the sequence's stream (write_through_pages gains prior_event, plumbed into _batched_migrate as extra_prior_event). A full host tier skips the copy silently -- the copy-on-free fallback covers such blocks. - The per-sequence stash of write-through slots is consumed at close() and suspend(): adopt_stale_copies moves each written-through page to the host tier without another copy. The vacated arena slot's reclaim gate merges the page's own last consumers with the write-through copy's finish event (the copy reads the slot); the page's ready event becomes the copy's finish event, so future consumers wait for its validity. Blocks without a write-through copy keep the existing copy-on-free path; unconsumed stash slots are returned to the host pool. - Rate limiting: the natural rate is one record per tokens_per_block steps per sequence (the §4.3 cost-model baseline); aggregate PCIe throttling is deferred to the executor integration (risk NVIDIA#4). Default behavior is untouched: ON_FREE remains the default and the v2 regression suite result is identical before/after (only pre-existing TestSSMSupport failures, reproduced on main). The ON_COMMIT test class subclasses the end-to-end suite so every arena scenario re-runs under the new policy, plus dedicated tests verify host copies exist at commit time while the blocks are still live on GPU, that close() adopts exactly the pre-copied host slots (no duplication, data verified via reuse), and that suspend() adopts committed blocks while moving the uncommitted tail (60/60 on H100). Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>

…E. All 8 acceptance criteria hold at runtime against... QA verdict: APPROVE (weighted_score 8.6) Change vs main: 7 files changed, 1376 insertions(+), 8 deletions(-) Build iterations run: 13 Code-quality verification: clean after 2 round(s) - round 1 (REJECT): ROUND MIS-SCOPED — patch not actually reviewed. This round's /simplify, checklist, and /code-review executed with cwd = the agent-flow framework repo (branch kleinc/code_quality_improvement_v1) and reviewed the code-quality-stage feature itself, NOT this task's PR patch. The ChatGLM3 code is on branch agent-team/chatglm3-6b-bringup-v1 in the TensorRT-LLM snapshot repo (a different directory), so `/simplify`/`/code-review` (which diff the cwd's HEAD) never touched it. Root cause: the Quality agent runs in the framework repo's cwd, not the snapshot repo; another round will misfire identically unless the Quality agent runs with the TensorRT-LLM task branch as its working directory. I re-reviewed the ACTUAL patch (git diff main in TensorRT-LLM). Real, unaddressed findings: - TESTS (duplication): `_patch_chatglm3_hf_tied_compat()` duplicated in all 4 new test files; `_find_cuda_graph_runner()` and `_assert_cuda_graph_hard_path()` duplicated across test_chatglm3_gsm8k.py and test_chatglm3_replay.py; CHATGLM3_CKPT const, skip decorators, and RuntimeCfg class repeated. Consolidate to a shared util/conftest (keep all coverage). - modeling_chatglm.py (simplification): redundant post-normalization reads of attention_bias/mlp_bias/head_dim via `getattr(...) or ...` after normalize_chatglm_config already set them; normalize_chatglm_config returns config but both callers discard it; a few verbose comments. - ALTITUDE (larger, behavior-sensitive; flagged not applied): custom imperative load_weights vs framework BaseWeightMapper; ChatGLM-specific normalization special-cased in shared config_utils.py; partial_rotary_factor/rope_fusion hardcoded. No fixes applied: the round did not validly review this patch, and edits to modeling/test code cannot be validated against the GSM8K accuracy criterion in this environment (risking a criterion regression). Recommend fixing the Quality agent's working directory (must be the snapshot repo/task branch) before the next round, then applying the test-dedup and post-normalization simplifications above. - round 2 (REJECT): Round 2 finally assessed the ACTUAL task patch (git diff main...agent-team/chatglm3-6b-bringup-v1 in the TensorRT-LLM repo), unlike round 1 whose skill steps mis-scoped to the agent-flow framework repo (cwd). Note: the automated stage still mis-scopes — /simplify and /code-review diff the cwd (agent-flow), not this task's snapshot repo — so I targeted the correct repo manually. /code-review (correctness): CLEAN — returned []. No correctness bugs survive verification. The high-risk items are validated: partial-RoPE, QKV/GQA split sizes (32 heads / 2 KV / head_dim 128), and transformer.* weight-name mapping are exercised by the GSM8K-parity acceptance gate QA approved; the fused-residual path is completed by the final self.norm(hidden_states, residual); embedding matches the Qwen sibling and load_weights handles attention-DP (tp_size = 1 if enable_attention_dp else mapping.tp_size). Speculative "other ChatGLM variant" concerns (_set_missing fill-if-None, head_dim None if kv_channels absent, missing dense/output biases) are not reachable for chatglm3-6b and are guarded by strict _take/consumed accounting. /simplify (round 2): applied one behavior-preserving cleanup (dropped the unused `return config` from normalize_chatglm_config, annotation -> None); it was subsequently reverted in the working tree, so the patch is unchanged. OUTSTANDING (why REJECT): the checklist's test item is not satisfied — heavy test-helper duplication across the 4 new test files (`_patch_chatglm3_hf_tied_compat` in all 4; `_find_cuda_graph_runner`/`_assert_cuda_graph_hard_path`/`RuntimeCfg` across gsm8k+replay; repeated CHATGLM3_CKPT/skip decorators), ~100+ duplicated lines to consolidate into a shared util; plus minor post-normalization config-read redundancy in modeling_chatglm.py. NOT applied here because tensorrt_llm is not importable in this environment and there is no GSM8K/test rig, so a test-file move cannot be validated for collection/pass (risking the acceptance criterion), and the working-tree edit I did make was reverted. RECOMMENDATION: this cannot be resolved by another automated round as currently wired — fix the Quality agent's working directory so /simplify and /code-review target the snapshot repo/task branch, then apply the test-helper de-duplication with the test rig available (or the author applies it directly). Correctness and framework-conformance are clean; the residual is a non-behavioral maintainability cleanup. - round 1 (REJECT): Code-quality round 1 for the chatglm3-6b bringup patch (modeling_chatglm.py, config_utils.py, _torch/models/__init__.py + 4 new test files). /simplify applied (safe, behavior-preserving, verified against repo convention): - modeling_chatglm.py: consume the normalized config fields instead of re-deriving booleans inline — `bias=config.attention_bias`, `dense_bias=config.mlp_bias`, GatedMLP `bias=config.mlp_bias`, and `head_dim=config.head_dim` at both sites (attention + load_weights). Confirmed no base infra reads these fields and that the model already hard-depends on normalize_chatglm_config having run; matches Llama/Phi3/Nemotron/Parakeet. - Both integration tests: removed the unreachable BFS-over-object-graph fallback in `_find_cuda_graph_runner`, keeping the deterministic `_executor.engine.model_engine.cuda_graph_runner` chain (verified in source; the runner is at a fixed location in single-process mode and in a subprocess otherwise, where BFS also fails). Codex checklist pass: the working tree carried minor docstring/comment trims in the ChatGLM3 test files (removed helper docstrings on `_patch_chatglm3_hf_tied_compat`, `_pick_metric_key`, `greedy_generate`, and an inline comment); no behavioral change. /code-review (high) found 10 items. Fixed this turn: (NVIDIA#4) `_pick_metric_key` in test_chatglm3_gsm8k.py now raises loudly instead of silently gating on an arbitrary non-exact_match metric — strict hardening, no change to any real gsm8k run. OUTSTANDING (why not auto-fixed): - CORRECTNESS, latent/out-of-target: (NVIDIA#1) load_weights only handles chatglm3-6b's bias layout — it unconditionally fetches the QKV bias and never loads dense/MLP biases, so ChatGLM2/3 variants with add_qkv_bias=False or add_bias_linear=True would KeyError or trip the strict unconsumed-weight ValueError. (NVIDIA#2) normalize_chatglm_config never maps ChatGLM `rope_ratio` to `rope_theta` (get_hf_rope_theta reads config.rope_theta), so chatglm3-6b-32k/128k silently get theta=10000 and degrade at long context. Both are no-ops for base chatglm3-6b (add_qkv_bias=True/add_bias_linear=False/rope_ratio=1) so QA passed; fixing them adds untestable variant paths (no GPU/checkpoint here) and is scope creep on a base bringup — hand to the coder to fix+test or narrow the module's "ChatGLM2/3" claim. - TEST ROBUSTNESS: (NVIDIA#3) the CUDA-graph hard-path tests require TLLM_WORKER_USE_SINGLE_PROCESS=1 but only name it in an assert string and never set it; without it, ENABLED-cfg tests hard-fail on `assert runner is not None` and BASELINE silently skips its checks. Not auto-fixed because gating/skipping would weaken the author's intended hard-path assertion — needs a CI-env or fixture decision. - ALTITUDE/REUSE (larger refactors, unsafe without test execution): (NVIDIA#5) hand-rolled load_weights + consumed-set + ignorable-suffix lists reimplement BaseWeightMapper/register_mapper (cf. Glm4WeightLoader in modeling_glm.py); (NVIDIA#9) normalize special-cased in shared load_pretrained_config + duplicated in __init__ vs a ChatGLMConfigLoader; (NVIDIA#6) `_patch_chatglm3_hf_tied_compat` (4 copies) + RuntimeCfg/_build_llm/_find_cuda_graph_runner/_assert_cuda_graph_hard_path (2 copies) duplicated across two separate, GPU-gated test trees; (NVIDIA#10) _GreedyHFLM._model_generate hand-rolls HF greedy decode. - EFFICIENCY: (NVIDIA#7) test_chatglm3_source_activation_replay reloads the 6B HF model + rebuilds TRT model for both parametrized scenarios though only the trailing cuda_graph block differs (use the _HFRef cache pattern). - SCOPE CREEP: (NVIDIA#8) the diff reflows unrelated bart/minimaxm3/qwen imports (+ a mistral import in config_utils.py); left as-is because formatter ownership is ambiguous (line 33 is 89 chars vs isort/yapf's 80) and reverting risks fighting the pre-commit hook. Rejecting: a substantive edit was made this turn and several actionable findings remain (test env-var robustness is the most important; the two latent correctness gaps and the altitude/reuse refactors should be triaged by the coder). - round 2 (APPROVE): Code-quality round 2 for the chatglm3-6b bringup patch. /simplify (round 2) applied 3 trivial, behavior-neutral cosmetic cleanups in modeling_chatglm.py: inlined the single-use `head_dim` local in ChatGLMAttention (`head_dim=config.head_dim` passed directly to super()), inlined the single-use `head_dim` temp in normalize_chatglm_config, and dropped the noise `: set` annotation on `consumed`. Reuse/efficiency/altitude agents found nothing new safely-actionable and confirmed the two "small fold" hypotheses don't hold (base `skip_modules` matches model module names, not the checkpoint keys `_IGNORABLE_*` filters; the second normalize_chatglm_config call is load-bearing for the direct-ModelConfig test path) — the only remaining altitude/reuse items are large refactors. Codex checklist pass (round 2): no net changes — independent verification confirmed the only diffs since round 1 were the 4 quality edits above/below, all behavior-neutral. /code-review (high) re-surfaced the same stable finding set (the substantive code was unchanged; no coder commits landed between rounds). I acted on the most actionable one this turn: - FIXED (NVIDIA#3, test robustness): the CUDA-graph hard-path tests required TLLM_WORKER_USE_SINGLE_PROCESS=1 (their ENABLED config hard-asserts the in-process runner) but never set it, so they would spuriously fail/silently-skip when run with a checkpoint outside a harness that exports it. Added an autouse `monkeypatch.setenv` fixture to both integration files (test_chatglm3_gsm8k.py, test_chatglm3_replay.py). Verified the env is read at runtime (utils.py:380, os.environ.get inside a function) so the fixture takes effect before _build_llm; the fix aligns each test with its own documented requirement and cannot regress the already-validated path. - Previously fixed (round 1): _pick_metric_key now fails loudly instead of gating on an arbitrary metric; consumed bias/head_dim inline derivation replaced with normalized fields; BFS runner-finder fallback removed. RESIDUAL (documented follow-ups — not safely resolvable by the automated quality loop): - Out-of-target latent correctness: (NVIDIA#1) load_weights only handles the chatglm3-6b/chatglm2-6b bias layout (add_qkv_bias=True/add_bias_linear=False) — it unconditionally fetches the QKV bias and never loads dense/MLP biases, so ChatGLM variants with add_qkv_bias=False (KeyError) or add_bias_linear=True (unloaded bias + unconsumed-weight ValueError) fail; (NVIDIA#2) rope_ratio is never mapped to rope_theta, so long-context chatglm3-6b-32k/128k (rope_ratio>1) silently use theta=10000. Both are no-ops for the bringup target (base chatglm3-6b/chatglm2-6b, rope_ratio=1); fixing needs variant checkpoints to test and expands scope beyond the bringup — for the coder/a follow-up. - Larger altitude/reuse/efficiency refactors needing test-execution access: (NVIDIA#4/NVIDIA#7) adopt BaseWeightMapper/register_mapper for load_weights + a ChatGLMConfigLoader/config-subclass to remove the shared-runtime normalize special-case and its second call; (NVIDIA#5) consolidate the 4x/2x-duplicated test helpers across the two GPU-gated test trees; (NVIDIA#6) cache the 6B HF/TRT model in the parametrized attention test; (NVIDIA#8) the _GreedyHFLM hand-rolled reference decode. APPROVE rationale: base chatglm3-6b is fully correct (QA-approved, acceptance criteria met); two thorough rounds applied all safe in-scope polish culminating in the NVIDIA#3 test-robustness fix; the quality loop has converged (round 2 reproduced round 1's findings with no coder changes); and every remaining finding is either out-of-target latent correctness for untested ChatGLM variants or a larger refactor requiring test-execution access — none of which another automated /simplify→/code-review round can safely resolve. Residuals are surfaced above for human/coder triage. QA report: DECISION: APPROVE. All 8 acceptance criteria hold at runtime against the current committed code (HEAD 2323424, iter11). Reference is the checkpoint's own trust_remote_code modeling_chatglm.py — an independent HF reference sharing zero code with the TRT-LLM path. PER-CRITERION (runtime evidence): 1. Bootstrap chain — PASS. slurm/core_gpu_tests.4147161.log: "Built target build_wheel_targets", "Successfully installed tensorrt_llm-1.3.0rc21 transformers-5.5.4" (pip install -e .[devel]), "Claude Code successfully installed!", "claude-agent-sdk 0.2.93"; all inside srun pyxis container with --container-mounts=<repo>:<repo> --container-workdir=<repo> under set -eo pipefail (tests ran only because bootstrap rc=0). 2. config_registration_and_weight_accounting — PASS (TIER1 rc=0). Real ckpt loads unedited, arch resolves to ChatGLMForCausalLM, TRTLLM backend + KVCacheManagerV2, strict state-dict accounting (raises on missing/unexpected/shape-mismatch; only rotary_pos_emb.inv_freq tolerated). 3. source_activation_replay — PASS both graph modes. layer0 cosine=1.000000 mean_abs=0.00000, layer27 cosine=0.999999; decode graph-vs-eager max_abs=0.0 cosine=1.0. Companion partial_rope_boundary_and_theta PASS (dims [0:64] rotate, [64:128] pass-through, theta=10000^(-i/32)). 4. source_logit_replay — PASS both configs. trt_tok==hf_tok (5231/30910/4802), cosine~0.99999; enabled num_captured_graphs=8, baseline 0. 5. generation_parity — PASS both configs. 5 prompts x 32 steps, token_mismatches=0, min per-step cosine~0.99998. 6. llm_api_smoke — PASS both configs (nonempty deterministic; V2 + TRTLLM asserted; enabled enabled=True graphs=8, baseline graphs=0). 7. gsm8k accuracy_canary — PASS (slurm/gsm8k_gate.4145862.log): HF=40.00, baseline=40.00 gap 0.00, enabled=40.00 gap 0.00; enabled num_captured_graphs=32. 8. gsm8k full_trtllm_eval — PASS: HF=52.77, baseline=52.99 gap 0.23, enabled=53.07 gap 0.30 (both < 2.0 tol), enabled captured 32 CUDA graphs; RESULT: PASS. Absolute ~53% matches published ChatGLM3-6B GSM8K, ruling out a shared-bug artifact. PROVENANCE: Core log 4147161 started 05:55, after the 05:33 HEAD commit -> definitively iter11 (criteria 1-6). git diff iter10->iter11 (2d58cc7..2323424) touched ONLY the two test files (additive CUDA-graph hard-path assertions); tensorrt_llm/ is byte-identical. GSM8K log 4145862 prints the iter11-only [cuda_graph_hard_path] output -> it ran the current test file against the unchanged runtime code (live bind-mount imported iter11 mid-bootstrap). Criteria 7-8 are therefore verified against current code. INDEPENDENT RERUN THIS TURN: I resubmitted both jobs. Core (4149172) matched the prior post-commit run and was cancelled as redundant. My GSM8K resubmit (4149173) FAILED in ~1.5 min inside build_wheel.py (FileNotFoundError: inherit_graph_6.md5 in copy_resolving_symlink) — a concurrent-build race from my own simultaneous submission of two bootstrap-heavy jobs on one live repo mount, NOT a ChatGLM defect. Accepted evidence is the sequential post-commit runs + full code/provenance review. RED-TEAM (all addressed): independent HF reference (no shared helpers); partial GPT-J RoPE covered analytically + behaviorally (layer0 exact); CUDA-graph silent fallback ruled out by in-process CUDAGraphRunner introspection (runner.enabled + >=1 torch.cuda.CUDAGraph; baseline 0); KV-cache/mask/scale drift ruled out by 32-step parity (0 mismatches) + GSM8K absolute-score match; transformers 5.x bridges (all_tied_weights_keys={}, max_length, num_hidden_layers) are loader-compat only and do not alter ChatGLM forward semantics. EVALUATION SCORES: - Functionality (2.0): 9 — every criterion passes with exact argmax parity, cosine >0.9999, GSM8K within 0.3 pts of HF on both configs. - Code Quality (1.0): 8 — clean, well-documented 373-line model reusing existing modules; narrow config hook; strict weight accounting; minor incidental import reformatting in __init__.py; a pre-existing (non-diff) ruff E501 at __init__.py:130. - Performance (1.5): 8 — parity-first bring-up (no perf gate in task.yaml); production TRTLLM backend + CUDA graph + overlap scheduler + fused QKV/gate-up + KVCacheManagerV2 all working. - Completeness (1.0): 9 — complete, buildable, no placeholders, full state-dict accounting, both runtime configs. - Technical Sophistication (1.0): 9 — correct unfused partial GPT-J RoPE with TRTLLM backend + CUDA-graph capture, compact MQA KV, config/transformers bridges. Weighted = (9x2.0 + 8x1.0 + 8x1.5 + 9x1.0 + 9x1.0)/6.5 = 56/6.5 = 8.6. STRENGTHS: genuine independent HF reference; real CUDA-graph hard-path proof (not just config flags); strict weight accounting; large GSM8K margin; minimal composable design (no new schemas). WEAKNESSES: multi-GPU/perf/chunked-prefill deferred (out of scope per task.yaml); harness must run the two jobs sequentially to avoid the concurrent-build race; pre-existing E501 worth a one-line fix. RECOMMENDATION: APPROVE — no code defect identified; task.yaml completion_criteria (GSM8K within 2 pts with and without CUDA graph + overlap scheduler) is met on both configs.</summary> </invoke>

… budget fraction Spilling canonical spans LAST is not spill protection: under sustained pressure at the arena's capacity edge, transient page needs (map-ahead margins, growth bursts) still consume the whole P3 registry -- and with it every zero-copy alias hit -- to cover needs worth a handful of pages. Measured on the 8B shared-prefix benchmark at memory-equalized quotas: all 15 spans (~60 pages ~= 1 GiB) pressure-spilled, alias hit rate 9% vs 87% when capacity allows the registry to survive, costing ~5 points of the ~7-point aliasing upside. Fix: exempt the registry from pressure spills up to ContiguousArenaConfig.protected_span_fraction of the page budget (default 5%; TRTLLM_KV_ARENA_SPAN_PROTECT_FRACTION overrides; 0 restores unprotected spilling). The excess above the floor stays spillable LRU-first, spans-after-ranges as before. Callers that cannot make progress back off like any failed allocation (DESIGN.md paragraph 4.6) instead of consuming the registry. The resume-utilization gate stops counting protected span pages as reclaimable-at-will: they are budget-retained (the charge-transfer accounting is unchanged) but pressure spills now refuse to touch them, and counting them available would admit resumes whose mapping then fails. Defaults are behavior-preserving where it matters: aliasing itself remains OFF, and without registered spans the floor is inert (arena regression suite unchanged; 214 passed / 12 skipped). Tests: floor survives arbitrary pressure while parked ranges spill; excess above the floor spills oldest-span-first; fraction=0 restores the previous behavior. Arena suite 103/103 on H100. [None][fix] Report the arena's physical page budget as total_quota, not VA CacheLevelStorage.total_quota sums pool byte sizes; arena pools are views over the VA reservation, so the un-overridden property reported virtual address space as allocated memory. The KV memory estimator credits total_quota back as "temporary pool that will be freed" (_util.py _cal_max_memory), so the arena's final KV quota was over-granted by the VA-vs-physical difference: on an 80 GiB H100 with free_gpu_memory_fraction=0.85, 60.81 GiB instead of the correct ~52.9 GiB (paged v2: 52.55) -- committing the device to within 80 MiB of its capacity at full budget fill, where cuMemCreate failure is a hard error, and silently giving every arena benchmark ~8 GiB more KV than its paged baseline. Fix: GpuArenaCacheLevelStorage overrides total_quota to page_budget.total_pages * phys_page_size. Unit test pins the physical semantics (VA extent 3x the physical quota must not leak through). Arena suite 100/100, pre-commit clean. Corrected, memory-equalized benchmark numbers (Llama-3.1-8B solo, 2500x(1024/1024)@conc256, both backends now ~52.5-52.9 GiB quota): paged 8,563 tok/s; arena+prefix-aliasing 7,816/7,844 (-8.6%), with alias hit rate collapsing to ~9% (was 87% at the inflated quota) and all 15 canonical spans pressure-spilled. The previously reported -1.3% depended on the over-granted quota: at honest capacity the level sits at the concurrency edge ((3,385-60 registry pages)/13 charged pages/seq = 256 exactly, before the per-sequence map-ahead margin), so pressure spills thrash the registry. Follow-ups tracked in the work log: span-aware spill policy (the registry is ~1 GB -- spilling it saves little and forfeits the aliasing win), map-ahead margin accounting at the capacity edge, and re-benchmarking at per-configuration best rather than equal-fraction quotas. [None][fix] Fall back to copy when a canonical span dies before first resume A registry hit is stored UNPINNED on the admission across the admission->first-resume window; under pressure, spill_canonical_spans can run in between and drop the span's handle references to zero. The first resume then alias-mapped freed handles (caught by the SharedPhysMem.handle refcount assert on the tinyhost pressure reproducer). Fix: _CanonicalSpan.is_live() (all pinned handles still referenced); the onboard alias branch re-validates a hit before alias-mapping a FRESH range and falls back to the host-copy path on a dead span. Signature-adopted ranges are immune -- their inherited mappings hold their own references -- and keep the fast path. Regression test reproduces the exact window: lookup hit at admission, full spill (registry pins dropped), then resume -- must not crash, must not alias, and must deliver correct bytes from the intact host copies. Arena suite 99/99, adapter 12/12, v2 regression 210 passed / 12 skipped, pre-commit clean. [None][fix] Alias partial canonical spans so shared-prompt matches engage The v1 aliasing lookup required a reuse match to cover the FULL registered span -- but spans are registered at owner close with the owner's whole committed chain (shared prompt + its unique continuation), while later admissions match only the shared prompt. Every lookup missed and aliasing never engaged: 8B shared-prompt solo benchmark came back +0.2% (noise) with the feature ON. Fix -- partial-span aliasing: - lookup_canonical_span identity-verifies only the MATCHED prefix and returns (key, span, usable_blocks): the match trimmed to whole chunks in every pool. A shorter-than-span match aliases exactly its own covered chunks; the adopter writes strictly past them, so the owner's tail chunks are never touched. - alias_span_into_range takes the extent and trims the shared handles per pool (SequenceArena.trim_shared_to_blocks; base-independent aliasable_span_blocks for pre-reserve arithmetic). - The adoption signature becomes (key, extent) -- extent-exact, so an admission with a shorter match can never adopt a longer-aliased range (whose extra aliased chunks its writes would corrupt). The owner's parked range carries (key, full_span) and is typically only spilled. - Alias hit/miss/spill counters now log at shutdown next to the ghost counters, so benchmark engagement is observable. New end-to-end test reproduces the benchmark shape: owner commits 128 blocks (2 chunks), an adopter matching only the shared 64 blocks aliases exactly 1 chunk and grows past it, and a full-chain match then adopts the owner's range with its tail chunk byte-intact -- proving a shorter alias cannot corrupt the shared span. Arena suite 98/98, adapter 12/12, v2 regression 209 passed / 12 skipped, pre-commit clean. [None][feat] Add physical-page aliasing for shared prefixes (P3 prototype) Zero-copy prefix reuse for the contiguous KV arena: cuMemMap the SAME physical pages holding a shared, fully-committed prefix into multiple sequences' VA ranges. Kernel-transparent (each sequence's offsets point into its own contiguous VA); the prefix's copies, budget charge, and DRAM/L2 duplication all vanish. Motivation (owner hypothesis, confirmed by discriminator): the 8B arena-vs-paged gap is dominated by prefix DUPLICATION, not copy cost -- paged shares prefix blocks by refcount and fits the GPU quota eviction-free, while arena's private copies oversubscribe it ~1.28x on the prefix-sharing benchmark. Reuse-off (identical footprint) collapses the gap from -6.1% to -1.6%. Design (P3_DESIGN.md; validated by bench_alias_map.py -- multi-mapping works on H100/580.x and an alias map is CHEAPER than a fresh one): - SharedPhysMem: refcounted pooled handle; VirtMem._page_map holds it uniformly, map_alias() maps an existing handle at a second VA, and unmap/destroy drop references -- the handle closes strictly on the last drop, so range reclaim stays uniform (no skip-unmap cases). - SequenceArena.alias_prefix() maps a canonical span at a range's start (before growth maps) charging NOTHING; charged_pages_in_range() excludes aliased chunks from park/adopt/spill/reclaim accounting. Only whole, fully-committed chunks alias (aliasable_prefix_blocks); the partial boundary page falls back to the copy path. - Canonical-span registry (ArenaPoolGroup, ghost-registry identity discipline): a closing owner's committed prefix pages are pinned (budget charge transfers range->registry, retained). Lookup requires the reuse match to cover the FULL span -- a shorter match would write into shared pages. - PREFIX-AFFINE ADOPTION (load-bearing, not an optimization): adoption recycles parked pages by overwriting them, which would corrupt every alias of a shared span. Parked ranges carry an alias signature and are adopted only by same-key admissions -- which want exactly those bytes and never write the prefix (reuse skips it). Same-key adoption is the zero-everything steady state: no map, no unmap, no copy. - Onboarding fast-path: registry-hit blocks skip the H2D/D2D copy entirely; the existing PrivateCommittedPage carries them (close drops without write-out, suspend swaps to canonical -- exactly the aliased semantics), so no new page type is needed. Consumers wait the span's ready event. - Spill: registry pins are reclaimable last (hottest bytes); live aliases keep their handles via refcounts and later admissions fall back to the host-copy path. Off by default: ContiguousArenaConfig.prefix_aliasing / TRTLLM_KV_ARENA_PREFIX_ALIASING=1. Tests (+7): VirtMem alias/refcount matrix (share bytes, survive owner unmap, destroy drops refs, collision asserts); SequenceArena zero-charge roundtrip + boundary-page trim; ArenaPoolGroup registry roundtrip (register/lookup/alias/park/affine-adopt/spill with exact charged accounting); manager-level end-to-end (zero-copy proven by corrupting every host copy, affine adoption across three same-prefix admissions, registry spill falls back to host). Arena suite 97/97, adapter 12/12, v2 regression 208 passed / 12 skipped (zero regressions), pre-commit clean. V1 scope notes (documented in code): single-life-cycle pool groups; aliasing engages at owner close (concurrent same-prompt admissions before the first close still copy); SSM/partial-block interplay unchanged. [None][fix] Degrade arena suspension to drop-and-recompute on host-tier exhaustion Host-slot exhaustion during a suspension write-out was process-fatal: OutOfPagesError escaped kv_cache.suspend() through the scheduler and killed the executor loop. The host tier CAN legitimately run out — its capacity fills with already-suspended sequences' HELD pages, which LRU eviction correctly refuses to drop at the last level (is_evictable), so "all host pages are evictable" stops holding exactly when suspension pressure is highest. Worse, suspend() mutated state (page-index buffers cleared, locks converted to holders, scratch freed) BEFORE the fallible host allocation, so the error could leave a half-suspended sequence. Fix, four layers: - _KVCache.suspend() is now atomic on failure: a pre-flight prepare_free_slots reservation (_arena_evacuation_requirements, side-effect-free mirror of the evacuation partition) runs before any mutation; OutOfPagesError leaves the sequence fully ACTIVE. The manager is single-threaded, so the reserved slots stay free until the per-pool-group offloads consume them. - Adapter suspend_request() returns bool and degrades to drop-and-recompute (v1-style preemption) via free_resources — close() is already exhaustion-safe — instead of raising. The two internal suspend sites degrade appropriately: the ctx pre-resize revert-suspend (an optimization) stays resident; the first-context-chunk backoff (no computed KV worth preserving) drops outright and prepare_context recreates the cache with a fresh reuse lookup. - V2 scheduler: a dropped victim still counts as evicted (its GPU pages are freed either way, keeping the deadlock guard satisfied); main and draft caches drop together; the request is flagged py_kv_dropped. - Executor: v2 loops now pause exactly the flagged drops, terminating first so the seq slot is released (a paused request re-enters as context and add_slot asserts on a duplicate ID); surviving canonical prefixes in the reuse tree still shorten the re-prefill. Tests: suspend raises atomically when the host tier is pinned by HELD pages (victim stays byte-identical, still grows, and close() never errors on the free path); stale DROPPABLE copies keep churning through LRU under the same pressure (suspend succeeds after eviction); adapter drop fallback end-to-end (first suspension offloads and stays resumable, second drops and cleans the map). Arena suite 90/90, adapter 12/12, v2 regression 201 passed / 12 skipped (identical to pre-change plus the new tests). E2E (Llama-3.1-8B, free_gpu_memory_fraction=0.42, host_cache_size=4GiB ~10x under suspension demand, conc 256): unfixed corrupts state at the first mass-suspension wave; with the fix the drop path executes 40+ drop-and-recompute cycles cleanly. KNOWN OPEN ISSUE (pre-existing, reproduced on the unfixed tree): a separate non-deterministic IMA can still fire at the first suspension wave in this extreme-churn regime (~2/3 of runs, ~25s, before any drop executes) — under investigation; reproducer: the config above. [None][feat] Adopt freed arena ranges instead of unmap/remap cycling Freed sequence ranges are now parked still mapped once their reclaim gates complete, and reserve_sequence hands a parked range whole to the next admitted sequence (LRU-first, size-fit; ranges with live ghost entries are skipped -- they are worth more as D2D reuse sources -- and stale ghost keys are pruned so they cannot block adoption forever). At steady state the arena issues no cuMemMap/cuMemUnmap at all: adopted ranges keep their mapped prefix and budget charge, and actual unmaps happen only under pressure spill and at teardown, where the (default on) pre-unmap quiesce is rare and bounded. Default on; TRTLLM_KV_ARENA_RANGE_ADOPTION=0 disables. This is the production fix for the confirmed driver-interference class: cuMemUnmap concurrent with in-flight kernels reading other, still- mapped ranges of the same VA reservation faults those reads at the device MMU level (H100, Llama-3.1-8B, overlap scheduler). The 42%-quota pressure reproducer crashes ~4/7 runs without the quiesce and passes 6/6 with it; adoption removes the churn instead of serializing it. A first adoption attempt failed validation 4/4 with the same fault signature; those crashes are re-attributed to the (since fixed) host block-offset staging race -- adoption widened that race's window by making admission instant. On top of the staging and evacuation-ordering fixes, adoption now passes the full pressure matrix: 42%-quota 300- request runs 3x plus a 2500-request endurance with defaults, and 2x with the quiesce disabled; at 85% quota / 2500 requests it improves plain-arena throughput ~5% over non-adoption on-demand mapping (8,041 vs 7,655 tok/s, Llama-3.1-8B) by eliminating admission map costs. Composes with the scheduling-pass blocking drain: drain_reclaim(wait= True) waits out a pending range's gates and parks it, the parked pages count as available to the resume utilization gate, and the retrying allocation can adopt the parked range directly. Lazy-retention D2D ghosts compose unchanged. Tests: adoption round-trip (zero new maps, budget/retained accounting), unfit/unready-range skip, adoption- disabled restores unmap-on-drain; parked-semantics assertions added to the drain/spill call sites across the arena and adapter suites. [None][fix] Order arena write-outs after the in-flight forward step A sequence's KV writes run on the forward stream, but the suspend and close write-out copies were ordered only by the pages' ready events and the sequence's manager-stream events -- neither covers a speculatively enqueued overlap-scheduler step that is still appending to the tail block when the scheduler evicts the sequence (or when a finishing request is freed). The batched-copy kernel reads its sources at execution time, so the host tier captured the tail block PRE-write: stale bytes restored on resume, and stale reuse copies from close. Demonstrated end to end -- a delayed tail-block write on a foreign stream is lost through a suspend/resume round-trip without this fix. Thread a caller-supplied prior event through suspend()/close() -> _arena_evacuate -> offload_arena_pages -> _batched_migrate (extra_prior_event), following the write_through_pages precedent. The adapter records the event on the executor thread's current stream (= the forward stream) at each suspend/close site, so it orders after every previously enqueued step; this iteration's forward is not yet enqueued and cannot reference the evicted/closed sequence. The range unmap was already safe (the risk-#3 reclaim fence is a forward-stream event); only the copy ordering was missing. Known remaining gap of the same class (out of scope): write-through-on- commit (opt-in, non-default) records its prior on the sequence's manager stream; the committed block's last tokens may likewise still be in flight. Needs the same treatment via the commit path. New tests: a storage-level test proves extra_prior_event gates the batched migrate against a pending foreign-stream write (with a warmup copy -- the kernel's first-call module-load latency exceeds the race window and masks it), and an end-to-end test proves a still-in-flight tail-block write survives a suspend/resume round-trip under both write-through policies. [None][fix] Recover from arena page exhaustion inside a scheduling pass Arena frees are event-gated AND fence-gated: a freed range is unmapped (returning its pages to the budget) only after its evacuation copies finish and a fence recorded on the execution stream at a drain point covers it (risk #3). Both normally happen at prepare_resources -- which runs AFTER scheduling. So pages freed by the V2 scheduler's own evictions could never be reclaimed within the same scheduling pass: _try_evict_for_gen suspends a victim, retries allocation, fails again (unfenced ranges are never reclaimable and the utilization gate still counts the freed pages as used), and moves to the next victim -- a mass-suspension cascade. Terminal state: every generation request suspended, all pages parked in unfenced pending-reclaim, every resume() refused by max_util_for_resume, nothing left to evict, and the scheduler's deadlock guard raises fatally. Under on-demand mapping (margin=1) with 256 growing 8B sequences this hit ~1 in 4 of 2500- request runs at 85% quota and every run at 42% quota. Add a blocking variant of the reclaim drain (wait=True through ArenaPoolGroup / cache-level storage / StorageManager / KVCacheManager): with a fence supplied, it synchronizes each pending range's fence, free event, and released-slot gate events instead of skipping in-flight ranges, then reclaims. Ranges with outstanding slots are skipped (their release is host-side state, not an event). The wait is bounded by the just-enqueued evacuation copies plus any in-flight forward step, and is only entered when the alternative is a futile cascade or a fatal deadlock. try_allocate_generation retries once through the blocking drain on both failure paths (resume refused, resize failed). Recording the fence mid-scheduling is safe by the same argument as _reclaim_fence: the scheduler runs on the executor thread whose current stream is the forward stream, so the event orders after every previously enqueued step, and this iteration's forward (not yet enqueued) cannot reference freed sequences. Eviction now frees pages in-pass, so pressure suspends the minimum number of victims instead of the whole batch. Validation: new pool-group tests cover the blocking drain (waits out a pending gate event; must not bypass outstanding-slot or unfenced protection) and a new adapter test reproduces the exact terminal state (single request holding 100% of a two-page budget, suspended, bare resume() refused) and shows try_allocate_generation recovers. At 42% GPU quota / concurrency 256 on Llama-3.1-8B (2.4x oversubscription, constant preemption), the unfixed tree fails 2/2 while the fixed tree completes 3x 300 and a 2500-request run cleanly. [None][fix] Revert "Pre-map whole arena ranges at admission by default" This reverts commit 669384b3efc385c2bd12b1a6db927fb1f4b9bbf9, restoring on-demand demand-paging (map-ahead margin 1) as the arena default. Pre-mapping was introduced as the fix of record for the 8B arena illegal-memory-access on the theory that growth-time cuMemMap/cuMemSetAccess into a range concurrently read by in-flight kernels triggers driver-level MMU-fault interference. The preceding commit re-attributes that crash to a CPU-side host-buffer reuse race in copy_batch_block_offsets (nvbug 6293536 class): a discriminator matrix on the reliable crash config showed the staging fix passes 3x at 300 requests and at 2500-request scale on every historical crash flavor (linear and plain, 16 and 64 MiB pages) with on-demand mapping enabled, while a masking-control variant with identical per-call CPU cost but a persistent (still racy) staging buffer crashes with the original signature. Pre-mapping "worked" by converting the stale-row illegal access into silent reads of valid mapped memory, at a ~3 point throughput cost from full-range page-budget charging at admission. On-demand mapping is required for production use (sequence lengths vary too widely to pre-charge every sequence's maximum), and with the staging fix in place it is no longer implicated in the crash. [https://nvbugs/6293536][fix] Stage v2 KV block offsets through fresh host buffers Feature-branch adaptation of the standalone fix proposed for main in PR #16088 (cherry of 3bdee3b3695c): copy_batch_block_offsets_to_device is an asynchronous cp.async gather kernel that reads its host inputs at execution time, and two of those inputs are reused in place across iterations (host_kv_cache_block_offsets rows rebound on slot reuse, and copy_idx, a slice of the IndexMapper's persistent copyIndex_ buffer). Under the overlap scheduler the CPU runs an iteration ahead and can clobber either input before the previous iteration's still-pending kernel drains, so the kernel gathers another batch's physical blocks: silently wrong reads in classic paged mode, and illegal memory accesses in arena mode when a stale row points into unmapped arena VA. Snapshot the rows each call needs into a fresh pinned buffer, feed the kernel an identity index, and retire each buffer only once a CUDA event recorded after the kernel completes (the caching host allocator's keep-alive does not cover custom kernels reading host memory). Differences from the main-branch PR: the gather stays enqueued on the caller's current (forward) stream per this branch's earlier ordering fix, so the retirement event is recorded on that same stream. On this branch the race is the confirmed root cause of the 8B arena illegal-memory-access previously attributed to driver-level cuMemMap/cuMemSetAccess interference: with this fix, the reliable on-demand-mapping crash config (margin=1, quiesce disabled, 300 and 2500 requests, linear and plain flavors, 16 and 64 MiB pages) passes with zero IMAs, while a masking-control variant with identical per-call CPU cost but a persistent (still racy) staging buffer crashes with the original signature. The new regression test fails without the fix and passes with it. [None][fix] Pre-map whole arena ranges at admission by default Scale testing showed the reclaim quiesce alone is not sufficient: the same H100 MMU-fault interference that reproduces at 300 requests with small map-ahead margins also reproduces at 2500 requests even with unmaps quiesced. Cross-referencing every configuration tested (margins 1/4/18, quiesced and unquiesced unmaps, with and without linear kernels, 16/64/128 MiB pages, 100 to 2500 requests) leaves a single consistent trigger: growth-time cuMemMap/cuMemSetAccess into a range that in-flight kernels are concurrently reading. Configurations that pre-map each sequence's whole range at admission never crashed at any tested scale (3x300 plus 4x2500 request runs); every configuration with growth-time maps eventually did. TRTLLM_KV_ARENA_MAP_AHEAD_PAGES therefore now defaults to -1 = pre-map the sequence's whole reserved extent on first touch (the mapping plan clamps the margin to the range). Admission-time maps target ranges no kernel reads yet and remain safe. Explicit small margins stay available for experiments; measured cost of pre-mapping on a memory-bound Llama-3.1-8B configuration is ~3% throughput versus demand paging, which is the right default trade for correctness. [None][fix] Quiesce the GPU before arena reclaim unmaps Reproduced on H100 (Llama-3.1-8B bf16, contiguous KV arena, overlap scheduler, requests > concurrency): decode kernels hit device-side MMU faults on ranges that are mapped and correctly fenced. A discriminator matrix isolated the trigger to cuMemMap/cuMemUnmap churn concurrent with in-flight kernels reading other ranges of the same VA reservation: - control (reclaim unmaps concurrent with kernels): crashes reproducibly at ~40s; - pre-mapping each range fully at admission (no growth maps): passes 3/3 plus two 2500-request endurance runs; - cuCtxSynchronize immediately before each reclaim unmap batch, with growth maps left concurrent: passes 3/3 (plus one more control run); - CUDA_LAUNCH_BLOCKING also masks it, and event-level ordering fixes do not help - consistent with driver-level TLB interference, not an ordering bug in this code. Unmapping is now preceded by a context synchronize by default (TRTLLM_KV_ARENA_SYNC_BEFORE_UNMAP=0 disables it for experiments). Reclaim happens once per admission wave, not per iteration, so the cost is bounded; correctness wins for an opt-in feature. Users can additionally set TRTLLM_KV_ARENA_MAP_AHEAD_PAGES to cover a whole range (e.g. 18 at 16 MiB pages for a 272 MiB range) to avoid growth-time maps as well. [None][fix] Gate arena reclaim on an iteration fence and order offset-table copies on the forward stream Two independent ordering hazards under the overlap scheduler, found while debugging an arena-mode illegal memory access on Llama-3.1-8B: 1. Reclaim fence (risk #3 of the deferred-reclaim design): with the overlap scheduler, step N+1 is enqueued speculatively before step N's sampled tokens reveal that a request finished; the already enqueued step still reads the finished request's range, and the free/gate events recorded at close do not order against it. Every freed SequenceRange now carries a reclaim fence - an event recorded on the execution stream at an iteration-boundary drain after the free - and is never unmapped before that fence completes. Internal pressure drains pass fence=None and can only reclaim previously fenced ranges; the pressure paths (resize_context, resume-and- restore, extend_capacity_for_tokens) fence and retry so capacity estimation cannot livelock when prepare_resources is skipped. 2. Offset-table copy stream: the v2 adapter enqueued its block-offset table refresh on the manager's side stream with no ordering against the forward stream (v1 uses a plain .copy_() on the current stream). A forward could read stale rows pointing at freed ranges: silently wrong reads in classic paged mode, illegal accesses in arena mode once the range is unmapped. The copy kernel is now enqueued on the caller's current stream, restoring v1-equivalent ordering. Tests: arena suite gains a reclaim-fence gating test (80 total) and the adapter suite three env-guard tests (10 total); drain call sites updated to the fenced signature. [None][feat] Add linear KV page addressing to XQA generation kernels (P1) Contiguous-arena sequences have consecutive KV cache pages: block-offset table entry j equals entry 0 + j * (num_layers_in_pool * kv_factor). Let the XQA HMMA generation kernel exploit this: load each sequence's base page index once and compute per-tile page indices arithmetically instead of loading every entry from the device page list (design 4.9 phase 1). - xqa device code: PAGED_KV_LINEAR / PAGED_KV_IDX_STRIDE compile-time flags; getLinearBasePage + getPageLinear helpers; K/V loadPages use them at beam width 1. Beam search, QGMMA, MLA and mmha are untouched. - Host plumbing: XQAParams.linear_kv_page_stride -> tllmXqaJitContext -> NVRTC macros (HMMA only); stride included in the cubin cache key and in AttentionOp::data() so op-cache entries cannot collide across models with different pool shapes; stride computed at op creation in the torch op (before runner->prepare()) from the pool mapping, so the prepared cubin matches the runtime dispatch key. - Opt-in via TRTLLM_KV_ARENA_LINEAR_KERNELS=1; both KV cache managers raise at construction if it is set without use_contiguous_kv_arena (the classic paged manager does not allocate consecutive pages). - Standalone xqa unit tests: RefCheck passes with the flag on; kernel perf -4.0% @ seq 512, -1.2% @ 2048, 0% @ 4096 (H100, identity page list in both builds, so the delta is purely the removed table loads). - End-to-end: token parity with paged and arena baselines (greedy, reuse round included, linear cubin verified active via compile log); throughput neutral on Llama-3.2-1B @ conc 256 on both sharing and unique-prompt datasets, where decode steps are host-bound and the attention kernel is a small slice of step time. The kernel-level win should surface for attention-heavier (larger/longer-context) decode. [None][fix] Fix mypyc-strict typing and compile issues in contiguous KV arena code Fixes in code this branch introduced, found while restoring the v2 package's mypyc build: - parameterize the bare cast(TypedIndexList, ...) for arena ranges; - rename loop variables that reused a narrower-typed name (slot in the onboarding copy loop and the OutOfPagesError rollback); - give arena_last_consumer the base CachedCudaEvent type via cast (its NULL initializer is the _NullCudaEvent subclass); - read _finish_event through an annotated local in close()/suspend() and assert suspend's entry invariant on a local: narrowing the member to None persists for the whole function under mypy(c), which cannot see _record_event's side effect, and compiled code then unboxes the later read as literal None; - move the _ArenaClosingPages NamedTuple to module level (mypyc cannot compile class definitions nested in a class body); - narrow arena config validation per layer with isinstance so the sliding_window_size access typechecks on the union; - type the batched-map-sweep bookkeeping (_pending_maps entries, exact delta charge) and the ghost registry's rawref. No behavior change; arena suite 79/79, adapter 7/7, v2 regression 61 passed/12 skipped, all unchanged. [None][perf] Track per-range mapped frontier in contiguous KV arenas A sequence range's mapped chunks are always the prefix [range start, frontier): every mapping plan extends from the current frontier, ranges are page-disjoint, and nothing maps arena pages outside ensure_mapped. Exploit this with a per-range per-pool mapped high-water frontier so growth planning (previously an O(pages) is_mapped rescan from the range start on every resize), retained-page accounting, and reclaim all work in O(1) chunks per pool. The _missing_runs/_unmap_mapped_run scanners are removed; the frontier advances as each map_range succeeds, so the partial-failure budget rollback is unchanged. Sharing benchmark (Llama-3.2-1B, H100, 5000x(1024/1024)@conc256, on_free/16MiB): 31,987 -> 32,022 tok/s; gap to same-node paged baseline -5.6% -> -4.8%. [None][feat] Add batched per-iteration map sweep to contiguous KV arenas (§4.2) Implement the design's batched map sweep: growth maps charge the page budget at resize time (admission control is byte-identical to synchronous mapping) but the cuMemMap/cuMemSetAccess calls are deferred and executed back-to-back once per iteration. Flag-gated by ContiguousArenaConfig.batched_map_sweep (default False) so direct users of the manager keep synchronous mapping semantics; the executor adapter enables it via TRTLLM_KV_ARENA_BATCHED_MAP and owns the flush points: prepare_resources (main and draft managers), add_dummy_requests (warmup and CUDA-graph capture run forwards without prepare_resources), and extend_capacity_for_tokens (runs after the per-iteration sweep). suspend() and close() flush before their write-outs read the sequence's pages, and reuse onboarding maps synchronously (its copies run immediately). Queue entries whose range is freed before the flush release their charge; per-range frontier dedup charges exact deltas for repeated growth within one iteration. Honest benchmark finding (shared-system-prompt workload, 5000 requests, concurrency 256, on_free/16MiB): the sweep is performance-neutral (31,918 vs 31,987 tok/s synchronous; paged baseline 33,885), as is a map-ahead ablation that reduces cuMemSetAccess to one call per sequence (31,775). Both match the microbenchmark arithmetic -- per-request driver savings are bounded well under the run-to-run noise. Together with the lazy-retention null result (99.7% D2D hit rate, no speedup), the residual arena gap on prefix-heavy workloads is distributed host-side bookkeeping across per-request paths rather than any concentrated sink; the remaining levers are mypyc compilation of the v2 package and moving the (now single-point) flush to a helper thread. The adapter therefore defaults the sweep OFF until the helper thread lands; the machinery is the ready-made offload point for it. 79/79 manager tests (defer-until-flush with delta charging, freed-before-flush charge release, onboarding-stays-synchronous); v2 regression suite identical to main; the sweep-enabled benchmark completed 5000/5000 requests, validating the flush protocol end to end. [None][feat] Add lazy GPU retention to contiguous KV arenas (P2, §4.4 phase 2) Freed sequence ranges now stay mapped on a retained LRU instead of being unmapped at drain time (opt-in: ContiguousArenaConfig.lazy_gpu_retention / TRTLLM_KV_ARENA_LAZY_RETENTION=1). Reuse hits whose blocks are still resident copy D2D from the retained bytes instead of H2D from the host tier; the page budget reclaims retained ranges only under pressure. - The close-path invariants are unchanged: canonical pages still move to the host tier (offload or write-through adoption), so retained ranges are pure 'ghost' caches -- a per-pool-group registry maps a canonical host page to the retained (range, ordinal) still holding its bytes. Sequence-private reuse copies re-register as fresh ghosts when their owner closes, so a hot prefix's D2D source is LRU-refreshed by every user rather than dying with its original committer's range. - Pressure paths spill LRU-first: page exhaustion in ensure_mapped drains then spills in chunks; VA exhaustion at reserve time spills before giving up. A range with an in-flight D2D onboard cannot be spilled (the copy's finish event gates its reclaim). The resume utilization gate counts retained pages as available, mirroring classic 'evictable' accounting (same livelock class as the estimation-phase fix). Teardown spills everything after synchronizing gates. - Ghost hit/miss/spill counters are logged at executor shutdown. Verification: 76/76 manager tests (retention matrix under both write-through policies; D2D proven by corrupting host copies and requiring original bytes; ghost refresh proven by spilling the oldest range; both spill paths; VA-pressure recovery). v2 regression suite identical to main. Honest benchmark finding: on the shared-system-prompt workload (5000 requests, conc 256) retention achieves a 99.7% ghost hit rate (20316/54) yet throughput is unchanged versus no-retention (-6.0% vs -5.6% against classic paged) -- proving the residual arena gap on prefix-heavy workloads is NOT the reuse copy transfer. The actual per-request admission cost (~1.8ms host time) is under separate investigation; retention remains the right mechanism for reuse-copy volume and its memory-traffic benefits. [None][fix] Handle intra-batch commit conflicts in contiguous KV arena mode (P0) Benchmarking with shared system prompts at concurrency 256 crashed arena mode within a few requests: when two in-flight sequences commit the same tokens, the loser's _commit_block finds an existing full tree block. Seq rebasing (adopting the other sequence's pages) is disabled in arena mode -- it would share GPU pages across arenas -- so the sequence fell into VIRTUAL_STOP, and commit()'s block loop, which only checks the commit state at entry, asserted on the next block. - _commit_block gains an arena branch for the conflict: keep this sequence's own (already-written) pages as sequence-private committed copies referencing the existing tree block, and keep committing. No copy is needed, the canonical entry is untouched, and the private copies are dropped at close as usual -- the same §4.4 principle applied intra-batch. UncommittedPage.convert_to_private_committed implements the conversion (identical to convert_to_committed, minus block.storage registration). - commit()'s block loop now re-checks the commit state each iteration and stops cleanly on a mid-loop virtual stop (remaining tokens stay virtually committed, matching the entry guard). This also hardens the classic path, where a mid-loop virtual stop would previously assert. With this fix the benchmark (5000 requests, ISL/OSL 1024/1024, 15 shared system prompts, concurrency 256) runs to completion in arena mode. The v2 regression suite result remains identical to main; manager suite 66/66 on H100, including a regression test where two sequences commit identical tokens and the loser must keep committing with no host-copy duplication and the canonical entry still reusable. [None][fix] Fix contiguous KV arena issues found in end-to-end validation (P0) Running Llama-3.2-1B through the LLM API with use_contiguous_kv_arena=True surfaced four gaps; with these fixes, arena outputs match classic paged v2 token-for-token (4 prompts x 200 greedy tokens, identical batch shapes, both write-through policies), and the reuse round matches exactly under ON_FREE. - _KVCache.resume() drains deferred range reclaim before evaluating the max_util_for_resume gate. Arena frees are event-gated and deferred, so page utilization can stay pinned above the gate after other sequences finished; a caller retrying resume() in a loop then livelocks, because nothing is schedulable and no other path drains either. This is exactly what the KV-capacity estimation phase hit: its exact-fit quota let two max-length dummy requests pin utilization at ~96%, the third dummy's resume was refused forever, and configure_kv_cache_capacity never got its responses. Classic mode recovers because slot frees are synchronous; the asymmetry was the bug. Regression test fills the whole page budget, closes without draining, and requires the next resume to succeed. - ArenaPoolGroup.destroy() synchronizes pending reclaim gates and drains before its leak assertion: the estimation flow tears the temporary executor down immediately after its dummy requests finish, with their ranges legitimately still in the deferred-reclaim queue. Teardown is allowed to block. - The post-warmup NaN/Inf KV scan is skipped in arena mode: it views the whole pool as a dense tensor, but an arena pool base is a sparse VA reservation that is not CUDA-registered while unmapped (and warmup pages are already unmapped by the time it runs). - get_buffers raises a descriptive NotImplementedError in arena mode; its consumers (FlashInfer-style backends, MLA latent append) are later-phase work and must fail loudly instead of dereferencing sparse VA. v2 regression suite result remains identical to main (only pre-existing TestSSMSupport failures); manager suite 64/64 on H100. [None][feat] Wire contiguous primary KV cache into the PyTorch executor (P0) Expose the contiguous-arena prototype through the executor stack (DESIGN §5, pyexecutor row): - KvCacheConfig gains use_contiguous_kv_arena (prototype status, following the use_kv_cache_manager_v2 precedent); setting it auto-enables the v2 manager. Tuning knobs stay on env vars while the feature is a prototype: TRTLLM_KV_ARENA_PHYS_PAGE_SIZE_MB, TRTLLM_KV_ARENA_MAP_AHEAD_PAGES, TRTLLM_KV_ARENA_WRITE_THROUGH (on_free | on_commit). - KVCacheManagerV2 adapter: builds ContiguousArenaConfig (requires a host cache tier -- the stale tier of §4.3; the cuMemHostRegister GPU-only fallback re-raises rather than silently dropping it), sizes each request's VA reservation to max_blocks_per_seq * tokens_per_block (matching offset-table sizing, covering extra KV and reserved draft tokens), drains deferred range reclaim every prepare_resources call, and reports physical capacity from the page budget in get_num_free_blocks (page-index bounds describe VA in arena mode). - clamp_max_seq_len_for_mem gains an arena branch that plans over physical pages against the shared budget (§4.6) instead of per-pool-group slot counts, which describe VA in arena mode. The v2 scheduler needs no changes: its allocate-or-suspend loop composes with arena backpressure (page exhaustion -> resize returns False -> preemption), and suspend/resume_request land on the §4.5 paths. Default behavior is untouched: the v2 regression suite result is identical before/after (only pre-existing TestSSMSupport failures, reproduced on main); api_stability passes (61/61) and all existing executor v2 tests pass (201/201). New tests cover the config field and validator, env knobs, real adapter construction in arena mode, physical capacity accounting, and the request lifecycle through adapter-facing APIs, plus page-based clamp_max_seq_len_for_mem math in the manager suite under both write-through policies. [None][feat] Add write-through-on-commit for contiguous primary KV cache (P0) Implement WriteThroughPolicy.ON_COMMIT (DESIGN §4.3): committed blocks are immutable, so their D2H write-out can happen the moment they commit, making the common free path copy-free. - _commit_block (arena mode, ON_COMMIT, host tier present) copies the just-committed block's pages to fresh host slots. The sources are still locked, so their ready events do not cover in-flight writes: the copy is ordered after an event recorded on the sequence's stream (write_through_pages gains prior_event, plumbed into _batched_migrate as extra_prior_event). A full host tier skips the copy silently -- the copy-on-free fallback covers such blocks. - The per-sequence stash of write-through slots is consumed at close() and suspend(): adopt_stale_copies moves each written-through page to the host tier without another copy. The vacated arena slot's reclaim gate merges the page's own last consumers with the write-through copy's finish event (the copy reads the slot); the page's ready event becomes the copy's finish event, so future consumers wait for its validity. Blocks without a write-through copy keep the existing copy-on-free path; unconsumed stash slots are returned to the host pool. - Rate limiting: the natural rate is one record per tokens_per_block steps per sequence (the §4.3 cost-model baseline); aggregate PCIe throttling is deferred to the executor integration (risk #4). Default behavior is untouched: ON_FREE remains the default and the v2 regression suite result is identical before/after (only pre-existing TestSSMSupport failures, reproduced on main). The ON_COMMIT test class subclasses the end-to-end suite so every arena scenario re-runs under the new policy, plus dedicated tests verify host copies exist at commit time while the blocks are still live on GPU, that close() adopts exactly the pre-copied host slots (no duplication, data verified via reuse), and that suspend() adopts committed blocks while moving the uncommitted tail (60/60 on H100). [None][feat] Add suspend/resume to per-sequence arenas for contiguous primary KV cache (P0) Wire preemption into arena mode (DESIGN §4.5): - suspend() evacuates the sequence's arena ranges after the classic lock->holder conversion: canonical committed pages and uncommitted pages migrate to the host tier (the block chain's holders follow their pages, so the logical chain survives suspension unchanged); sequence-private reuse copies swap back to holding their canonical radix entry and are dropped (or migrate like canonical pages if the entry is gone). The whole VA reservation is then released, gated on the sequence's finish event. - resume() after suspend takes the same path as a fresh resume: reserve fresh ranges (the base block index may change; offset tables are rebuilt from the locks) and onboard the resident state. _arena_onboard_matched now serves both cases: committed sources are copied into private pages (canonical entries untouched), uncommitted tail pages are moved (the page object relocates into its arena slot and the host slot is freed). Two correctness fixes that surfaced while testing: - close() had been passing a NULL event as the range-reclaim gate because _record_event clears _finish_event on exit; masked in tests by full-device synchronization, but a real use-after-unmap risk under overlap. close() and suspend() now capture the finish event inside the recording block. - PrivateCommittedPage is now never evictable (is_evictable returns False): a lingering local reference could delay a private copy's holder death past any exclusion pass, leaking the page into the GPU eviction queue and making its range unreclaimable. Keeping private copies out of eviction queues entirely removes the timing dependence; their slots free whenever the object dies. The paired orphan-exclude in _PageHolder.__del__ is guarded on scheduled_for_eviction accordingly (behavior-neutral for pages that were scheduled). Default behavior is untouched: the v2 regression suite result is identical before/after (only pre-existing TestSSMSupport failures, reproduced on main). New tests cover the suspend/resume round-trip (committed + uncommitted blocks written out, VA released, budget drained to zero, byte-identical restore at consecutive indices, growth continuing the run) and the suspend x reuse interaction (private copies dropped without host duplication, re-onboarded on resume with data verified) -- 51/51 on H100. [None][feat] Add reuse onboarding into per-sequence arenas for contiguous primary KV cache (P0) Wire the stale->active reuse path (DESIGN §4.4) into arena mode, re-enabling KV cache reuse for contiguous sequences: - PrivateCommittedPage: a sequence-private copy of a committed block. It shares the canonical page's content and tree-block association but is never registered in block.storage, and its destruction does not unset the canonical radix entry (Block.unset_page is not identity-guarded). - On a reuse hit, first resume() copies the matched committed prefix into the sequence's arena ranges: one explicit-destination migrate per (pool group, source level) with update_src=False, so the canonical entries stay valid wherever they live (host, disk, or another sequence's arena via D2D -- committed blocks are immutable, so concurrent reads are safe). Consecutive destination ordinals let the migrate path coalesce the copies. Dropping the match holders returns the canonical pages to eviction control at their tier. - close() drops private copies instead of offloading them -- the canonical stale copy already covers reuse -- so a reuse hit never duplicates host blocks. The committed-page collector moved into a helper so its loop locals cannot outlive the eviction-exclusion pass (a lingering lock-chain reference delayed a holder's death past the pass, leaking the page into the GPU eviction queue and making the range unreclaimable). - Arena mode matches full committed blocks only for now (partial-block matching disabled); partial reuse follows with the deferred private-copy path adapted to arena slots. The match must fit max_capacity, validated in create_kv_cache (raising inside _KVCache.__init__ would leave a partially constructed object for __del__). - _batched_migrate allows same-level copies for explicit destinations. Default behavior is untouched: the v2 regression suite result is identical before/after (only pre-existing TestSSMSupport failures, reproduced on main). New tests drive the full reuse round-trip with data verification -- commit in one sequence, close to host, match in a second sequence, onboard into consecutive arena indices, verify bytes on GPU, grow past the prefix, and confirm no host-block duplication after close (50/50 on H100). [None][feat] Add per-sequence arena growth to _KVCache for contiguous primary KV cache (P0) Wire the sequence lifecycle into the arena storage layer (DESIGN §5, _core/_kv_cache.py row), flag-gated by KVCacheManagerConfig.contiguous_arena: - create_kv_cache gains max_capacity (required in arena mode): the request's capacity ceiling in tokens, sizing its contiguous VA reservation. - First resume() reserves one contiguous block-index range per pool group; VA exhaustion backs off like any resource shortage (returns False). - resize() growth demand-maps physical pages covering the new block frontier (draining deferred reclaim and retrying once on page exhaustion before backing off) and issues slots at range base + ordinal, so the sequence's kernel-visible page indices are consecutive while all downstream slot/page/ offset-table machinery is unchanged. resize() beyond max_capacity fails loudly. - close() migrates the sequence's committed GPU pages to the host tier (active->stale copy-on-free; dropped instead if the host tier cannot take them) and queues the whole range for reclaim gated on the sequence's finish event. P0 scope gates, each lifted by a follow-up: reuse onboarding (stale->active copy) is skipped at creation and seq rebasing is disabled in arena mode; suspend() raises; SSM, sliding-window layers, and SWA scratch reuse are rejected at config validation. Also: prepare_free_slots with all-zero GPU requirements is a no-op in arena mode (empty lock batches during fresh resume), and offload_arena_pages un-schedules pages before migration as _batched_migrate requires. Default behavior is untouched: the v2 regression suite result is identical before/after (only pre-existing TestSSMSupport failures, reproduced on main). New end-to-end tests drive create -> resume -> grow -> commit -> close on a real KVCacheManager, covering consecutive page indices, host write-out, budget-exhaustion backoff with event-gated reclaim recovery, max_capacity enforcement, and the scope gates (49/49 on H100). [None][feat] Add storage-manager arena mode for contiguous primary KV cache (P0) Wire the flag-gated contiguous-arena storage path (DESIGN §5, _storage_manager.py row) into StorageManager: - Construction: KVCacheManagerConfig.contiguous_arena now plumbs through KVCacheManager into StorageManager; the GPU level constructs GpuArenaCacheLevelStorage. Arena VA extents are sized at startup (overcommit over quota/record_stride or an explicit per-pool VA cap, clamped to the int31 kernel-offset ceiling) and per-pool-group page-index scales are derived from buffer attributes for the int31 startup check. - Retired paths guarded: scattered GPU slot allocation, GPU-level slot eviction (prepare_free_slots) and GPU pool-group rebalancing raise LogicError in arena mode; host/disk tiers are unchanged. - Per-sequence pass-throughs: reserve_gpu_sequence, take_gpu_sequence_slot, ensure_gpu_mapped, free_gpu_sequence, drain_gpu_reclaim, gpu_page_budget. - GPU utilization reports the physical page-budget fraction (feeds max_util_for_resume gating). - _batched_migrate gains an explicit-destination mode for stale->active onboarding and suspend/resume placement; copies into explicit destinations coalesce byte-adjacent (dst, src) runs into fewer, larger transfers. - Active->stale helpers: offload_arena_pages (copy-on-free) and write_through_pages (write-through on commit). Default behavior (contiguous_arena=None) is untouched; the v2 regression suite result is identical before/after (only pre-existing TestSSMSupport failures, reproduced on main). New tests cover construction and sizing, guards, the sequence lifecycle, utilization accounting, and real data-movement round-trips for offload, explicit-destination onboarding, and write-through (44/44 on H100). [None][feat] Add arena-backed GPU storage seam for contiguous primary KV cache (P0) Second P0 chunk for the contiguous-in-VA active KV cache prototype in KVCacheManagerV2 (flag-gated; default paged behavior untouched): - _sequence_arena.py: generalize BlockRangeAllocator and SequenceArena to multi-pool groups (one block-index space per pool group, one sparse VirtMem per pool; range alignment = lcm across pools so no two sequences share a physical page in any pool). Add PageBudget, the level-wide physical-page quota that replaces per-pool-group slot partitioning; ensure_mapped charges it atomically and raises OutOfPagesError with nothing mapped on exhaustion. - _storage/_core.py: add ArenaSlotPool (per-pool address view), SequenceRange (per-sequence handle gating reclaim on outstanding slots, released-slot ready events, and the last-consumer event), ArenaPoolGroup (per-sequence reserve/take_slot/ensure_mapped/free_sequence/drain_reclaim; slot ids are base+ordinal so downstream Slot/Page/offset-table machinery is unchanged; scattered allocate() retired), and GpuArenaCacheLevelStorage (shared phys pool + PageBudget; int31 index-width check wired into construction). - setup_mypyc.py: add _sequence_arena.py to the mypyc module list. - tests: cover PageBudget, multi-pool alignment, atomic budget enforcement, arena pool group slot issue/release/reclaim gating, cross-pool-group budget sharing, and the int31 startup check. [None][feat] Scaffold contiguous primary KV cache (P0) for KVCacheManagerV2 Foundation for the contiguous-in-VA active KV cache (vAttention-style: per- sequence VA-contiguous blocks backed by CUDA VMM demand paging). Flag-gated; default behavior (classic paged v2) is unchanged. - _cuda_virt_mem.py: add sparse mapping to VirtMem (map_range/unmap_range/ is_mapped/num_sparse_chunks) alongside the existing tail-only stack, with a per-chunk page table, one cuMemSetAccess per contiguous run, and atomic OOM rollback. Sparse and tail-stack modes are mutually exclusive per reservation. - _sequence_arena.py (new): check_index_width (int31 kernel-offset ceiling), BlockRangeAllocator (page-aligned first-fit over the arena block-index space), and SequenceArena (demand paging with map-ahead margin + event-gated deferred reclaim). Remaining P0 integration seams are documented inline. - _config.py: ContiguousArenaConfig + WriteThroughPolicy; opt-in via KVCacheManagerConfig.contiguous_arena (None => unchanged). - tests: pure-logic coverage for the allocator, index-width check, and config; CUDA-gated coverage for sparse VirtMem and SequenceArena. Design and rationale: contiguous_primary_kvcache/DESIGN.md. [None][infra] Waive 1 failed cases for main in pre-merge 45709 (#15843) [None][infra] Upgrade NIXL to v1.3.0 (#15694) [https://nvbugs/5955773][test] Unwaive test_no_kv_cache_reuse for DeepSeekV3Lite (#15826) [None][perf] Eliminate torch.where host sync in multimodal fuse_input_embeds path (#14384) [https://nvbugs/6342840][fix] Add spec_metadata=None kwarg to SpecWorkerBase._apply_force_accepted_tokens… (#15797) [TRTLLM-13168][feat] test best ucx env for cache transceiver (#14933) [None][infra] Waive 10 failed cases for main in post-merge 2814 (#15825) [None][feat] add CI-enforced stability gate for trtllm-serve CLI options (#15643) [None][feat] Support DSA cross-layer indexer top-k sharing (e.g. GLM-5.2) (#15574) [TRTLLM-8997][feat] Add encoder_max_batch_size & encoder_max_num_tokens and deterministic dummy sizing (#13503) [None][fix] compressed-tensors fp8: resolve config group by name and flatten per-channel weight_scale (#15415) [None][feat] Support update weight for nemotron-h (#15573) [None][test] Waive 1 failed cases for main in post-merge (#15811) [None][infra] Check in most recent lock file from nightly pipeline [TRTLLM-13409][feat] hard-exit on HangDetector fire + cross-rank propagation (#15612) [https://nvbugs/6196614][infra] Upgrade aiperf to 0.8.0 (#15769) [TRTLLM-13630][test] Reuse MPIPoolExecutor across MoE multi-GPU test cases (#15744) [https://nvbugs/6162506][test] AutoDeploy: fix NVFP4 MoE fp4 weight scale in test + unwaives (#15748) [None][feat] Handle empty-block ranks with Helix CP (#15383) [None][feat] Add prefix-aware scheduling config flag to support opt-out (#15526) [https://nvbugs/6331421][fix] Fix TRTLLM-GEN backend multiCtasKv counter clear (#15761) [None][chore] Allow fp8 per-tensor base weights for MoE LoRA (#15528) [TRTLLM-12200][feat] WideEP FT: add active_rank_mask to NVLink AlltoAll kernels (1a.2) (#13404) [TRTLLM-12622][feat] Reland: Add native post-processing hook to trtllm-serve (#15631) [https://nvbugs/6336747][ci] Unwaive nemotron nvfp4 E2E test (#15561) [https://nvbugs/6250439][fix] Correct speculative XQA attention sinks (#15739) [None][fix] Honor Qwen Image quant ignore list (#15599) [None][feat] VisualGen: enable CUDA graph capture with torch.compile (#15603) [None][perf] set ncclConfig graphUsageMode=1 on communicator init (#13511) [None][feat] Sharding IR default (auto-detect) + Eagle/MTP draft sharding + qwen3_next (#14786) [https://nvbugs/6266306][fix] Fix testing harness (#15558) [None][chore] Autodeploy disable the pipeline cache by default (#15530) [None][feat] Cosmos3 reasoner only support (#15117) Co-authored-by: Chang Liu <9713593+chang-l@users.noreply.github.com> [https://nvbugs/6150288][fix] Use persistent per-stream workspace in cublas_mm for CUDA-graph safety (#15534) [TRTLLM-13580][test] Add model-derived PyTorch attention backend test suite (#15536) [None][test] waive tests (#15778) [None][infra] Waive 5 failed cases for main in pre-merge 45198 (#15779) [None][test] Waive 8 failed cases for main in post-merge (#15774) [None][feat] kv cache manager v1 use fabric memory to enable mnnvl transfer for kv cache transfer in disaggregated serving (#12490) [None][test] remove two model tests from test list (#15701) [None][test] Waive hang issues (#15764) [None][test] Waive 8 failed cases for main in QA CI (#15733) Co-authored-by: xinhe-nv <200704525+xinhe-nv@users.noreply.github.com> [None][chore] Bump version to 1.3.0rc21 (#15742) Co-authored-by: TensorRT LLM Release <90828364+tensorrt-cicd@users.noreply.github.com> [None][feat] Add 'kimi_k2.5_fp4' to TRUST_REMOTE_CODE_MODELS in performance tests (#15751) [None][docs] Add VisualGen engineering criteria (#15225) [None][infra] Check in most recent lock file from nightly pipeline [TRTLLM-13543][feat] WideEP FT: add EPLB mask-only reconfigure (1b.1) (#15525) [https://nvbugs/6316980][fix] Added a runtime guard in FlashInferTrtllmGenAttention.is_supported using the… (#15496) Co-authored-by: Dongfeng Yu <dongfengy@nvidia.com> [https://nvbugs/6273845][fix] Fix FMHA kernels are not found for GPTOSS + SM120 (#15230) Co-authored-by: Dongfeng Yu <dongfengy@nvidia.com> [https://nvbugs/6287834][chore] Unwaive GPT-OSS disagg perf sanity case (#15600) [None][test] Enable single-GPU LPIPS CI protection for VisualGen models (#15564) [None][test] AutoDeploy: Fix standalone linear simple on B200 (#15675) [TRTLLM-13613][test] Trim duplicated and dead multimodal accuracy tests from pre-merge CI (#15615) [None][feat] Parallelize host KV cache pool prefault and add THP control (#15431) Co-authored-by: Thor Johnsen <41591019+thorjohnsen@users.noreply.github.com> [None][ci] Disable home mount for sbatch L0 tests (#15713) [None][test] Waive 1 failed cases for main in QA CI (#15719) Co-authored-by: xinhe-nv <200704525+xinhe-nv@users.noreply.github.com> [None][test] Waive 4 failed cases for main in post-merge (#15721) Co-authored-by: xinhe-nv <200704525+xinhe-nv@users.noreply.github.com> [TRTLLM-13319][fix] fix output distribution correctness for Eagle3 dynamic-tree rejection sampling (#15098) [None][feat] Bind NUMA-aware CPU affinity in visual_gen workers (#13590) [None][feat] add CI-enforced stability gate for trtllm-serve HTTP schema (#15644) [None][feat] Optimize trtllmgen moe routing (#15656) [None][infra] Waive 5 failed cases for main in post-merge 2811 (#15705) [None][test] Waive 4 failed cases for main in QA CI (#15686) Co-authored-by: xinhe-nv <200704525+xinhe-nv@users.noreply.github.com> [None][infra] Check in most recent lock file from nightly pipeline [None][test] Waive 2 failed cases for main in QA CI (#15667) Co-authored-by: xinhe-nv <200704525+xinhe-nv@users.noreply.github.com> [None][test] record per-case hostname in perf result CSV (#15663) Co-authored-by: Ruodi Lu <ruodil@users.noreply.github.com> [None][infra] Waive 9 failed cases for main in post-merge 2810 (#15700) [TRTLLM-12751][feat] visual-gen /metrics iteration stats producer (#14246) [None][test] Waive 1 failed cases for main in QA CI (#15688) [None][feat] Cache LTX2 merged LoRA weight (#14984) [None][test] Waive 1 failed cases for main in QA CI (#15678) Co-authored-by: xinhe-nv <200704525+xinhe-nv@users.noreply.github.com> [None][chore] Decrease the qwen3.5 ci test time (#15646) [None][chore] split unittest/_torch/visual_gen (#15670) [None][perf] DSv4 follow-up: autotuner updates (#15626) Co-authored-by: Mingyang <mhao1999@outlook.com> Co-authored-by: Mingyang Hao <mingyangHao@users.noreply.github.com> Co-authored-by: Yukun He <23156053+hyukn@users.noreply.github.com> [None][fix] GLM-5.1 NVFP4 fallback to AR-Norm fusion for unquantized dense layers (#15659) [TRTLLM-11715][infra] Upgrade dependencies for dlfw 26.04 stack (#12643) Co-authored-by: chenfeiz0326 <chenfeiz@nvidia.com> [TRTLLM-13265][feat] Fuse LTX-2 Gate + Residual + Norm + AdaLN modulation(ShiftScale) + Quant kernels (#15102) [None][feat] DSv4: model, tokenizer, and integration coverage (#15414) Co-authored-by: Fanrong Li <lfr-0531@users.noreply.github.com> For merge to unblock some hard dependency. Fanrong will monitor the new ToT in case there are anything unexpected. [TRTLLM-13629][test] Optimize MoE CI test-db (#15624) [None][fix] Fall back to NVLink P2P when NVLS fabric is unprovisioned (#15302) [None][feat] Support inflight weight update (#14815) [TRTLLM-12721][fix] Bound V2 context transfer polling (#15356) [None][infra] Check in most recent lock file from nightly pipeline [None][feat] Converge VisualGen LPIPS and VBench test generation (#14654) [None][feat] DSv4: sparse MLA attention backend (#15409) Co-authored-by: Fanrong Li <lfr-0531@users.noreply.github.com> [https://nvbugs/6301807][fix] Fix FP8 rowwise linear reference precision (#15421) [None][infra] Check in most recent lock file from nightly pipeline [https://nvbugs/6062416][fix] Cache NCCL window allocation failures by size (#15596) Co-authored-by: tensorrt-cicd <90828364+tensorrt-cicd@users.noreply.github.com> [https://nvbugs/6344612][test] relax GPT-OSS GPQA references due to high variance in random sampling (#15567) [None][feat] add MXFP8 weight format + CUTLASS W8A8 Linear and MoE (#14962) [https://nvbugs/6269778][fix] Fix overallocation of draft KV cache (#15017) [https://nvbugs/6248837][chore] waive memory polluters (#15665) [TRTLLM-13370][perf] LTX2 + WAN: Fuse MLP up-GEMM + bias + GELU(tanh) + NVFP4-quant into CuteDSL epilogue (#15299) [https://nvbugs/6293536][fix] Stage KV block offsets through a fresh host buffer (#15546) [None][fix] User/tjohnsen/evict empty blocks first (#11685) Co-authored-by: coderabbitai[bot] <136622811+coderabbitai[bot]@users.noreply.github.com> [https://nvbugs/6248783][test] Unwaive Qwen3 skip softmax test (#15652) [TRTLLM-13712][feat] Add Qwen-Image-Bench evaluator (#14837) [https://nvbugs/6239637][fix] Unwaive Qwen3.5 cases on A100 platform (#15481) [None][chore] Update .gitattributes (#15606) [None][doc] Add deploy guide for Minimax M3 (#15587) [None][infra] take test durations into account to determine cbts splits num (#15614) [None][test] Add Qwen3.5-397B-A17B-NVFP4 B200 aggregated perf-sanity tests (#15650) [TRTLLM-12982][chore] improve multi-item scoring request validation (#15627) [None][test] add GLM nvfp4 stress test (#15437) Co-authored-by: GitLab CI Bot <gitlab-ci@nvidia.com> [TRTLLM-13575][feat] Add EPLB support for Qwen3.5 (#15543) [TRTLLM-12762][test] Add Test coverage for MiniMax Model with multi-node, M2.5 checkpoints eval (#15361) [#15613][fix] Gemma4 multimodal: fix vision TP and xgrammar startup crashes (#15566) [#15565][fix] AutoDeploy: Fix Super MTP IMA introduced by checkpointing replay (#15622) [None][test] Add modularized perf tests (attention + MoE discrete/continuous) (#15541) Co-authored-by: Ruodi Lu <ruodil@users.noreply.github.com> [None][infra] Check in most recent lock file from nightly pipeline [https://nvbugs/6368480][fix] Cache the SM count once in FmhaDispatcher's constructor and reuse the cached… (#15611) [TRTLLM-13247][feat] Wave 2: stage Linear and Attention transforms (#15288) [TRTLLM-13444][test] Add Qwen-Image text-to-image unit tests (#15580) …

Reduces the pre-merge disagg PerfSanity set from 8 tests to 4 by moving these back to post-merge: * NVIDIA#2 gpt-oss-120b-fp4 TEP4 gen_only 1k1k con64 (GB200) * NVIDIA#4 deepseek-r1-fp4 TEP8 gen_only 1k1k con1 (GB200) * NVIDIA#5 kimi-k25-thinking-fp4 DEP8 e2e 1k1k con4096 (GB200) * NVIDIA#6 glm-5-fp4 TEP4 gen_only 1k1k con1 (GB300) Remaining 4 pre-merge FUNCTIONAL-ONLY stages: * gpt-oss-120b-fp4 DEP2 e2e 1k1k con2048 (GB200, 2n/8g) * kimi-k25-thinking-fp4 TEP4 gen_only 1k1k con4 (GB200, 2n/8g) * glm-5-fp4 DEP8 e2e 1k1k con4096 (GB300, 3n/12g) * deepseek-r1-fp4 DEP8 gen_only 1k1k con2048 (B200, 2n/16g) Corresponding Jenkins adjustments: 2 GB200 pre-merge stages deleted, 2 GB200 Post-Merge testCounts bumped, 1 GB300 pre-merge stage renamed back to Post-Merge. Signed-off-by: Chenfei Zhang <chenfeiz@nvidia.com>

Follow-up: user requires each pre-merge pick to already be actively running in L0 CI post-merge on origin/main — i.e. no QA-lane borrowing, no uncommenting of previously-disabled lines, no new config yamls, no first-time-in-L0 e2e/gen_only combinations. Prior revision's two picks that failed that bar: * gpt-oss 8k1k con512 e2e dep2 — active in QA lane only; the e2e mode of this 8k1k dep2 config was not in L0 anywhere. * kimi 8k1k con4 tep8 mtp3 gen_only GB200 — was commented out in the GB200 12-GPU 3-node post_merge block since April. Replaced with: * gpt-oss 8k1k con1024 tp4 e2e MTP0 — was active in L0 post-merge on GB200 ctx1_node1_gpu1_gen1_node1_gpu4 (8 GPUs, 2 nodes). Gen-worker parallelism shifts dep→tp for this slot. * kimi 8k1k con4 tep8 mtp3 gen_only — was active in L0 post-merge on GB300 ctx1_node1_gpu4_gen1_node2_gpu8 (12 GPUs, 3 nodes). GPU family shifts GB200→GB300 for this slot. File changes: - l0_gb200_multi_nodes_perf_sanity_ctx1_node1_gpu1_gen1_node1_gpu2.yml: reverted to origin/main state (no pre_merge block). - l0_gb200_multi_nodes_perf_sanity_ctx1_node1_gpu1_gen1_node1_gpu4.yml: new pre_merge block with gpt-oss 8k1k con1024 e2e; test line removed from post_merge to avoid double-run. - l0_gb200_multi_nodes_perf_sanity_ctx1_node1_gpu4_gen1_node2_gpu8.yml: reverted to origin/main state (kimi 8k1k stays commented out). - l0_gb300_multi_nodes_perf_sanity_ctx1_node1_gpu4_gen1_node2_gpu8.yml: new pre_merge block with kimi 8k1k con4 tep8 mtp3 gen_only; same test removed from post_merge. Jenkins: - Dropped stages: FUNCTIONAL-ONLY-CTX1-NODE1-GPU1-GEN1-NODE1-GPU2 (GB200) and FUNCTIONAL-ONLY-CTX1-NODE1-GPU4-GEN1-NODE2-GPU8 (GB200 kimi variant). - New stages: FUNCTIONAL-ONLY-CTX1-NODE1-GPU1-GEN1-NODE1-GPU4 (GB200 gpt-oss) and FUNCTIONAL-ONLY-CTX1-NODE1-GPU4-GEN1-NODE2-GPU8 (GB300 kimi variant). - Post-merge decrement: GB200-8_GPUs-2_Nodes CTX1-NODE1-GPU1-GEN1- NODE1-GPU4 7→6; GB300-12_GPUs-3_Nodes CTX1-NODE1-GPU4-GEN1-NODE2- GPU8 5→4. Others revert to origin/main testCounts. Shape matrix (all 4 pre-merge slots now at 8k1k): NVIDIA#1 DSR1 B200 8k1k con1536 dep e2e=NO gen_only MTP✓ NVIDIA#2 gpt-oss GB200 8k1k con1024 tp e2e=YES MTP✗ NVIDIA#3 glm-5 GB300 8k1k con1024 dep e2e=YES MTP✓ NVIDIA#4 kimi GB300 8k1k con4 tep e2e=NO gen_only MTP✓ Signed-off-by: Chenfei Zhang <chenfeiz@nvidia.com>

kaiyux added 3 commits October 10, 2023 23:22

Update

027cd51

Update

6c3bfd2

Update

3a87f27

kaiyux marked this pull request as draft October 11, 2023 07:17

kaiyux added 2 commits October 11, 2023 00:42

Update

39d574a

Update .a libs

cf3028d

kaiyux marked this pull request as ready for review October 11, 2023 07:58

kaiyux requested review from jdemouth and juney-nvidia October 11, 2023 07:59

juney-nvidia approved these changes Oct 11, 2023

View reviewed changes

kaiyux merged commit 4941ad2 into main Oct 11, 2023

juney-nvidia deleted the kaiyu/update branch October 19, 2023 00:43

zxs789 mentioned this pull request Jun 4, 2024

H20 Using random weights to infer llama2-13B results in a divide-by-zero error. #1717

Closed

4 tasks

wu1du2 pushed a commit to wu1du2/TensorRT-LLM that referenced this pull request May 11, 2025

Merge pull request NVIDIA#4 from NVIDIA/kaiyu/update

2c80726

Update TensorRT-LLM

Superjomn added a commit that referenced this pull request Jun 1, 2025

fix: llmapi-launch add add trtllm-bench test with engine building (#4… (

55170ec

#4550) Signed-off-by: Superjomn <328693+Superjomn@users.noreply.github.com>

v-shobhit pushed a commit to v-shobhit/TensorRT-LLM that referenced this pull request Jul 7, 2025

Merge pull request NVIDIA#4 from Tom-Zheng/update_sm100f_cubins

0d426f0

Update sm100f cubins

yuxianq added a commit to yuxianq/TensorRT-LLM that referenced this pull request Jul 18, 2025

Fix OOM issue for fp8 resmooth. (NVIDIA#4)

0286391

Signed-off-by: Yuxian Qiu <142763828+yuxianq@users.noreply.github.com>

litaotju pushed a commit to litaotju/TensorRT-LLM that referenced this pull request Jul 24, 2025

Fix OOM issue for fp8 resmooth. (NVIDIA#4)

3ee97f7

Signed-off-by: Yuxian Qiu <142763828+yuxianq@users.noreply.github.com> Signed-off-by: Fanrong Li <23290157+lfr-0531@users.noreply.github.com>

yuxianq added a commit to yuxianq/TensorRT-LLM that referenced this pull request Jul 28, 2025

Fix OOM issue for fp8 resmooth. (NVIDIA#4)

d83cc25

Signed-off-by: Yuxian Qiu <142763828+yuxianq@users.noreply.github.com> Signed-off-by: Fanrong Li <23290157+lfr-0531@users.noreply.github.com>

dmvevents mentioned this pull request Dec 15, 2025

[Bug] ABI incompatibility causes segfault when mixing rc4/rc5 binaries #10016

Open

justremotephone mentioned this pull request Feb 5, 2026

[Bug]: Random "CUDA error: an illegal instruction was encountered" for gpt-oss 120B at Nvidia Spark. #11313

Closed

5 tasks

thorjohnsen mentioned this pull request Mar 16, 2026

[TRTLLM-11513][fix] eliminate dangling reference and TOCTOU races in KVCacheManager #12251

Closed

3 tasks

chienchunhung mentioned this pull request May 5, 2026

[https://nvbugs/6104831][fix] Disaggregated KV transfer: lifecycle, cancellation, and quiescence hardening #13713

Closed

1 task

xxi-nv mentioned this pull request May 8, 2026

[None][refactor] MoEScheduler split + MegaMoE EPLB / multi-chunk / CI integration #13908

Merged

2 tasks

longcheng-nv mentioned this pull request May 21, 2026

[None][perf] DSv4 indexer Top-K: align kFTarget=kK, uniform kSeqSmall=4096, diagnostic env knobs #14413

Merged

zhenhuaw-me mentioned this pull request May 29, 2026

[TRTLLM-12491][feat] Align VisualGen serve request schema with VisualGenParams #14733

Merged

limin2021 mentioned this pull request Jun 1, 2026

[TRTLLM-35882][feat] Add cute dsl gvr top-k decode kernel #14602

Merged

1 task

thorjohnsen mentioned this pull request Jul 4, 2026

[None][feat] Contiguous primary KV cache (P0): per-sequence VMM arenas for KVCacheManagerV2 #15931

Draft

Wanli-Jiang mentioned this pull request Jul 15, 2026

[TRTLLM-14026][feat] BREAKING: Remove C++ modules for legacy TRT backend #16369

Merged

1 task

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

Update TensorRT-LLM#4