Skip to content

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

Draft
thorjohnsen wants to merge 38 commits into
NVIDIA:mainfrom
thorjohnsen:feat/contiguous-primary-kvcache-p0
Draft

[None][feat] Contiguous primary KV cache (P0): per-sequence VMM arenas for KVCacheManagerV2#15931
thorjohnsen wants to merge 38 commits into
NVIDIA:mainfrom
thorjohnsen:feat/contiguous-primary-kvcache-p0

Conversation

@thorjohnsen

@thorjohnsen thorjohnsen commented Jul 3, 2026

Copy link
Copy Markdown
Collaborator

Summary

Prototype implementation (P0) of the contiguous primary KV cache for KVCacheManagerV2: each sequence's active KV blocks are contiguous in virtual memory, backed by physical pages mapped on demand via CUDA VMM, while reusable (stale) state lives in the host tier.

Flag-gated end to end: KvCacheConfig(use_contiguous_kv_arena=True) (prototype status). Default behavior is untouched — with the flag off, no code path changes.

Design

  • One large VA reservation ("arena") per (pool group, pool); each sequence gets a contiguous, page-aligned block-index range sized for its per-request maximum, so slot_id = base + ordinal and the sequence's kernel-visible page indices are consecutive. Arena base = pool base pointer → zero kernel changes, CUDA-graph safe.
  • Physical 2 MiB pages (configurable super-pages) map on demand against a level-wide page budget; freed ranges are parked still-mapped and adopted whole by the next admitted sequence (default ON), so steady state performs zero cuMemMap/cuMemUnmap; unmaps happen only under real pressure (spill) or teardown, behind a device quiesce.
  • Active→stale: committed blocks copy out to the host tier on free, or eagerly via an opt-in write-through-on-commit policy that makes the free path copy-free (adoption of the pre-copied host slot).
  • Stale→active: a reuse hit copies the matched prefix into the new sequence's range as sequence-private pages; canonical radix entries are untouched. Suspend/resume evacuates and re-onboards through the same paths.
  • Prefix aliasing (opt-in, prefix_aliasing / TRTLLM_KV_ARENA_PREFIX_ALIASING=1): a closing canonical owner's fully-committed prefix pages are pinned (refcounted) in a canonical-span registry, and later same-prefix admissions cuMemMap the SAME physical pages into their own ranges — zero-copy, zero-budget reuse, with prefix-affine adoption (parked ranges carry an (span, extent) signature and are only handed to admissions matching it, since adoption overwrites page content). This recovers paged's refcount-sharing economics for the arena layout.
  • Capacity planning (clamp_max_seq_len_for_mem, scheduler backpressure) moves from slot counts to physical pages; the v2 scheduler needed no changes.

Full design doc, VMM microbenchmarks (H100 + GH200), and per-commit work log are tracked in the development directory (contiguous_primary_kvcache/DESIGN.md); can be attached or moved into docs/ on request.

P0 scope gates (fail loudly, lifted in later phases)

Sliding-window/VSWA layers, SSM, SWA scratch reuse, partial-block reuse matches, seq rebasing, beam search (unsupported in v2 anyway), FlashInfer-style dense pool views (get_buffers), disaggregated serving.

Correctness campaign (8B scale): three distinct bugs, all root-caused and fixed

Scaling validation from 1B to Llama-3.1-8B exposed an intermittent illegal-memory-access that turned out to be three independent bugs (two of which also affect or inform main):

# Bug Class Fix
a KV block-offset staging race: the offset-gather kernel could read pinned staging buffers already overwritten by the next iteration under the overlap scheduler → stale offsets → IMA (arena) / silent wrong reads (classic) CPU-side buffer reuse (nvbug-6293536 class); affects main 4d49f1d66e here; standalone main fix in PR #16088
b Scheduler capacity deadlock: arena frees are event- and fence-gated, and the fence was only assigned after scheduling — pages freed by the scheduler's own evictions were unreclaimable in-pass → suspension cascade → guard raise Arena-specific capacity model 4fefc34db5 — blocking fenced drain inside the pass + allocation retry
c cuMemUnmap racing in-flight kernels: unmapping a VA range that running kernels concurrently read intermittently produces a Warp MMU fault Driver-level; being escalated to the driver team with a repro matrix b271d635fd (freed-range adoption removes steady-state unmaps entirely); TRTLLM_KV_ARENA_SYNC_BEFORE_UNMAP (default ON) quiesces the rare pressure/teardown unmaps

Evidence for (c), with (a) and (b) fixed: on a deterministic pressure reproducer (free_gpu_memory_fraction=0.42 ≈ 2.4× KV oversubscription @ conc 256), no-quiesce crashes ~4/7 across mapping margins (growth maps exonerated), while sync-before-unmap runs 6/6 clean (p ≈ 0.6% against the ~57% base rate). Two further ordering fixes landed from the same campaign: arena write-outs (suspend/close) are now ordered after the in-flight forward step via a caller-supplied prior event (52a311ed99 — silent stale-tail restore, proven both directions), and on-demand mapping (margin=1) was restored as the default (53f7167405) after pre-mapping proved to be a ~3-point stopgap for what was mostly bug (a).

A fourth bug was root-caused and fixed after the table above (51e092e3ed): host-tier exhaustion during suspension. The host tier legitimately fills with already-suspended sequences' HELD pages (which LRU eviction correctly refuses to drop at the last level), and suspend() mutated sequence state before its fallible host allocation — so exhaustion either escaped as a process-fatal OutOfPagesError or, raced by the in-flight overlap step against the half-suspended state, poisoned the device as an IMA. Fix: an atomic pre-flight host reservation in suspend() (failure leaves the sequence fully intact) plus graceful degradation to v1-style drop-and-recompute through the adapter, scheduler, and executor (surviving canonical prefixes still shorten the re-prefill). Discriminator (Llama-3.1-8B, 0.42 quota, 4 GiB host ≈ 10× under suspension demand, conc 256): unfixed 4/4 dead (2 IMA + 2 clean FATAL — two faces of the same trigger); fixed 8/8 clean at 300/300 with 40+ drop-and-recompute cycles per run.

Endurance under production defaults on the pressure reproducer: 300-req ×3 + 2500-req runs all clean — zero IMA, zero deadlock.

