[None][chore] split unittest/_torch/visual_gen#15670
Merged
Merged
Conversation
Signed-off-by: Tyler Burt <195370667+tburt-nv@users.noreply.github.com>
Collaborator
Author
|
/bot run --stage-list "DGX_B200-8_GPUs-PyTorch-1, DGX_B200-8_GPUs-PyTorch-2, DGX_B200-8_GPUs-PyTorch-3, DGX_B200-8_GPUs-PyTorch-4" |
Contributor
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (2)
📝 WalkthroughWalkthroughUpdates B200-related test configuration and selection: the DGX_B200-8 PyTorch Slurm split count is raised to 4 with a new split group, and the VisualGen multi-GPU integration list is expanded into explicit module paths. ChangesB200 test split and list updates
Estimated code review effort🎯 2 (Simple) | ⏱️ ~10 minutes Possibly related PRs
Suggested reviewers
🚥 Pre-merge checks | ✅ 5✅ Passed checks (5 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
dpitman-nvda
approved these changes
Jun 26, 2026
dpitman-nvda
left a comment
dpitman-nvda
left a comment
Collaborator
There was a problem hiding this comment.
LGTM, just one comment for my own understanding
Collaborator
|
PR_Github #56118 [ run ] triggered by Bot. Commit: |
Collaborator
|
PR_Github #56118 [ run ] completed with state |
chzblych
approved these changes
Jun 29, 2026
evezhier
pushed a commit
to evezhier/TensorRT-LLM
that referenced
this pull request
Jul 2, 2026
Signed-off-by: Tyler Burt <195370667+tburt-nv@users.noreply.github.com>
thorjohnsen
added a commit
to thorjohnsen/TensorRT-LLM
that referenced
this pull request
Jul 8, 2026
… 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) …
BrianLi23
pushed a commit
to BrianLi23/TensorRT-LLM
that referenced
this pull request
Jul 9, 2026
Signed-off-by: Tyler Burt <195370667+tburt-nv@users.noreply.github.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Description
The
unittest/_torch/visual_gensuite takes a long time. This PR splits it into smaller pieces so that it can be distributed between shards better. Also adds another shard to be safe.Test Coverage
PR Checklist
Please review the following before submitting your PR:
PR description clearly explains what and why. If using CodeRabbit's summary, please make sure it makes sense.
PR Follows TRT-LLM CODING GUIDELINES to the best of your knowledge.
Test cases are provided for new code paths (see test instructions)
If PR introduces API changes, an appropriate PR label is added - either
api-compatibleorapi-breaking. Forapi-breaking, includeBREAKINGin the PR title.Any new dependencies have been scanned for license and vulnerabilities
CODEOWNERS updated if ownership changes
Documentation updated as needed
Update tava architecture diagram if there is a significant design change in PR.
The reviewers assigned automatically/manually are appropriate for the PR.
Please check this after reviewing the above items as appropriate for this PR.
GitHub Bot Help
To see a list of available CI bot commands, please comment
/bot help.Summary by CodeRabbit