Testing

  • tests/unittest/kv_cache_manager_v2_tests/test_contiguous_arena.py: 108 tests — allocator/budget/int31 pure logic, sparse VirtMem, arena storage seam, manager-level end-to-end (growth, reuse round-trip with data verification, suspend/resume byte-identical restore, page-exhaustion backoff, page-based capacity planning), lazy-retention matrix, batched-sweep charging, freed-range adoption semantics, blocking-drain deadlock terminal state, host-exhaustion suspension atomicity (HELD-pinned raise + DROPPABLE churn), prefix aliasing (refcount matrix, zero-charge accounting, registry roundtrip, and two manager-level end-to-ends where zero-copy is proven by corrupting every host copy — including the partial-span shape where the owner's chain outlives a shorter adopter's match, and the pressure window where the registry spills between an admission's lookup and its first resume — dead spans fall back to the host-copy path), and write-out ordering (with explicit copy-kernel warmups — the first batchedCopy call pays ~200 ms module load and masks races otherwise), run under both write-through policies via class inheritance. Newest additions: span-spill protection (floor survives full-pressure spill while parked ranges die, excess spills oldest-first, fraction=0 restores unprotected spilling), map-ahead margin degradation (frontier-only mapping at the budget edge, deterministic batched-sweep charge replay, delta-degrade excess release), and owner-live aliasing (owner-charged span lifecycle at the pool-group seam; a manager-level end-to-end where the consumer admits while the owner is mid-generation and physical sharing is proven by writing through the owner's VA and reading the change through the consumer's).
  • tests/unittest/_torch/executor/test_kv_cache_v2_contiguous_arena.py: 12 adapter tests (config field/validator, env knobs, real adapter construction, physical capacity accounting, request lifecycle, terminal-state allocation retry, drop-and-recompute fallback).
  • tests/unittest/_torch/executor/test_kv_block_offset_overlap_race_v2.py: regression test for the offset-staging race (a).
  • v2 regression suite: 219 passed / 12 skipped (211 inherited + 8 added), identical to main before/after every commit.
  • api_stability 61/61; existing executor v2 tests 201/201.

Throughput benchmarks — Llama-3.1-8B (H100, solo runs on an idle node)

trtllm-bench throughput, 2500 requests, ISL/OSL 1024/1024, concurrency 256, v2 manager + block reuse + 64 GiB host tier, prefix-sharing dataset (15 shared system prompts, avg 512 tokens). Arena = production defaults (adoption ON, on-demand mapping, on_free, 16 MiB pages, linear kernels off).

Memory-equalized runs (both backends at ~52.5–52.9 GiB KV quota — a total_quota accounting bug that silently granted the arena ~8 GiB more than paged was found and fixed, c645628fcc):

config tok/s vs paged
paged v2 (baseline) 8,563
arena (adoption) 7,746 −9.5%
arena + prefix aliasing 7,816 / 7,844 −8.6%
arena + aliasing + span-spill protection (+ margin degrade, owner-live, dedup remap) 8,125–8,278 over 7 runs −4.5% ± 0.5 (paged 8,617–8,631, ±0.05%)

At this quota the shared-prefix workload sits exactly at the arena's capacity edge even WITH aliasing ((3,385 pages − 60 registry)/13 charged pages/seq = 256), so pressure spills thrashed the canonical-span registry: alias hit rate dropped to ~9% and all 15 spans got spilled — aliasing recovered only ~0.9 points. Span-spill protection (463f45cee9) fixes this: the registry is exempt from pressure spills up to protected_span_fraction of the page budget (default 5%; the excess stays spillable LRU-first), and the resume-utilization gate stops counting protected pages as reclaimable. Re-run at the same equalized quota: zero spans spilled, 2,278 alias hits with zero lookup misses, recovering ~4 points (−8.6% → −4.5%). An error-bar note on that row: an interleaved A/B bisect across the night's runs showed the ARENA config wanders ~±0.7% across time windows on an otherwise idle node while paged holds ±0.05%, so the honest pooled number for the full stack is −4.5% ± 0.5 — single run-pairs at this scale cannot resolve sub-point deltas. Within that resolution the three follow-on commits are individually throughput-neutral at this benchmark shape, each by an understood mechanism: best-effort map-ahead margin (812ef73a32) converts capacity-edge growth failures into frontier-only mappings (this workload rarely hits that path); owner-live aliasing (35260d6afc) registers spans at context-end commit rather than close — at saturated concurrency admissions happen exactly at closes, so close-time registration was already timely, and its win shape is staggered/bursty same-prompt arrivals; pressure-time dedup remap (530da2a633) frees duplicate private prefix copies (commit-conflict losers, pre-span copy-onboards) by remapping their VA onto the canonical physical pages, as a reclaim-ladder rung between the parked-range spill and the registry-excess spill — one device quiesce per batch, paid only when the alternative is a preemption; the benchmark's steady state never reaches that rung (every admission aliases, so nothing is duplicated).

Capacity sensitivity (the same runs before the quota fix, i.e. arena inadvertently at 60.8 GiB — retained as the headroom data point):

config tok/s vs paged (52.5 GiB)
arena (adoption) 8,051 / 8,058 −6.1%
arena + prefix aliasing 8,471 / 8,461 −1.3% (87% alias hit rate)
arena + aliasing + linear kernels 8,434 / 8,439 −1.6%

Read together: the aliasing mechanism works — with ~8 GiB of headroom the spans stay resident and it recovers the full duplication gap (−1.3%, beating the −1.6% reuse-off bound) — but the current spill policy forfeits it exactly when memory is tight: the registry is only ~1 GB, and spilling it costs ~5 points. Attribution (reuse-off discriminator, identical footprint both modes): the arena-vs-paged gap is dominated by prefix duplication (paged refcount-shares and fits the quota eviction-free; arena's private copies oversubscribe it), with −1.6% of sharing-independent overhead underneath. Follow-ups: span-aware spill protection, map-ahead-margin degradation, and owner-live aliasing have all since landed (463f45cee9, 812ef73a32, 35260d6afc — see the equalized table above). Owner-live registration pins the span at the owner's context-end commit (matching paged, which shares context blocks the moment they commit), so same-prefix admissions alias while the owner is still generating; the entry stays owner-charged (no registry charge, spill-exempt) until the owner's free_sequence flips it to the registry-charged state. Pressure-time dedup remap (530da2a633) has also landed: a stateless scan over live sequences finds committed prefixes held as private duplicates of a registered span and re-backs their VA with the canonical physical pages (same-VA swap — slot ids, offsets, and page holders are untouched, since page objects are slot-indexed), freeing the duplicates' budget charge from LIVE sequences with zero information loss before the reclaim ladder resorts to spilling the registry or preempting.

Aliasing status: prototype, default OFF (ContiguousArenaConfig.prefix_aliasing / TRTLLM_KV_ARENA_PREFIX_ALIASING=1). Correctness-validated at production-shaped load (multiple 2×2500 runs, zero errors) and across a pressure discriminator matrix (adoption-off, tight-quota-only, registry-spill fallback). One rare first-wave IMA on the extreme pressure reproducer (2.4× KV oversubscription + 10× undersized host tier) remains under investigation — a dedicated 84-round multi-GPU grind reproduced it on BOTH arms, with the no-alias arm crashing at 5× the aliasing arm's rate (8/36 vs 2/48; the no-alias crashers fire with an empty registry and zero drop/OutOfPages events), confirming it is not the aliasing machinery but a rare residual of the pre-hardening IMA class (the driver-interference family being escalated separately).

Gap trajectory as the campaign's fixes landed (same workload; node-shared runs up to adoption, solo after; memory-equalized from c645628fcc): pre-map default −12.8% → margin=18 −9.4% → margin=1 −7.4% → adoption −6.8% → equalized adoption −9.5% → + prefix aliasing −8.6% → + span-spill protection (and follow-ons) −4.5% ± 0.5 (−1.3% with 8 GiB of KV headroom).

Throughput benchmarks — Llama-3.2-1B (H100)

5000 requests, ISL/OSL 1024/1024, concurrency 256, v2 manager + block reuse + 8 GiB host tier in all runs. Two datasets: shared (every prompt prefixed by one of 15 shared system prompts, avg 512 tokens — maximizes prefix reuse) and unique (fully unique prompts — isolates intrinsic overhead). Note both are cache-oversubscribed: ~320k committed blocks vs a ~65k-block GPU quota, so classic paged also offloads D2H at steady state (it runs copy-free only for the first ~15% of requests); copy volume is comparable between modes.

dataset paged v2 arena on_free / 16 MiB arena on_commit / 16 MiB arena on_free / 2 MiB
unique 30,908 tok/s 30,785 (−0.4%) 28,147 (−8.9%)
shared 33,885 tok/s 31,987 (−5.6%) 30,423 (−10.2%) 28,689 (−15.3%)

Findings:

  • Arena reaches parity with paged (−0.4%, noise) on the unique-prompt workload with the default on_free policy and 16 MiB super-pages. Demand paging, VA management, deferred reclaim, and write-out volume are all in the noise at this scale.
  • The apparent "intrinsic" −8.9% was entirely the on_commit write-through's per-block dispatch granularity (~320k small migrations vs classic's large batched evictions) — the rate-limiting concern the design flags as risk Update TensorRT-LLM #4. on_free (the default) batches per request. A batched/async write-through dispatcher is the P1 fix before on_commit is recommendable.
  • All runs completed 5000/5000 requests — sustained stress of demand paging, deferred reclaim, and the page budget at high concurrency.
  • With policy and page size held at their best settings: arena is at parity (−0.4%) without prefix sharing and −4.4% to −5.6% behind with heavy sharing (−4.4% with 256 system prompts and a 64 GiB host tier); 2 MiB paging adds ~10 points of churn at this request rate (16 MiB super-pages recommended); on_commit write-through as currently implemented adds ~5-8 points of per-block dispatch (on_free is the default).

Attributing the residual gap (P2 + ablations)

The branch also implements P2 lazy GPU retention (§4.4 phase 2, opt-in): freed ranges stay mapped on a retained LRU, reuse hits copy D2D from the resident bytes (private copies re-register as fresh sources so hot prefixes survive LRU spills), and the page budget reclaims retained ranges only under pressure — plus the §4.2 batched map sweep (opt-in): growth maps charge the budget at resize time but execute in one batched pass per iteration. (Freed-range adoption, above, grew out of this parking machinery and is now default-ON in all modes.)

Three controlled ablations on the prefix-heavy workload all came back null, which bounds the residual precisely:

ablation mechanism tested result
lazy retention reuse copy transfer (99.7% D2D hit rate measured) no change
batched sweep driver-call placement no change
map-ahead driver-call count (1 cuMemSetAccess/seq) no change

Together with py-spy profiles (arena-vs-paged on identical workloads), the residual ~5% is distributed host-side Python bookkeeping across the arena's per-request paths (onboarding, close write-out, reclaim scans) plus associated GPU-wait — no concentrated sink. Both benchmark modes ran the pure-Python v2 package; the arena executes more Python per request, so mypyc compilation is the largest remaining lever, followed by moving the (now single-point) map flush to a helper thread. The sweep defaults off until then. The P1 linear-addressing kernels remain the upside this storage layout exists to enable. Profiles and per-run reports are archived with the work log.

Super-page size curve and follow-up optimizations

Sweeping TRTLLM_KV_ARENA_PHYS_PAGE_SIZE_MB on the prefix-heavy (shared) workload shows page count — not map-call count — scales several per-request costs at once (teardown unmap chunks, budget arithmetic, mapped-range scans):

super-page size vs paged mapped/seq (68-block range) tail waste @ 256 seqs
16 MiB −5.6% 80 MiB ~3 GB
32 MiB −3.5% 96 MiB ~7 GB
64 MiB −3.0% 128 MiB ~15 GB

Recommended default for this workload shape: 32 MiB (16→32 gains 2.1 points; 32→64 only 0.5 more while doubling tail waste). Long-context workloads can go larger, memory-tight deployments smaller. (With adoption ON, steady-state map traffic is zero, so page size matters mainly for spill/teardown and tail waste.)

Follow-ups landed since:

  • Per-range mapped high-water frontier (Track per-range mapped frontier in contiguous KV arenas): ensure_mapped previously rescanned each range from its start on every growth call (O(blocks × pages) per sequence of pure-Python is_mapped checks — the inefficiency the page-size curve exposed). Ranges now track a mapped high-water frontier, making growth planning, retained-page counts, and reclaim O(1)/single-unmap. Recovers ~0.8 points at 16 MiB (−5.6% → −4.8% against a same-node paged baseline) without 32 MiB's tail-waste cost, and removes the two scanning helpers outright (net −18 lines).
  • mypyc compilation (the biggest lever per profiling): main no longer mypyc-compiles at all; PR [None][fix] Restore KVCacheManagerV2 mypyc compilability #15990 restores compilability (pre-existing issues, separate PR), and this branch carries the arena-side typing fixes. With the package hybrid-compiled (20/22 modules; two excluded for mypyc codegen bugs documented in [None][fix] Restore KVCacheManagerV2 mypyc compilability #15990), both modes speed up (paged +2.3%, arena +2.8%) and the gap narrows to −4.3%. That is a lower bound on full compilation's effect: the single hottest module (_core/_kv_cache.py, resize/commit per token) is one of the two still interpreted.
  • P1 linear-addressing kernels (opt-in, TRTLLM_KV_ARENA_LINEAR_KERNELS): XQA reads KV pages by linear arithmetic instead of table loads. E2e: +1.8% at short sequences (ISL/OSL 256/256, the µbench sweet spot), null at 1024/1024 — including under prefix aliasing in the KV-headroom regime (8,436 vs 8,466, −0.35%), so deduplicated, L2-friendlier prefix reads do not change the verdict. Graduation case is short-context serving regimes; default off.

Benchmarking also surfaced and fixed an arena-mode crash (commit "Handle intra-batch commit conflicts"): two in-flight requests committing the same system-prompt block asserted, since seq rebasing is disabled in arena mode; the loser now keeps its own pages as sequence-private committed copies referencing the existing tree block, and commit()'s block loop tolerates mid-loop virtual stops (hardening the classic path as well).

Hard-reclaim redesign + mappability-weighted canonical promotion (P3 v3)

Two commits (c9ac413fae, 7fd533b628) rework pressure-time reclaim and fix a pathology in how the reuse tree tracks residency.

The canonical-residency trap. Once a canonical block is evicted to host and its span pin dies, every later same-prefix admission copy-onboards from host, holds a sequence-private duplicate, and the tree never re-learns that a GPU copy exists — the prompt class copies from host forever. Fix: mappability-weighted promotion at both duplicate-creating seams (commit-conflict and onboarding). If the incumbent canonical is unmappable (no live registered span identity-covers it) and droppable, the GPU-resident challenger becomes the canonical instead of a private duplicate, and it adopts the incumbent's host slot as its pre-valid host backing — a future eviction is copy-free (verified byte-level: a corruption pattern written to the host copy survives the promoted owner's close, proving no D2H copy ran). The first descendant pays the copy once; its context-end commit registers a live span; every later descendant aliases.

Hard reclaim. hard_reclaim_gpu(min_pages) replaces the ladder's three lower rungs: a sync-free exit when nothing is reclaimable anywhere, otherwise ONE device quiesce under which everything harvestable is taken in destruction order — dedup remap of ALL flagged sequences first (duplicates have zero retention value), then parked terminated ranges, then registry excess above the protection floor; pausing stays the caller's fallback. A per-request dirty set (flagged at the duplicate-creating sites, cleared on scan attempt) gates the scan at O(1). The old ladder short-circuited at the first progressing rung, so the dedup rung never fired on any real workload, and the parked-range spill re-quiesced per range; both are gone (a unit test pins the discipline: zero syncs when idle, exactly one per harvest).

Validation (Llama-3.1-8B, H100, solo, interleaved):

shape config throughput dedup remapped pages
capacity edge (free_gpu_memory_fraction=0.7) pre-v3 arena+aliasing 7763.6 / 7785.9 0 / 0
capacity edge (0.7) v3 arena+aliasing 7852.9 / 7839.7 / 7885.3 (+0.9..1.4%) 187 / 184 / 150
saturated honest memory (0.85) paged 8593.1 / 8583.1
saturated honest memory (0.85) v3 arena+aliasing 8210.0 / 8198.9 (−4.47%, = pre-v3 pooled −4.5%±0.5) 31 / 24

Net: wins at the capacity edge where reclaim pressure is real, neutral at saturation; zero IMA / drops / alias misses across all runs. The second commit fixes a probe bug found during validation: the promotion mappability check originally reused the admission-path span lookup, whose verify loop invalidates a span on identity mismatch — probe chains legitimately diverge from the span owner's path past the shared prefix, so live spans were sporadically forfeited (47/2279 admissions missed). The probe is now a non-destructive O(1) identity check.

Mean-batch gap at saturation: root cause and fix (map-ahead margin)

Per-iteration instrumentation of the saturated 0.85 shape (scheduler batch, page-budget composition, and admission-refusal sites, correlated per iteration) decomposed the arena's −4.4% against paged into an admission-capacity story:

  • New context requests are admitted through resume(), whose utilization gate (max_util_for_resume, default 0.95) refuses while (used − reclaimable)/total > 0.95. On refusing iterations the gate-visible utilization sits exactly at 0.951 — the system is a regulator pinned at its setpoint, holding the arena at ~234 active sequences while paged (whose evictable blocks reclaim via synchronous free-list ops and whose gate essentially never fires — 566 checks failed in a whole run vs ~197k) runs client-bound at 256.
  • Of the 212 MB charged per active sequence at the pin, 187 MB is true KV capacity; the rest is the demand-paging map-ahead margin plus super-page rounding. The margin is charged to the page budget per sequence, so at the utilization ceiling it converts one-for-one into fewer admitted sequences — while its latency-hiding value does not measure, because at super-page granularity a sequence crosses a page boundary only every page_size/bytes_per_token tokens (128 tokens at 16 MiB for an 8B model).
  • Steady-phase means (excluding the identical ramp/drain phases, which had masked the effect) on Llama-3.1-8B, H100, 2500×(1024/1024)@conc 256, 0.85 fraction, 64 GiB host, prefix aliasing on:
config steady mean batch at 256 cap throughput
paged 232.9 77.8% 8615–8621
arena+aliasing, margin 1 (old default) 218.5–218.8 13% 8238–8251
arena+aliasing, margin 0 228.3 14% 8383 (+1.7%)
arena+aliasing, margin 0 + max_util_for_resume: 0.98 232.9 77.9% 8473 / 8485 (+2.8%)

This PR flips the map-ahead default to 0 (TRTLLM_KV_ARENA_MAP_AHEAD_PAGES, ContiguousArenaConfig.map_ahead_pages; the knob remains). Margin 0 is already routinely exercised — the mapping planner degrades the margin to 0 under pressure before failing a charge — and the flip also reduces pressure-machinery activity (gate refusals drop 4×, −5.4% iterations to completion, allocation failures unchanged, zero IMA across all runs).

With the margin honest, raising max_util_for_resume to 0.98 in the serving config makes the arena client-bound with the same steady batch and cap-residency as paged, leaving −1.6%, which profiling attributes to per-step host bookkeeping (the mypyc item above), with no capacity component. Note the ordering matters: 0.98 without the margin fix ran the pressure machinery hot (5× dedup harvests, ~1000 iterations with blocking mid-scheduling drains) and measured slower (−1.4%); the gate default is therefore left at 0.95 and 0.98 is a recommended config for arena workloads rather than a default change (the field is shared with classic v2).

Status

Ready for review. Since the last body update: P3 v3 hard-reclaim redesign + mappability-weighted canonical promotion (c9ac413fae + probe fix 7fd533b628), and the saturation mean-batch root cause + map-ahead default flip (a4fa141fe7, sections above; the flip closes the batch gap's dominant term and combined with max_util_for_resume: 0.98 in serving configs reaches batch parity with paged). Remaining follow-ups tracked for later phases: batched on_commit dispatcher (its write-through prior-event is on the manager stream — needs the same forward-stream treatment as 52a311ed99 before default use), helper-thread map flush, partial-block reuse, VSWA/SSM support, disaggregated serving (P2), driver-team escalation of the unmap-vs-in-flight-kernels fault (repro matrix ready).

Summary by CodeRabbit

  • New Features

    • Added experimental contiguous KV arena mode for KV cache management.
    • Introduced new arena-aware capacity, mapping, reclaim, and write-through behaviors.
    • Added public configuration options and APIs for arena-backed cache handling.
  • Bug Fixes

    • Improved handling of cache lifecycle operations in arena mode, including resume, suspend, close, and capacity growth.
    • Tightened validation to prevent unsupported configuration combinations when arena mode is enabled.
  • Tests

    • Added extensive CUDA unit coverage for contiguous arena allocation, mapping, reclamation, and end-to-end cache workflows.

…agerV2

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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
… 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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
… 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).

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
… 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).

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
…uous 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).

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
… 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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@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>
…or (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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
…tion (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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
…a 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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
… 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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
…nas (§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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
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%.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
…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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
@thorjohnsen
thorjohnsen marked this pull request as ready for review July 6, 2026 22:25
@thorjohnsen
thorjohnsen requested review from a team as code owners July 6, 2026 22:25
@thorjohnsen
thorjohnsen requested a review from hchings July 6, 2026 22:25
@thorjohnsen

Copy link
Copy Markdown
Collaborator Author

/bot run

@coderabbitai

coderabbitai Bot commented Jul 6, 2026

Copy link
Copy Markdown
Contributor

Review Change Stack

📝 Walkthrough

Walkthrough

This PR introduces an experimental "contiguous KV cache arena" mode that reserves per-sequence contiguous virtual-address block ranges and maps physical pages on demand via a shared page budget, replacing scattered GPU slot allocation. Changes span configuration, CUDA virtual memory, storage layers, StorageManager/KVCacheManager APIs, per-sequence lifecycle logic, pyexecutor wiring, and new tests.

Changes

Contiguous KV Arena Mode

Layer / File(s) Summary
Configuration and public exports
tensorrt_llm/llmapi/llm_args.py, tensorrt_llm/runtime/kv_cache_manager_v2/_config.py, tensorrt_llm/runtime/kv_cache_manager_v2/__init__.py
Adds use_contiguous_kv_arena (forces use_kv_cache_manager_v2), ContiguousArenaConfig/WriteThroughPolicy, KVCacheManagerConfig.contiguous_arena with post-init compatibility checks, and re-exports.
Sparse VirtMem mapping
tensorrt_llm/runtime/kv_cache_manager_v2/_cuda_virt_mem.py
Adds _page_map-backed sparse chunk mapping (map_range/unmap_range/is_mapped/num_sparse_chunks) exclusive of tail-stack mode, with teardown cleanup.
SequenceArena allocator
tensorrt_llm/runtime/kv_cache_manager_v2/_sequence_arena.py, tensorrt_llm/runtime/kv_cache_manager_v2/setup_mypyc.py
New PageBudget, BlockRangeAllocator, and SequenceArena implementing per-pool-group reservation, demand mapping, and deferred/immediate reclaim; module added to mypyc compile list.
Private committed page model
tensorrt_llm/runtime/kv_cache_manager_v2/_page.py
Adds convert_to_private_committed and PrivateCommittedPage for sequence-private reuse copies, and refines eviction-exclusion logic in _PageHolder.__del__.
Arena pool group and cache-level storage
tensorrt_llm/runtime/kv_cache_manager_v2/_storage/_core.py
Adds ArenaSlotPool, SequenceRange, ArenaPoolGroup (with reserve/take/ensure_mapped/queue_mapping/drain_reclaim/ghost registry), and GpuArenaCacheLevelStorage.
StorageManager arena API
tensorrt_llm/runtime/kv_cache_manager_v2/_storage_manager.py
Adds arena-mode public API (reserve/take/ensure/free/drain/spill/queue/flush, ghosts, onboarding, offload, write-through), guards against scattered GPU allocation/eviction/rebalancing, and updates utilization/eviction logic.
KVCacheManager wiring
tensorrt_llm/runtime/kv_cache_manager_v2/_core/_kv_cache_manager.py
Forwards arena_config, validates/threads max_capacity, disables rebasing/partial-match in arena mode, adds drain_gpu_reclaim/flush_gpu_mappings, and updates clamp_max_seq_len_for_mem for page-budget accounting.
_KVCache arena lifecycle
tensorrt_llm/runtime/kv_cache_manager_v2/_core/_kv_cache.py
Reworks close/suspend/resume/resize/_commit_block for arena mode: reserved-range mapping, write-through, onboarding, and private-committed-copy conflict handling.
pyexecutor integration
tensorrt_llm/_torch/pyexecutor/kv_cache_manager_v2.py
Adds _arena_enabled detection, _build_arena_config() from environment variables, and arena-mode overrides across buffer access, free-block accounting, resource preparation, dummy requests, invalid-value checks, and shutdown.
Tests
tests/unittest/_torch/executor/test_kv_cache_v2_contiguous_arena.py, tests/unittest/kv_cache_manager_v2_tests/test_contiguous_arena.py
New unit and integration test modules covering config wiring, adapter lifecycle, arena primitives, storage layers, and end-to-end KV manager behavior under arena mode.

Estimated code review effort: 5 (Critical) | ~120 minutes

Sequence Diagram(s)

sequenceDiagram
  participant KVCacheManager
  participant KVCache as _KVCache
  participant StorageManager
  participant ArenaPoolGroup
  participant SequenceArena
  KVCacheManager->>KVCache: create_kv_cache(max_capacity)
  KVCache->>StorageManager: reserve_gpu_sequence(pg_idx)
  StorageManager->>ArenaPoolGroup: reserve_sequence(range)
  ArenaPoolGroup->>SequenceArena: allocate(range)
  KVCache->>StorageManager: ensure_gpu_mapped(range)
  StorageManager->>ArenaPoolGroup: ensure_mapped(range)
  ArenaPoolGroup->>SequenceArena: map physical pages
  KVCache->>KVCache: close()/suspend()
  KVCache->>StorageManager: free_gpu_sequence(range, last_consumer)
  StorageManager->>ArenaPoolGroup: free_sequence(range)
  ArenaPoolGroup->>SequenceArena: reclaim(range)
Loading

Suggested reviewers: joyang-nv, bo-nv, xinhe-nv, JunyiXu-nv, pengbowang-nv, yechank-nvidia

🚥 Pre-merge checks | ✅ 5
✅ Passed checks (5 passed)
Check name Status Explanation
Docstring Coverage ✅ Passed No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check.
Linked Issues check ✅ Passed Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check ✅ Passed Check skipped because no linked issues were found for this pull request.
Title check ✅ Passed The title matches the main feature and follows the repository's ticket/type pattern.
Description check ✅ Passed The description covers the summary, design, and testing well, with only the PR Checklist section left incomplete.
✨ Finishing Touches
🧪 Generate unit tests (beta)
  • Create PR with unit tests

Comment @coderabbitai help to get the list of available commands.

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 2

🧹 Nitpick comments (1)
tests/unittest/kv_cache_manager_v2_tests/test_contiguous_arena.py (1)

968-974: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Consolidate duplicated _block_floats/_host_floats helpers.

_block_floats is redefined verbatim in TestArenaKVCacheManagerEndToEnd and TestArenaLazyRetention, and _host_floats duplicates the ctypes-buffer pattern already used inline in TestStorageManagerArenaMode._slot_floats. Extracting these into a shared module-level helper (or a small mixin) avoids drift if the addressing logic ever changes.

♻️ Proposed consolidation sketch
+def _gpu_block_floats(storage, page_index: int, n_floats: int) -> "torch.Tensor":
+    addr = int(storage.slot_address(GPU_LEVEL, PoolGroupIndex(0), page_index, PoolIndex(0)))
+    return TestSparseVirtMem._tensor_at(addr, n_floats)
+
+
+def _host_slot_floats(storage, slot_id: int, n_floats: int) -> "torch.Tensor":
+    import ctypes
+
+    addr = int(storage.slot_address(CacheLevel(1), PoolGroupIndex(0), slot_id, PoolIndex(0)))
+    buf = (ctypes.c_float * n_floats).from_address(addr)
+    return torch.frombuffer(buf, dtype=torch.float32)

Then have TestArenaKVCacheManagerEndToEnd._block_floats and TestArenaLazyRetention._block_floats/_host_floats delegate to these.

Also applies to: 1339-1357

🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@tests/unittest/kv_cache_manager_v2_tests/test_contiguous_arena.py` around
lines 968 - 974, Consolidate the duplicated `_block_floats` and `_host_floats`
logic used in `TestArenaKVCacheManagerEndToEnd`, `TestArenaLazyRetention`, and
`TestStorageManagerArenaMode` by extracting shared module-level helper(s) or a
small mixin. Keep the addressing/buffer creation behavior in one place so
`_block_floats` delegates to the shared helper and `_host_floats` reuses the
same ctypes-buffer pattern instead of reimplementing it inline. Update the test
methods to call the shared helper(s) and remove the verbatim duplicate
implementations.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@tests/unittest/_torch/executor/test_kv_cache_v2_contiguous_arena.py`:
- Around line 103-152: Coverage for the arena adapter path is incomplete: add
tests in TestArenaAdapter that directly exercise get_buffers() raising
NotImplementedError when _arena_enabled, and check_invalid_values_in_kv_cache()
returning False without scanning in arena mode. Also add a test around
prepare_resources(), add_dummy_requests(), or extend_capacity_for_tokens() that
spies on self.mgr.impl to verify flush_gpu_mappings() and drain_gpu_reclaim()
are invoked. Finally, replace the manual kv_cache_map/index_mapper cleanup in
test_request_lifecycle_through_adapter with the public free_resources() path (or
add a separate test) so the full teardown flow and close() behavior are
validated through the production route.
- Around line 96-100: The defaults test for KVCacheManagerV2._build_arena_config
is not actually isolating the TRTLLM_KV_ARENA_* environment, so it can inherit
ambient values and become flaky. Update test_env_defaults in
test_kv_cache_v2_contiguous_arena to explicitly clear or override the relevant
arena env vars before calling KVCacheManagerV2._build_arena_config, rather than
using a no-op os.environ patch. Keep the assertions on phys_page_size and
write_through as the check for true default behavior.

---

Nitpick comments:
In `@tests/unittest/kv_cache_manager_v2_tests/test_contiguous_arena.py`:
- Around line 968-974: Consolidate the duplicated `_block_floats` and
`_host_floats` logic used in `TestArenaKVCacheManagerEndToEnd`,
`TestArenaLazyRetention`, and `TestStorageManagerArenaMode` by extracting shared
module-level helper(s) or a small mixin. Keep the addressing/buffer creation
behavior in one place so `_block_floats` delegates to the shared helper and
`_host_floats` reuses the same ctypes-buffer pattern instead of reimplementing
it inline. Update the test methods to call the shared helper(s) and remove the
verbatim duplicate implementations.
🪄 Autofix (Beta)

Fix all unresolved CodeRabbit comments on this PR:

  • Push a commit to this branch (recommended)
  • Create a new PR with the fixes

ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 0d56f573-44ed-4c74-add5-a96010929eb4

📥 Commits

Reviewing files that changed from the base of the PR and between b1ea26c and aa6d2eb.

📒 Files selected for processing (14)
  • tensorrt_llm/_torch/pyexecutor/kv_cache_manager_v2.py
  • tensorrt_llm/llmapi/llm_args.py
  • tensorrt_llm/runtime/kv_cache_manager_v2/__init__.py
  • tensorrt_llm/runtime/kv_cache_manager_v2/_config.py
  • tensorrt_llm/runtime/kv_cache_manager_v2/_core/_kv_cache.py
  • tensorrt_llm/runtime/kv_cache_manager_v2/_core/_kv_cache_manager.py
  • tensorrt_llm/runtime/kv_cache_manager_v2/_cuda_virt_mem.py
  • tensorrt_llm/runtime/kv_cache_manager_v2/_page.py
  • tensorrt_llm/runtime/kv_cache_manager_v2/_sequence_arena.py
  • tensorrt_llm/runtime/kv_cache_manager_v2/_storage/_core.py
  • tensorrt_llm/runtime/kv_cache_manager_v2/_storage_manager.py
  • tensorrt_llm/runtime/kv_cache_manager_v2/setup_mypyc.py
  • tests/unittest/_torch/executor/test_kv_cache_v2_contiguous_arena.py
  • tests/unittest/kv_cache_manager_v2_tests/test_contiguous_arena.py

Comment on lines +96 to +100
def test_env_defaults(self) -> None:
with mock.patch.dict("os.environ", {}, clear=False):
arena_cfg = KVCacheManagerV2._build_arena_config()
self.assertEqual(arena_cfg.phys_page_size, 2 * MiB)
self.assertEqual(arena_cfg.write_through, WriteThroughPolicy.ON_FREE)

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win

Isolate arena env vars explicitly instead of relying on ambient environment.

mock.patch.dict("os.environ", {}, clear=False) is a no-op for isolation — it neither clears nor guarantees absence of TRTLLM_KV_ARENA_* vars that might already be set in the runner's environment, so this "defaults" test can silently validate the wrong state or flake depending on the ambient shell/CI environment.

🧪 Proposed fix: explicitly clear the relevant keys
     def test_env_defaults(self) -> None:
-        with mock.patch.dict("os.environ", {}, clear=False):
+        keys = (
+            "TRTLLM_KV_ARENA_PHYS_PAGE_SIZE_MB",
+            "TRTLLM_KV_ARENA_MAP_AHEAD_PAGES",
+            "TRTLLM_KV_ARENA_WRITE_THROUGH",
+            "TRTLLM_KV_ARENA_LAZY_RETENTION",
+            "TRTLLM_KV_ARENA_BATCHED_MAP",
+        )
+        with mock.patch.dict("os.environ", {}, clear=False):
+            for k in keys:
+                os.environ.pop(k, None)
             arena_cfg = KVCacheManagerV2._build_arena_config()

As per path instructions for tests/**, coverage/robustness of environment-dependent tests should be verified explicitly rather than relying on ambient state.

📝 Committable suggestion

‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.

Suggested change
def test_env_defaults(self) -> None:
with mock.patch.dict("os.environ", {}, clear=False):
arena_cfg = KVCacheManagerV2._build_arena_config()
self.assertEqual(arena_cfg.phys_page_size, 2 * MiB)
self.assertEqual(arena_cfg.write_through, WriteThroughPolicy.ON_FREE)
def test_env_defaults(self) -> None:
keys = (
"TRTLLM_KV_ARENA_PHYS_PAGE_SIZE_MB",
"TRTLLM_KV_ARENA_MAP_AHEAD_PAGES",
"TRTLLM_KV_ARENA_WRITE_THROUGH",
"TRTLLM_KV_ARENA_LAZY_RETENTION",
"TRTLLM_KV_ARENA_BATCHED_MAP",
)
with mock.patch.dict("os.environ", {}, clear=False):
for k in keys:
os.environ.pop(k, None)
arena_cfg = KVCacheManagerV2._build_arena_config()
self.assertEqual(arena_cfg.phys_page_size, 2 * MiB)
self.assertEqual(arena_cfg.write_through, WriteThroughPolicy.ON_FREE)
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@tests/unittest/_torch/executor/test_kv_cache_v2_contiguous_arena.py` around
lines 96 - 100, The defaults test for KVCacheManagerV2._build_arena_config is
not actually isolating the TRTLLM_KV_ARENA_* environment, so it can inherit
ambient values and become flaky. Update test_env_defaults in
test_kv_cache_v2_contiguous_arena to explicitly clear or override the relevant
arena env vars before calling KVCacheManagerV2._build_arena_config, rather than
using a no-op os.environ patch. Keep the assertions on phys_page_size and
write_through as the check for true default behavior.

Source: Path instructions

Comment on lines +103 to +152
@requires_cuda
class TestArenaAdapter(unittest.TestCase):
def setUp(self) -> None:
self.mgr = _create_arena_manager()

def tearDown(self) -> None:
self.mgr.shutdown()
del self.mgr

def test_arena_mode_enabled(self) -> None:
self.assertTrue(self.mgr._arena_enabled)
self.assertTrue(self.mgr.impl._storage.is_arena_mode)

def test_capacity_accounting_is_physical(self) -> None:
# get_num_free_blocks counts what the page budget can back, not VA
budget = self.mgr.impl._storage.gpu_page_budget
page_size = self.mgr.impl._storage.arena_config.phys_page_size
expected_blocks = budget.total_pages * page_size // BLOCK_COLUMN_BYTES
self.assertEqual(self.mgr.get_num_free_blocks(), expected_blocks)
self.assertEqual(budget.total_pages, GPU_QUOTA // page_size)
# clamp is page-based too and must not exceed physical capacity
available = self.mgr.get_num_available_tokens(token_num_upper_bound=10**6)
self.assertLessEqual(available, budget.total_pages * page_size // BYTES_PER_TOKEN)
self.assertGreater(available, 0)

def test_request_lifecycle_through_adapter(self) -> None:
mgr = self.mgr
budget = mgr.impl._storage.gpu_page_budget
kv = mgr._create_kv_cache(1, None, None)
self.assertIsNotNone(kv)
# the VA reservation is sized for the per-request maximum
self.assertEqual(kv._max_capacity, mgr.max_blocks_per_seq * TOKENS_PER_BLOCK)
kv.cuda_stream = mgr._stream.cuda_stream
self.assertTrue(kv.resume(mgr._stream.cuda_stream))
self.assertTrue(kv.resize(3 * TOKENS_PER_BLOCK))
# the adapter pre-sizes index buffers to max_blocks_per_seq; only the
# first num_blocks entries are valid -- and they must be consecutive
indices = list(kv.get_base_page_indices(0))[:3]
self.assertEqual(indices, list(range(indices[0], indices[0] + 3)))
self.assertGreater(budget.used_pages, 0)
# suspend/resume via the adapter-facing kv methods
kv.suspend()
self.assertTrue(kv.resume(mgr._stream.cuda_stream))
kv.close()
mgr.kv_cache_map.pop(1)
mgr.index_mapper.remove_sequence(1)
torch.cuda.synchronize()
self.assertGreaterEqual(mgr.impl.drain_gpu_reclaim(), 1)
self.assertEqual(budget.used_pages, 0)

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🎯 Functional Correctness | 🟠 Major | 🏗️ Heavy lift

Coverage gap: adapter-level arena integration points untested here.

This file is described as covering "KVCacheManagerV2 adapter plumbing," but the new adapter-level arena behavior isn't directly exercised:

  • get_buffers() raising NotImplementedError in arena mode.
  • check_invalid_values_in_kv_cache() returning False and skipping the scan.
  • The flush_gpu_mappings()/drain_gpu_reclaim() calls newly added to prepare_resources(), add_dummy_requests(), and extend_capacity_for_tokens().

Additionally, test_request_lifecycle_through_adapter tears down via manual kv_cache_map.pop/index_mapper.remove_sequence instead of the public free_resources() path, so the arena close() behavior isn't validated through the same route production code uses.

Suggest adding to test_kv_cache_v2_contiguous_arena.py:

  • A test asserting get_buffers() raises NotImplementedError when _arena_enabled.
  • A test asserting check_invalid_values_in_kv_cache() returns False without raising in arena mode.
  • A test that drives a request through prepare_resources()/add_dummy_requests() (mirroring _KVCache growth) to confirm flush_gpu_mappings/drain_gpu_reclaim are actually invoked (e.g. via mock/spy on self.impl).
  • Replacing the manual teardown in test_request_lifecycle_through_adapter with mgr.free_resources(req) (or adding a second test that does), to validate the full production teardown path in arena mode.

As per path instructions, "Act as a QA engineer reviewing test changes and coverage... suggest concrete list file names and whether coverage is sufficient, insufficient, or needs follow-up outside the PR": coverage here is insufficient for the adapter-level integration points that are the stated purpose of this file.

🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

In `@tests/unittest/_torch/executor/test_kv_cache_v2_contiguous_arena.py` around
lines 103 - 152, Coverage for the arena adapter path is incomplete: add tests in
TestArenaAdapter that directly exercise get_buffers() raising
NotImplementedError when _arena_enabled, and check_invalid_values_in_kv_cache()
returning False without scanning in arena mode. Also add a test around
prepare_resources(), add_dummy_requests(), or extend_capacity_for_tokens() that
spies on self.mgr.impl to verify flush_gpu_mappings() and drain_gpu_reclaim()
are invoked. Finally, replace the manual kv_cache_map/index_mapper cleanup in
test_request_lifecycle_through_adapter with the public free_resources() path (or
add a separate test) so the full teardown flow and close() behavior are
validated through the production route.

Source: Path instructions

…(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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
…-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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
… host buffers

Feature-branch adaptation of the standalone fix proposed for main in
PR NVIDIA#16088 (cherry of 3bdee3b): 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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
@thorjohnsen
thorjohnsen requested a review from a team as a code owner July 8, 2026 16:59
@thorjohnsen
thorjohnsen requested a review from lfr-0531 July 8, 2026 16:59
…type)

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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
…gage

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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
… 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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
…ot 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.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
… 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)

…
The map-ahead margin (a latency-hiding hint, default 1 page/pool) was
charged atomically with the frontier: at the arena's capacity edge a
satisfiable growth failed OutOfPagesError because frontier+margin did
not fit, triggering reclaim drains, retained-range spills, and
ultimately scheduler backoff -- for pages nobody needs yet. The
shared-prefix benchmark sits exactly at that edge (256 concurrent
sequences x 1 margin page x 16 MiB = 4 GiB of transient margin charge
against an exactly-consumed budget).

Fix: the margin is now best-effort. SequenceArena.ensure_mapped
retries the plan with margin=0 on budget exhaustion and raises only
if the frontier alone does not fit. No unmaps are introduced (margin
pages already mapped are simply consumed as the frontier advances),
so the driver-bug mitigation surface (zero steady-state unmap) is
unchanged. Off the edge, behavior is identical: the degrade branch
only runs where the old code raised.

The batched map sweep charges at queue time and maps at flush time,
so the degrade decision is recorded on the pending entry (the charged
margin) and the flush replays exactly that plan -- preserving the
charge==plan determinism the flush relies on. A frontier bump on a
margin-charged entry that no longer fits degrades the entry and can
even RELEASE budget (margin wider than the added blocks).

Tests: sync degrade maps frontier-only at the edge and still raises
cleanly when the frontier itself does not fit; queued degrade flushes
exactly the charged pages; delta-degrade releases the excess exactly.
Arena suite 106/106, adapter 12/12, v2 regression 217 passed / 12
skipped on H100.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
…ve aliasing)

Close-time span registration (P3 v1) made the whole first concurrent
wave of a prompt class copy the shared prefix: nothing was aliasable
until the first same-prompt sequence CLOSED, one full request lifetime
after its prefix became immutable. Paged v2 shares context blocks the
moment they commit; this brings the arena to parity.

The owner's first commit burst (context end -- prefill commits all
input blocks at once) now registers the span, so same-prefix
admissions alias the prefix while the owner is still generating. The
entry is OWNER-CHARGED while the owner lives: the pages are its
active KV, so there is no budget transfer, no retained marking (the
resume gate must not count them reclaimable), and pressure spills
skip the entry (dropping the pin frees nothing -- the handles stay
alive through the owner's own mappings). free_sequence -- the single
choke point that close, evacuate, and error unwinds all pass through
-- flips the entry to the registry-charged state, producing exactly
the accounting close-time registration used to.

Registration walks the leading CANONICAL run from ordinal 0: a
commit-conflict loser's blocks are private copies and stop the walk
at block 0, so losers register nothing and the winner's entry serves
the prompt class. (Remapping losers onto the winner's pages -- a
batched, quiesced unmap+alias at the context->generation flip -- is
tracked as a follow-up.) Consumers order their first reads after a
commit-time event on the owner's stream, the same discipline as the
write-through copies. V1 scope unchanged: single-non-SSM-lc pool
groups, extent-exact adoption signatures, partial-span lookups.

Tests: owner-charged lifecycle at the pool-group seam (no budget
movement, spill-exempt, lookup/alias live, consumer-first close does
not flip, owner close flips to close-time-identical accounting); a
manager-level end-to-end where the consumer admits while the owner is
mid-generation and physical sharing is proven by writing through the
owner's VA and reading the change through the consumer's (a copy
could not see it). Arena suite 108/108, adapter 12/12, v2 regression
219 passed / 12 skipped on H100.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
…v2 remap)

Sequences holding private duplicates of a registered canonical prefix
-- commit-conflict losers (concurrent same-prompt prefills) and
sequences that copy-onboarded before a span existed -- keep those
duplicate pages for their whole lifetime. Under pressure the arena
would spill the registry or preempt a sequence while reclaimable
duplicates sat in live ranges.

Fix: a pressure-time dedup remap, as a new rung in the reclaim ladder
between the parked-range spill and the registry-excess spill:

  drain reclaim -> spill parked ranges -> DEDUP REMAP ->
  spill registry excess -> back off (preempt)

The scan is STATELESS (no commit-time loser tracking): walk live
sequences' leading committed blocks; a dedupable run is
PrivateCommittedPage holders whose radix entries are live canonical
pages covered by a registered span (identity-verified). The remap is a
same-VA swap: unmap the private chunks and cuMemMap the span's handles
at identical positions, so slot ids, block offsets, page holders, and
frontiers are all untouched -- page objects are slot-indexed, not
physical-handle-indexed, so no per-sequence state changes at all. The
duplicates' budget charge is released (memory freed from LIVE
sequences with zero information loss) and the range gains the alias
signature (prefix-affine parking at close). Largest-savings-first,
stops at min_pages.

Sync accounting: the whole batch runs behind ONE quiesce
(quiesce_before_unmap, factored out of the reclaim path), paid only
when the ladder's earlier rungs freed nothing -- exactly when the
alternative is a preemption (evacuation copies + resume onboarding),
which costs strictly more. Memory saving only matters under pressure,
so deferring dedup to the pressure path loses nothing. If re-aliasing
fails after an unmap, the pool's original handles are restored
(references held across the swap) -- a range is never left with holes.

The scan probes the span registry without touching the admission
hit/miss counters (count_stats=False); remapped pages are reported at
shutdown ("dedup remapped pages=N").

Tests: pool-group same-VA swap (bytes flip to canonical, write-through
-A-visible-through-B physical proof, budget released, signature
stamped, growth still works); manager-level commit-conflict loser
dedup (counters clean, idempotent second scan); a reclaim-ladder
end-to-end at a 2-page quota where growth succeeds ONLY via the dedup
rung -- registry unspilled, no preemption. Arena suite 111/111,
adapter 12/12, v2 regression 222 passed / 12 skipped on H100.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
@nvpohanh
nvpohanh requested a review from lowsfer July 9, 2026 06:30
@nvpohanh

nvpohanh commented Jul 9, 2026

Copy link
Copy Markdown
Collaborator

[by Codex] @lowsfer Could you review this PR? Thanks!

…romotion (P3 v3)

Replaces the reclaim ladder's three lower rungs with hard_reclaim_gpu:
fenced drain, dedup remap of ALL flagged sequences, parked-range spill,
and registry-excess spill -- destruction-ordered (duplicates first: they
have zero retention value) behind ONE device quiesce, with a sync-free
exit when nothing is reclaimable anywhere. Fixes the ladder short-circuit
that starved the dedup rung and the per-range re-sync in the parked spill.

Per-request dedup dirty set: sequences are flagged at the duplicate-
creating sites (commit-conflict loss to a mappable incumbent, copy/resume
onboarding that stays private) and scanned ONLY while flagged; flags
clear on the scan attempt.

Mappability-weighted promotion (both duplicate-creating seams): when the
incumbent canonical is unmappable (no live span identity-covers it) and
droppable, the GPU-resident challenger becomes the canonical instead of
a private duplicate, adopting the incumbent's host slot as its pre-valid
host backing (future eviction is copy-free). This breaks the canonical-
residency trap: previously, once a canonical was evicted to host and its
span pin died, every same-prefix admission copied from host forever.

Tests: promotion breaks the trap end-to-end (slot-identity adoption,
zero-copy descendant aliasing, copy-free close), commit-seam promotion,
promotion-window race healing into productive dedup, dedup-preserves-
parked-ranges ordering, and the sync discipline (zero syncs idle,
exactly one per harvest).

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
The P3 v3 promotion probe reused lookup_arena_canonical_span, whose
verify loop invalidates the span registry entry on any identity
mismatch. That is correct for admissions (matched pages share the span
owner's tree path, so a mismatch means staleness) but destructive for
promotion probes, whose chains legitimately diverge from the owner's
path past the shared prefix -- observed as spurious alias misses
(47/2279 admissions) on the f070 validation shape.

Replace it with ArenaPoolGroup.span_covers: an O(1) identity check of
span.page_refs[ordinal] against the incumbent, keyed by the chain's
head canonical, touching nothing. Also drops the now-unneeded
chain-walk helper.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
The map-ahead margin is charged against the GPU page budget per
sequence, so near the admission-utilization ceiling it converts
one-for-one into fewer admitted sequences, while its latency-hiding
value does not measure: at super-page granularity a sequence crosses a
page boundary only every page_size/bytes_per_token tokens, and the
mapping planner already degrades the margin to zero under pressure
before failing a charge.

Measured on Llama-3.1-8B (H100, 2500x(1024/1024) @ concurrency 256,
free_gpu_memory_fraction 0.85, 16 MiB pages, prefix aliasing on),
margin 1 -> 0: steady-phase mean batch 218.8 -> 228.3 (paged: 232.9),
throughput 8241 -> 8383 tok/s (+1.7%), iterations to completion -5.4%,
admission-gate refusals -4x, allocation failures unchanged, no IMA.
With max_util_for_resume: 0.98 additionally set in the serving config,
the arena reaches steady-batch parity with the paged baseline (232.9,
77.9% of iterations at the 256 cap); that gate default is left
unchanged because the field is shared with classic V2 and raising it
without this margin fix measured slower (pressure-machinery churn).

The deadlock-recovery test sized its 2-page budget assuming margin 1
mapped the second page; it now spans both pages with the request
itself (33 blocks at 32 blocks/page), margin-independent.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
…offloads

Two arena-mode fixes for drop-heavy host-exhaustion shapes, found by
instrumenting the free-count invariant on the tinyhost reproducer:

1. close() kept strong references to the pages it had just handed to
   the host eviction controller via adopt_stale_copies. Under host
   exhaustion (everything else HELD by suspended sequences) those
   just-adopted pages are the only droppable pages on the level, so
   the subsequent offload's prepare_free_slots evicted one that could
   not die: no slot was freed and _prepare_free_slots's free-count
   assert fired mid-run. With on_free write-through the adopt set is
   populated only by the P3 v3 promotion's stolen host slots, which is
   why this surfaced after promotion landed. Drop the references as
   soon as the controller owns the pages.

2. offload_arena_pages could offload pages whose radix-tree blocks had
   died mid-call: its own prepare_free_slots eviction may collapse a
   subtree containing blocks of pages the same call is about to
   offload. The now-treeless page was migrated to host and scheduled
   for eviction anyway -- an orphan the shutdown sweep's tree-driven
   exclusion can never find (the long-standing clear_reusable_blocks
   num_evictable_pages shutdown assert) and an unreachable-for-reuse
   host copy wasting a slot. Skip droppable committed pages whose
   block is gone or orphaned; held suspension state is never skipped.
   A page skipped this way also ends the prefix-aliasing span
   registration run in close().

One new manager-level test pins both: reverting fix 1 reproduces the
exact production free-count assert; reverting fix 2 reproduces the
shutdown assert. Validated e2e on the tinyhost reproducer (6/6 clean
runs across both aliasing arms, previously ~1/4 fatal no-alias and
2/2 shutdown asserts with aliasing).

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
@nvpohanh

Copy link
Copy Markdown
Collaborator

[by Codex] @lowsfer Friendly reminder to review this PR as the primary KV-cache-manager reviewer when you have a chance. Thanks!

The trtllm-gen context/generation FMHA builds a flat KV-pool view in
buildFlashinferTrtllmGenPagedKvCacheBuffers via torch::from_blob's
options-only overload, which resolves the tensor device by probing the
base pointer with getDeviceFromPtr. For a classic paged pool (one large
cudaMalloc) the base is always resident, so this succeeds.

The contiguous KV arena instead reserves one large VA range and maps
physical pages on demand, so the pool base VA (slot 0) is not guaranteed
to be backed at context-preprocess time -- e.g. a KV-cache-estimation
pool sized with zero slack at small max_batch_size leaves slot 0
unmapped. getDeviceFromPtr then reports the base as host memory and the
forward aborts with "The specified pointer resides on host memory and is
not registered with any CUDA device" (reproducible at max_batch_size=4;
8+ has enough estimation slack to keep slot 0 mapped).

The base pointer is only an addressing origin: the kernel indexes into
mapped blocks via block offsets and never dereferences offset 0 unless a
block table points there. Build both KV-pool views (kv_pool and the fp4
block-scale pool) with an explicit target_device so the tensor takes the
current CUDA device directly and skips the residency probe. Classic
fully-resident pools are unaffected (identical tensor); the arena no
longer requires slot-0 residency.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
@thorjohnsen
thorjohnsen requested review from a team as code owners July 17, 2026 02:36
@mikeiovine
mikeiovine removed the request for review from pcastonguay July 20, 2026 15:49
@thorjohnsen
thorjohnsen marked this pull request as draft July 20, 2026 19:59
A canonical span whose physical pages a live sequence still aliases must
not be spilled: dropping the registry pin frees no memory (the aliases
keep the handles resident) and only evicts a still-usable prefix from the
reuse lookup, forcing later same-prefix admissions to re-copy pages that
are still on the GPU. Skip such spans in spill_canonical_spans (positive
external refcount) so a span leaves the reuse lookup only when it is
actually reclaimable, i.e. once its last aliaser frees.

Tests: add test_spill_skips_span_with_live_alias and update the two tests
that asserted the previous spill-while-aliased behavior.

Signed-off-by: thorjohnsen <41591019+thorjohnsen@users.noreply.github.com>
@thorjohnsen
thorjohnsen force-pushed the feat/contiguous-primary-kvcache-p0 branch from ef2a77e to af7b5f6 Compare July 21, 2026 21:00
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants