diff --git a/.agents/engine-matrix.md b/.agents/engine-matrix.md
index 835adf787..3f0980fcd 100644
--- a/.agents/engine-matrix.md
+++ b/.agents/engine-matrix.md
@@ -280,7 +280,7 @@ claims it.
 | `LOAD-SAFETENSORS` | General safetensors loading, stacked parameters, weights mapping. Binding memory scan finds all source mappings live through load plus a persistent CPU mirror of selected tensors; windowed progressive `madvise(DONTNEED)` release now drops each copied-then-dead source range during the copy loop so the mirror build no longer double-resides with the full source mmap | T0 | streamed target-device construction `vllm/model_executor/model_loader/base_loader.py:43-82`; incremental `safe_open` yield `vllm/model_executor/model_loader/weight_utils.py:905-954`; immediate parameter copy `vllm/model_executor/models/utils.py:170-180,252-279` | all mappings opened/retained `src/vllm/entrypoints/model_loader.cpp:47-63,303-312`; `MAP_PRIVATE` reader `src/vllm/model_executor/model_loader/safetensors_reader.cpp:43-70`; windowed release primitive/gate `src/vllm/model_executor/model_loader/safetensors_reader.cpp:285-340`; copy-helper instrumentation `src/vllm/model_executor/models/qwen3_5_weights.cpp:88,104,119,171,190,226,231` + `qwen3_5_dense_weights.cpp:64,80,109,159,164,263,323` | [reader contracts + windowed-release cases](../tests/vllm/test_safetensors.cpp#L68) green (smaps-Rss drop, byte-identity, gate semantics, neighbor safety); exact 27B accounting finds **1,155 tensors / 24,610,136,064 B (22.920 GiB)** persistent host bytes. **VmHWM A/B MEASURED** at `cb2d310` (root `~/work/vllm.cpp-windowed-load/cb2d310c…518/evidence`, one flock, single 27B load per arm): OFF **48,285,916 kB** vs ON **24,750,704 kB = VmRSS** (**−23.54 GB**, load transient eliminated), ON-arm smoke 6/6; ledger row 2026-07-15. Exact-grid memory axes remain `FAILED` (projected PASS); direct-device streaming gate still open | [specs/safetensors-windowed-load.md](specs/safetensors-windowed-load.md) | `PARTIAL` | - |
 | `ENG-LOAD-DIRECT-UPLOAD` | Cut the SECOND copy out of a safetensors load: a weight the device consumes VERBATIM VIEWS the read-only shard mmap (`OwnedBytes::Borrow`, keep-alive carried on `StTensor::mapping`) instead of being copied into an owned host buffer, so `ResidentWeight` uploads straight from the file mapping and the load moves those bytes ONCE. Backend- and arch-agnostic: it lives in the shared dense loader helpers, so every arch that routes through them inherits it. Additive leaf below `LOAD-SAFETENSORS`; the windowed source-page release is preserved and now also runs after the device upload. Non-verbatim paths (transpose, dtype conversion, dequant, concatenation, GGUF load-time repack) are unchanged and keep their copy. Also lands the measurement half of issue #150: per-phase load timing plus host-copy / borrowed / device-upload byte counters (`VT_LOAD_STATS`) | T1 | one move per byte: streamed shard iterator `vllm/model_executor/model_loader/weight_utils.py:905-954`; immediate copy into the destination parameter `vllm/model_executor/models/utils.py:252-279`; driver `vllm/model_executor/model_loader/base_loader.py:43-82` @ `555967922` | refcounted mapping `src/vllm/model_executor/model_loader/safetensors_reader.cpp:46-80,231-262`; byte counters `:337-400`; `BorrowStTensorBytes` + post-upload adoption/release `src/vllm/model_executor/models/qwen3_5_weights.cpp`; qualifying call sites `include/vllm/model_executor/models/dense_weight_loaders.h` (`LoadBf16Direct`, `LoadCtNvfp4W4A16`, `LoadCtMxfp4W4A16`), `src/vllm/model_executor/models/qwen3_5_dense_weights.cpp` (`LoadModelBf16Direct`, `LoadCtNvfp4Raw`), `src/vllm/model_executor/models/qwen3_5_weights.cpp` (35B `LoadBf16Direct`); upload counter `include/vllm/model_executor/models/dense_attn_block.h`; fp4 resident upload counter + post-upload residency `include/vllm/model_executor/models/dense_nvfp4_gemm.h` (`ResidentNvfp4`) and `src/vllm/model_executor/models/qwen3_5.cpp` (private `ResidentNvfp4`); shared checker text normalization `scripts/checker_text.py` | [mechanism gate](../tests/vllm/test_load_direct_upload.cpp#L1) 14/14, 183 assertions (178 through round 4) — the 6 borrow-mechanism cases, 4 post-upload residency cases that pin the adopt branch, its RELEASE-BEFORE-REASSIGN ordering and the `VT_ADOPT_DEVICE_BYTES` decoupling, 3 fp4-resident cases that pin `ResidentNvfp4`'s upload accounting, its `d_dev` publication, its adoption and the previously unexercised NON-borrowed/host-addressable regime, and 1 case that pins the GENERAL adopt branch's `MADV_DONTNEED` by OBSERVING RESIDENCY (`mincore()` over the host mirror's interior pages, glibc pinned to the sbrk arena so `free()` cannot return them by itself) — the RSS half of this row, which every value assertion was blind to. RED under eleven mutations, each applied alone with the binary rebuilt (a failed build aborts rather than re-running a stale one), tree md5-verified restored and re-GREEN after each. Mutations are named by the TEST CASE they land on, never by a line ref, because carried-forward line refs are how the counts below were twice recorded wrong: drop the size-identity check 1 case / 5 assertions; skip the borrow in `LoadBf16Direct` 2 / 7; delete the adopt branch 5 / 19; move the release after the `bytes` reassignment 5 / 7 (incl. the ordering assertions read from inside the munmap); move the release after the host-addressable early return but BEFORE the env one 2 / 3; move it after the `VT_ADOPT_DEVICE_BYTES` early return, i.e. after BOTH, 3 / 4; delete the general branch's `MADV_DONTNEED` 1 / 1; `ResidentNvfp4` drop both `AddDeviceUpload` 3 / 3; drop both `d_dev` publications 3 / 20; drop both `AdoptDeviceBytesAsHost` 3 / 16; all six at once 3 / 23. FOUR of those rows had been recorded wrong (1/1, 2/2, 3/7 and 3/3) — every one a count measured against an older suite and copied forward rather than re-measured; all eleven are measured on this tree. The `qwen3_5.cpp` duplicate of `ResidentNvfp4` is in an anonymous namespace no test can call, so it is held to the same invariant by [`scripts/check-fp4-resident-consistency.py`](../scripts/check-fp4-resident-consistency.py) (42-case mutation suite), which checks PER BUFFER inside each buffer's own `if (!w.d_<buf>)` upload block — body-wide matching let one surviving `AddDeviceUpload` satisfy both buffers, so dropping exactly one counter passed (reproduced against the real `qwen3_5.cpp` text: old checker exit 0, new exit 1). Nine drop-one/substitute-one mutations of the LIVE duplicate now go RED. Round 5 closed two holes in it. It matched RAW source, so a statement left behind as a COMMENT, inside `#if 0`, or inside `if (false)` read as present while the compiler saw a deletion; the clause matchers now run on [`scripts/checker_text.py`](../scripts/checker_text.py)'s `normalize_source`, which blanks all three IN PLACE so byte offsets and line numbers survive — `strip_comments` had existed as two byte-identical private copies (check-runner-routing-consistency.py, check-surface-coverage.py) and both now import the shared helper instead of a third copy being written. And the ordering clause required PUBLISH-before-ADOPT but not COPY-before-ADOPT, so a body whose adoption runs BEFORE its copy passed although the upload then reads pages the adoption already released; new clause (f) READ-FIRST covers it. WHICH MOTION produces that order is part of the claim, because the two are different mutations: SINKING `d.b.Copy(...)` to below `AdoptDeviceBytesAsHost` leaves the `d_dev` publication in place so ONLY clause (f) bites (MEASURED on the live duplicate, old exit 0 / new exit 1 — the 0/1 row below), whereas HOISTING the adoption above the copy also lifts it above the publication and trips the PRE-EXISTING clause (e), which the old checker already caught (MEASURED 1/1) and which is therefore no evidence for (f). The equivalent mutations of the SHARED `dense_nvfp4_gemm.h` copy are red at run time, MEASURED one at a time with the binary DELETED before each rebuild and the header md5-verified (`4665255f7af6f52254367f9118ad92ee`) before and after each: sink `packed`'s `Copy` 2 cases / 2 assertions, sink `scale`'s 2 / 2, sink BOTH 2 / 4 — the failures are `w.<buf>.bytes.data()[0] == kSrcPattern` and `AllBytesMatchPattern(...)` in `fp4 resident: ResidentNvfp4 COUNTS its upload…` and `fp4 resident: a host-addressable device adopts an OWNED fp4 mirror too`, TWO per mutated buffer, so an ODD count is not obtainable; hoist `packed`'s adoption 3 / 8 and hoist both 3 / 16, wider because a null `d_dev` makes the adoption return immediately and also takes the NON-host-addressable case. This clause had been recorded as **2 cases / 3 assertions**, which is the count of the *release after the host-addressable early return, BEFORE the env one* row carried onto a different mutation — the same failure mode as the four rows above, found by a fifth reviewer and re-measured here rather than copied. MEASURED on the LIVE `qwen3_5.cpp` on disk, one mutation at a time, md5-verified restored after each (35b5ea250f490105d579f4ffb573aa36 before and after all eleven) — old checker exit / new checker exit: delete the packed adoption 1/1, `//` it out 0/1, `/* */` it out 0/1, `#if 0` around it 0/1, `if (false)` around it 0/1, `//` the packed upload counter 0/1, `//` the packed `d_dev` publication 0/1, `#if 0` around that publication 0/1, `if (false)` around it 0/1, SINK the packed `Copy` to below its adoption 0/1. That gate is STRUCTURAL: it guards the duplicate against DELETION — including deletion disguised as a comment, an `#if 0` or a never-taken branch — against gross substitution and against mis-ordering, NOT against corruption, and it does not model the preprocessor (extending it to arbitrary `#ifdef` conditions is DECLINED: `#ifdef VT_CUTLASS_NVFP4` around the live adoption stays exit 0 in both checkers, because a build configuration is not a disguised deletion). Run-time proof exists only for the shared copy. The `mincore()` residency case now also carries a RUNTIME allocator guard — `#if __GLIBC__` proves the headers, not that glibc's allocator is running — which goes red (1 case / 1 assertion) when a purging allocator is simulated; the `mallopt` calls leave glibc's `no_dyn_threshold` set for the process either way, MEASURED as not specific to `M_TRIM_THRESHOLD` (glibc 2.39, `mallinfo2().hblks` probe: no mallopt LIVE, `M_TRIM_THRESHOLD` DISABLED, `M_MMAP_MAX` DISABLED), so it is recorded rather than removed. Round-4 re-run on the dev box, CLEAN Vulkan/llvmpipe Release build: `test_load_direct_upload` 14/14 (178), `test_safetensors` 34/34 (79), `test_qwen36_weights` 7/7 (45), `test_vulkan_backend` 35/35 (2107), `test_backend_cross_device` 11/11 (132), `test_opt_paged_engine` **6/6 token-exact (96/96), 0 declines, device type 3**. Round-5 re-run on the same box, branch rebased onto `a0fa12c7`, all identical except `test_load_direct_upload` **14/14 (183)** — the five added assertions are the RSS-reclaim case's runtime allocator guard, no case-count change — plus `test_checker_text` 33/33 and `scripts/gen-vulkan-spirv.py --check` `committed SPIR-V is up to date` run with the pinned `~/tools/glslang-16.5.0/bin/glslang` (without it the script exits 1 on this tree, so the green is a real compile-and-compare and not a skip). Round 6 changed NO compiled file — this row, the spec, the `(f)` clause docstring and one checker-test name — and re-ran the whole battery on the branch rebased onto `e3cc4f64`, CLEAN Release Vulkan build, 0 warnings: every count above reproduces byte-for-byte, `check-fp4-resident-consistency` rc 0, its suite 42, `test_checker_text` 33, `test_check_surface_coverage` 46, `test_check_runner_routing_consistency` 31, SPIR-V check up to date, preflight all green. GB10 Vulkan gates on the changed tree: `test_vulkan_backend` 35/35 (2650), `test_backend_cross_device` 11/11 (132), `test_opt_paged_engine` **6/6 prompts token-exact (96/96), 0 declines, device type 3**. GB10 CUDA (cutlass+triton) full `ctest` **383/393**, BOTH SACRED gates PASS (`test_qwen36_paged_engine`, `test_qwen27_paged_engine`); all 10 failures reproduce with the same signature on a clean `origin/main` CUDA build. MEASURED 27B bf16 load **1.54x warm / 1.61x cold**, bytes moved **100.196 -> 81.260 GiB** | [specs/load-direct-upload.md](specs/load-direct-upload.md) | `ACTIVE` | `CLAIM-ENG-LOAD-DIRECT-UPLOAD` |
 | `LOAD-SAFETENSORS-DIRECT-DENSE` | Layer-bounded target-device loading for ordinary plain-BF16 Qwen3.5 dense safetensors on discrete CUDA; additive leaf below `LOAD-SAFETENSORS`, preserving windowed source release and owned-shard lifetime. Plain weights, stacked raw-NK owners, tied logits, logical resident state and same-queue layer staging are implemented and locally CUDA-gated. H32 Triton AOT, plain-BF16 decode graphs and ratio-4 FA2 repair the transplanted hot path; the broader row remains speed-gating | T1 | target-device model construction `vllm/model_executor/model_loader/base_loader.py:43-82`; incremental safetensors yield `vllm/model_executor/model_loader/weight_utils.py:820-954`; stacked/tied parameters `vllm/model_executor/models/qwen3_5.py:276-303,483-492` | load/dispatch `src/vllm/model_executor/models/qwen3_5_dense_weights.cpp:52-133,187-246,334-472`; discrete/unified classifier `src/vt/cuda/cuda_backend.cu:251-265`; plain execution/residency and graph selection `src/vllm/model_executor/models/qwen3_5.cpp`, `src/vllm/model_executor/models/qwen3_5_dense.cpp`; H32 recurrence `src/vt/cuda/cuda_gdn.cu`; ratio-4 FA2 `src/vt/cuda/cuda_flash_attn_fa2.cu`, `src/vt/cuda/cuda_paged_attn.cu`; exact benchmark corpus/output capture `examples/bench/bench_core.h:61-68,357-386,409-589`; reference metrics/token collector `tools/bench/vllm_closed_loop_metrics.py`; guarded driver and summarizer `tools/bench/run_qwen35_4b_compare.sh`, `tools/bench/summarize_qwen35_4b_compare.py` | [Real 4B graph/direct ON/OFF/eager gate](../tests/vllm/models/test_qwen35_plain_weights.cpp#L80) passes **3/3, 1672/1672**; [H32/GDN tests](../tests/vt/test_ops_gdn.cpp#L1) **10/10** flag and **66/66, 4242/4242** full; [paged-attention tests](../tests/vt/test_ops_paged_attn.cpp#L1354) **25/25, 454474/454474**. Final 18-leg root `/tmp/qwen35-main-final-fa2-20260725` plus stable vLLM confirmation: direct ON/OFF/vLLM-0.25 total **5769.99/5660.70/5849.80 tok/s**, output **638.03/625.94/646.85**, TPOT/ITL **43.72/43.84/38.55 ms**, peak PSS **2.406/8.592/7.662 GiB**, stable PSS **0.759/8.589/4.029 GiB**, VRAM **12850.7/12843.3/12942.7 MiB**. ON=OFF output IDs 128/128 every pair; ON is **+1.93%** total and cuts peak/stable PSS **72.0%/91.2%**. H32 AOT/graph/FA2 A/B gains are **+4.59%/+0.39%/+1.60%**. Graph-node trace has 453 launches and 200972 child kernels; local FA2 is 180.28 us/call vs vLLM 178.40, so the remaining **0.9864x** throughput and TPOT gap is host/engine-side. Sanitizer availability and external 27B/35B remain open | [plain-BF16 direct-load spike](specs/qwen35-plain-bf16-direct-load.md); [2026-07-25 evidence](../docs/bench-evidence/qwen35-4b-main-repair-20260725.md) | `GATING` | - |
-| `ENG-MOE-HOSTFREE` | MoE Marlin resident host-weight release: after `BuildMoeMarlinResident` uploads+repacks the routed experts to the device Marlin resident, free the per-expert fp4 HOST mirror (`OwnedTensor` packed+scale bytes) and `madvise(MADV_DONTNEED)` the pages back to the OS — returns the ~16.9 GiB steady 35B host double-store (`LoadNvfp4Raw` `MakeOwned` copies kept resident forever). Guarded to the committed Marlin path (`MarlinMoeEnabled()`; retained for the `VT_NVFP4_MARLIN=0` wmma fallback that re-reads them); `VT_MOE_HOST_FREE=0` A/B rollback. Realizes `release_host_weights_after_upload` for the dominant host consumer. **Item-2 (2026-07-19, `CLAIM-BACKEND-PLATFORM-2`): the host-free decision is now CONSUMED from `GetPlatform(d.q.device.type).residency_policy()` via `vllm::platforms::ShouldReleaseHostWeights(policy, MarlinMoeEnabled()/*kernel-path*/, VT_MOE_HOST_FREE/*env*/)` — `CudaPlatform` flag flipped false→true (reproduces today EXACTLY); `MarlinMoeEnabled()` stays the orthogonal KERNEL-PATH safety gate.** | T0 | streamed target-device weight construction `vllm/model_executor/model_loader/base_loader.py:43-82`; residency capability `vllm/platforms/interface.py:134-229`; Marlin MoE in-place repack (no host mirror) `vllm/model_executor/layers/quantization/utils/marlin_utils_fp4.py:375-434` | free region `src/vllm/model_executor/models/qwen3_5.cpp:3743-3781`; `OwnedTensor::ReleaseHost` decl `include/vllm/model_executor/models/qwen3_5_weights.h:65` + impl (madvise+swap) `src/vllm/model_executor/models/qwen3_5_weights.cpp:24`; public hook `Qwen3_5Model::PrepareMarlinResident` `src/vllm/model_executor/models/qwen3_5.cpp:9130` | `tests/vllm/test_qwen36_weights.cpp:283` (`ReleaseHost` frees buffer+capacity) + `:314` (`PrepareMarlinResident` release under Marlin / retention under `VT_NVFP4_MARLIN=0`, DGX release 27/27 + retention 15/15); DGX A/B (`VT_MOE_HOST_FREE`) 35B STEADY serving PSS **20.17→3.53 GiB** (root `dgx:~/work/mem35-hostfree`); token-neutral 315/315 + 235/235; c2 smoke clean; memcheck clean. Whole-window load-phase PEAK bounded by the `ENG-MOE-LOADSTREAM` follow-up. Ledger [parity-ledger.md#L521](parity-ledger.md#L521) | [moe-marlin-host-free.md](specs/moe-marlin-host-free.md) | `DONE` | `ac77bec` |
+| `ENG-MOE-HOSTFREE` | MoE Marlin resident host-weight release: after `BuildMoeMarlinResident` uploads+repacks the routed experts to the device Marlin resident, free the per-expert fp4 HOST mirror (`OwnedTensor` packed+scale bytes) and `madvise(MADV_DONTNEED)` the pages back to the OS — returns the ~16.9 GiB steady 35B host double-store (`LoadNvfp4Raw` `MakeOwned` copies kept resident forever). Guarded to the committed Marlin path (`MarlinMoeEnabled()`; retained for the `VT_NVFP4_MARLIN=0` wmma fallback that re-reads them); `VT_MOE_HOST_FREE=0` A/B rollback. Realizes `release_host_weights_after_upload` for the dominant host consumer. **Item-2 (2026-07-19, `CLAIM-BACKEND-PLATFORM-2`): the host-free decision is now CONSUMED from `GetPlatform(d.q.device.type).residency_policy()` via `vllm::platforms::ShouldReleaseHostWeights(policy, MarlinMoeEnabled()/*kernel-path*/, VT_MOE_HOST_FREE/*env*/)` — `CudaPlatform` flag flipped false→true (reproduces today EXACTLY); `MarlinMoeEnabled()` stays the orthogonal KERNEL-PATH safety gate.** | T0 | streamed target-device weight construction `vllm/model_executor/model_loader/base_loader.py:43-82`; residency capability `vllm/platforms/interface.py:134-229`; Marlin MoE in-place repack (no host mirror) `vllm/model_executor/layers/quantization/utils/marlin_utils_fp4.py:375-434` | free region `src/vllm/model_executor/models/qwen3_5.cpp:3743-3781`; `OwnedTensor::ReleaseHost` decl `include/vllm/model_executor/models/qwen3_5_weights.h:65` + impl (madvise+swap) `src/vllm/model_executor/models/qwen3_5_weights.cpp:24`; public hook `Qwen3_5Model::PrepareMarlinResident` `src/vllm/model_executor/models/qwen3_5.cpp:9132` | `tests/vllm/test_qwen36_weights.cpp:283` (`ReleaseHost` frees buffer+capacity) + `:314` (`PrepareMarlinResident` release under Marlin / retention under `VT_NVFP4_MARLIN=0`, DGX release 27/27 + retention 15/15); DGX A/B (`VT_MOE_HOST_FREE`) 35B STEADY serving PSS **20.17→3.53 GiB** (root `dgx:~/work/mem35-hostfree`); token-neutral 315/315 + 235/235; c2 smoke clean; memcheck clean. Whole-window load-phase PEAK bounded by the `ENG-MOE-LOADSTREAM` follow-up. Ledger [parity-ledger.md#L521](parity-ledger.md#L521) | [moe-marlin-host-free.md](specs/moe-marlin-host-free.md) | `DONE` | `ac77bec` |
 | `ENG-MOE-LOADSTREAM` | 35B load-phase PEAK-PSS interleave (follow-up to `ENG-MOE-HOSTFREE`): the steady free returns the routed-expert host mirror only AFTER the whole model loads, so whole-window `peak_pss`/`peak_rss` (~19.8 GiB) is still set by all N layers' ~256 experts host-coexisting at load. DEFER the routed-expert host copies (`LoadQwen3_5Moe` loads each layer WITHOUT experts + installs a per-layer `load_layer_experts` streaming closure that owns the mmap'd shards) and materialize ONE layer's experts inside `PrepareMarlinResident` immediately before that layer's device Marlin build + host free — so at most one layer's experts coexist on the host (peak ~ one layer, not all N). Device residents byte-identical (same source bytes, same per-layer build order). Non-CUDA/`VT_NVFP4_MARLIN=0`/no-Marlin build falls back to bulk host materialization (their forward reads the host bytes). 27B is a different loader (`LoadQwen3_5Dense`, true-W4A4) → unaffected; GGUF/synthetic/borrowed pass no shards owner → eager. **Item-2 (2026-07-19, `CLAIM-BACKEND-PLATFORM-2`): the per-layer interleave gate is now CONSUMED from `GetPlatform(queue.device.type).residency_policy()` via `vllm::platforms::ShouldInterleaveLoadStream(policy, MarlinMoeEnabled())` — reproduces the old `queue.device.type != kCUDA` OR `!MarlinMoeEnabled()` gate EXACTLY (unified/CPU retain-host ⇒ policy false ⇒ materialize-all fallback); the ~4 GiB load-peak win is preserved.** | T0 | streamed target-device construction `vllm/model_executor/model_loader/base_loader.py:43-82`; in-place Marlin repack, no host mirror `vllm/model_executor/layers/quantization/utils/marlin_utils_fp4.py:375-434`; residency capability `vllm/platforms/interface.py:134-229` | deferred field `include/vllm/model_executor/models/qwen3_5_weights.h:1147` (`Qwen3_5MoeWeights::load_layer_experts`); loader defer + closure `src/vllm/model_executor/models/qwen3_5_weights.cpp:331,399` (`LoadMoeExpertsInto`/`LoadQwen3_5Moe`); shared shards owner `include/vllm/model_executor/models/model_registry.h:60` + `src/vllm/model_executor/models/model_registry.cpp:219` (`ModelSource::FromSafetensorsOwned`) + `src/vllm/entrypoints/model_loader.cpp:2556` (`LoadedEngine::FromModelDir`); per-layer interleave + `MaterializeAllDeferredExperts` `src/vllm/model_executor/models/qwen3_5.cpp:4498,4508` (`PrepareMarlinResident`) | CPU coexistence-bound contract `tests/vllm/test_qwen36_weights.cpp:324` ("deferred routed-expert load: move-safe closure + bounded coexistence", peak==1 across the per-layer materialize→free loop, move-safe closure); clean `-Werror` CPU build 0 warn, full ctest (3 HTTP/engine-proc parallel-port flakes pass isolated), tools 164/164. **DGX PROVEN** (`~/work/vllm.cpp-mem35-loadstream` new vs `-parent` 7a1a6d6 eager, production flags CUTLASS sm120a+Marlin+FA2 sm_121a, one flock): 35B load-to-ready **peak RSS (VmHWM) 21.43 GiB → 4.19 GiB (−17.24 GiB / −80%, below vLLM 13.3 GiB)**; token BYTE-IDENTICAL both binaries — 35B `test_qwen36_paged_engine` 315/315 + 27B `test_qwen27_paged_engine` 235/235; **27B UNAFFECTED** (peak RSS 24.8 GiB, matches its baseline — dense loader, no deferral); `compute-sanitizer memcheck` on the deferred load path **0 errors / 315 assertions** (no use-after-free of the freed host bytes); weights unit 127 assertions (CPU coexistence peak==1 + DGX residency). `benchmark_binding=false` — the orchestrator re-grids the binding `peak_pss`/`peak_rss` axes to confirm the FAIL→PASS flip | [moe-expert-load-stream.md](specs/moe-expert-load-stream.md) | `ANCHOR-BACKFILL` | CLAIM-MEM35-LOADSTREAM |
 | `LOAD-GGUF` | GGUF reader, dequantization, Qwen name transforms, embedded vocabulary | T0 | Pinned vLLM has no GGUF loader: `vllm/model_executor/model_loader/__init__.py:31-65`; compatibility reference is llama.cpp | `src/vllm/model_executor/model_loader/gguf_reader.cpp:302`; `src/vllm/model_executor/model_loader/gguf_dequant.cpp:223`; `src/vllm/model_executor/models/qwen3_5_gguf_weights.cpp:432`; `src/vllm/entrypoints/model_loader.cpp:240` | `tests/vllm/test_gguf.cpp:53`; `tests/vllm/test_gguf_dequant.cpp:25`; `tests/vllm/test_gguf_qwen36_loader.cpp:153`; real 35B `tests/parity/test_qwen36_gguf_engine.cpp:145` | `planned: specs/gguf-loader.md` | `PARTIAL` | - |
 | `LOAD-GGUF-MMPROJ` | A SECOND, `clip`-architecture GGUF projector file beside the language file, and the Qwen3-VL vision tower loaded out of it. **W1 LANDED (#821):** `EngineParams::mmproj_path` / `vllm_model_params.mmproj_path` (C ABI v22) / the server `--mmproj` flag name the file, `clip_mmproj_gguf.h` reads its `clip.*` metadata and `v.*` / `mm.*` tensors into the SHARED `multimodal::Qwen3VLVisionWeights`, and `LoadedEngine::vision_tower()` holds the result. The two-tensor temporal patch embedding (`v.patch_embd.weight` + `v.patch_embd.weight.1`) is INTERLEAVED per channel into the `[out, C*T*p*p]` conv3d operand the tower reads, and a file carrying only the first half is refused by name — the MuseGlimmer condition enforced rather than assumed. MuseGlimmer's own `MuseGlimmerRefuseMmproj` now has a PRODUCTION caller (routed on `clip.projector_type == "muse-glimmer"`), which is a change to MuseGlimmer's behaviour and is stated as one. **NOT supported:** no forward consumes the loaded tower yet — there is no multimodal request path on the C ABI and no GGUF image/video driver — and the COMMITTED 334-name manifest with its CI accounting is owed by `QUANT-QWEN38-27B-GGUF-ARM`, because the live confirmation below is env-gated on a NAS file and CI reads only the synthetic fixture. Auto-discovery of a sibling `mmproj*.gguf` is deliberately out of scope: a directory holding two unrelated models must not silently fuse them. First consumer is `Qwen3.8-27B` ([#821](https://github.com/mudler/vllm.cpp/issues/821)), whose `mmproj-BF16.gguf` (334 tensors, `clip.projector_type = qwen3vl_merger`, header-verified 2026-08-18: data end == file size 931,146,432) ships BOTH halves; MuseGlimmer's lacks the second | T0 | Pinned vLLM has no GGUF loader at all (`555967922`, `model_loader/__init__.py:33-49`), so the compatibility reference is llama.cpp `b10451` = `10bf611e5` (`PROJECTOR_TYPE_QWEN3VL`; the previously recorded `tools/mtmd/clip-impl.h:330` was read at the SUPERSEDED local fork `237ad9b96` and is owed re-anchoring, [#1003](https://github.com/mudler/vllm.cpp/issues/1003)) | the flag `EngineParams::mmproj_path` ([model_loader.h:1](../include/vllm/entrypoints/model_loader.h#L1)), the reader [clip_mmproj_gguf.cpp:1](../src/vllm/model_executor/models/clip_mmproj_gguf.cpp#L1), the open + refuse + read site `src/vllm/entrypoints/model_loader.cpp::LoadedEngine::FromModelDir` (the `.gguf` branch, after the device-fit refusal and BEFORE the tokenizer), the reader `src/vllm/model_executor/models/clip_mmproj_gguf.cpp::LoadQwen3VLVisionFromClipMmproj` + `::ClipMmprojVisionConfig` + `::RefuseUnsupportedClipMmproj`, the holder `include/vllm/entrypoints/model_loader.h::vision_tower` on `LoadedEngine`, the C face `include/vllm.h` `vllm_model_params.mmproj_path` (ABI v22) wired in `src/capi/vllm_c.cpp`, the server flag `src/vllm/entrypoints/openai/server_main.cpp` `--mmproj`, and the now-reached refusal `src/vllm/model_executor/models/muse_glimmer_gguf_weights.cpp::MuseGlimmerRefuseMmproj` | [test_clip_mmproj_gguf.cpp:1](../tests/vllm/models/test_clip_mmproj_gguf.cpp#L1) (9 cases / 272 assertions hermetic: the `clip.*` config mapping, the DeepStack discovery, the per-position patch-embedding interleave, every block/merger slot, and four refusals; the ninth is the LIVE confirmation, which skips loudly unless `VLLM_CPP_QWEN38_27B_MMPROJ` names the real `mmproj-BF16.gguf` and adds 43 assertions when it does — 334 consumed names == 334 shipped with nothing unread, the `clip.*` geometry, and the join checked at all 1,769,472 positions against F32 bytes read straight from the mmap. Renaming a tensor in the reader AND the fixture together leaves the hermetic gate green at 9/9 and reds only the live case, which is what the live case buys) and [test_gguf_mmproj_reach.cpp:1](../tests/vllm/entrypoints/test_gguf_mmproj_reach.cpp#L1) (6 cases / 19 assertions hermetic, all through `LoadedEngine::FromModelDir`; the sixth is the LIVE confirmation that the ENGINE holds the tower, and it skips loudly unless BOTH `VLLM_CPP_QWEN38_27B_GGUF` and `VLLM_CPP_QWEN38_27B_MMPROJ` name the real files, so the 19 assertions belong to the five hermetic cases). The two are separate targets on purpose: deleting the `LoadQwen3VLVisionFromClipMmproj` call site in `model_loader.cpp` reds the SECOND and leaves the FIRST fully green, which is the measured difference between gating a class and gating a capability. The port target for the vision config mapping is `src/vllm/model_executor/models/minimax_h3_vision_gguf.cpp::MiniMaxH3EncoderVisionConfig`, which builds the same `Qwen3VLVisionConfig` from `visual.*` | [quantized arms of Qwen3.8-27B](specs/qwen38-27b-quant-arms.md) | `PARTIAL` | - |
diff --git a/.agents/issues/BACKEND-TENSTORRENT-QWEN35/ISSUE-LOCAL-01M3JXEFQKSZP23PP2HWY9G0VQ.md b/.agents/issues/BACKEND-TENSTORRENT-QWEN35/ISSUE-LOCAL-01M3JXEFQKSZP23PP2HWY9G0VQ.md
new file mode 100644
index 000000000..4fe339f5f
--- /dev/null
+++ b/.agents/issues/BACKEND-TENSTORRENT-QWEN35/ISSUE-LOCAL-01M3JXEFQKSZP23PP2HWY9G0VQ.md
@@ -0,0 +1,378 @@
+ID: ISSUE-LOCAL-01M3JXEFQKSZP23PP2HWY9G0VQ
+Title: 27B serve: TT_FATAL writes during trace capture from CaptureSafeReshape tiled reshape
+Row: BACKEND-TENSTORRENT-QWEN35
+State: OPEN
+Kind: bug
+GitHub: -
+Mirror: PENDING
+Availability: FULL
+Created: 2026-09-28
+Updated: 2026-09-28
+Closed: -
+
+## Problem
+
+At main 8b5435bb0, the Qwen3.8-27B-Q4_K_M served arm (2x128/32 c2, both VT_TT_KEEPQUANT_INT8DOT=0 and =1) crashes in Qwen3_5DenseDecodeGraph::Step during trace capture: TT_FATAL 'Writes are not supported during trace capture' (tt-metal fd_mesh_command_queue.cpp:826). Chain: EnsureDevice2D -> CaptureSafeReshape -> ttnn::reshape (tiled) -> ReshapeViewTiledProgramFactory::create_program_artifacts -> ttnn::to_device host write mid-capture. Logs: /tmp/int8dot-leg0.log, /tmp/int8dot-leg1.log (2026-09-28). Suspects: the capture-warmup redesign #3321 or the GDN state-binding fix #3327 changed the capture shape; or the pinned tt-metal (9161e8fdb27+4) tiled-reshape path now materializes at artifact creation. The 27B serve arm has not run since those landed; the 9B did. The 92/92 suite does not cover this shape.
+
+## Resolution
+
+- 2026-09-28 (worktree row/tt-27b-capture-write, base 40990825d): root cause is
+  `EnsureDevice2D`'s exact-rows/cols arm in
+  `src/vt/tenstorrent/tenstorrent_residency.cpp` running a DIFFERENT chain in
+  the capture pass than the eager pass. A decode op commits a rank-3 device
+  result under a flat 2D slot record (`CommitDeviceLogical2D`), so the next
+  2D `EnsureDevice2D` lands on the arm whose logical shape mismatches. The
+  eager pass warms `to_layout(ROW_MAJOR) -> reshape (free view) ->
+  to_layout(TILE)`; the `tt_capture_active()` branch instead ran a bare
+  `ttnn::reshape` on the TILED rank-3 shadow — a spec the eager pass never
+  warmed — so the first decode capture created
+  `ReshapeViewTiledProgramFactory::create_program_artifacts`
+  (tt-metal reshape_tiled_program_factory.cpp:275), whose unconditional
+  `to_device` of the page-mapping tensor is the mid-capture write that
+  fatals at fd_mesh_command_queue.cpp:826. Fix: run ONE chain in both
+  passes (the W4 doctrine), deleting the capture-active branch. Red:
+  new doctest `EnsureDevice2D rank-3 reshape is capture-safe` reproduces the
+  exact TT_FATAL on the old branch (/tmp/red-focused.log) and passes with the
+  fix (/tmp/green-focused.log). Suite: 93/93 (/tmp/green-suite.log). 27B
+  serve leg: /tmp/leg-27b-green.log.
+- 2026-09-28 UPDATE: the 27B serve leg still fatals after the arm fix — a
+  SECOND, distinct divergence. With VT_TT_TRACE_DEBUG=1 (/tmp/leg-27b-diag.log)
+  the capture pass hits EnsureDevice2D's same-numel arm with spec
+  {1,10240} -> {2,5120}, a reshape the eager pass never ran (its arm726 specs
+  were {96,128}->{2,6144} and {3072,128}->{64,6144}); the reshape_tiled program
+  is created mid-capture and the same TT_FATAL fires. This is the "slot state
+  differs between passes" class: the producer commits {1,10240} in the capture
+  step where the warmup step committed a different shape. Issue stays OPEN for
+  that second site; the arm-710 unification and its red/green doctest stand as
+  committed.
+- 2026-09-28 (worktree row/tt-27b-capture-write, ab7cdb359 + follow-up): site 2
+  ROOT CAUSE: `MemsetDeviceIfCapture`'s fresh-slot lane in
+  `src/vt/tenstorrent/tenstorrent_residency.cpp` was CAPTURE-ONLY for the
+  shadow install. Under capture, a fresh-slot `DBuf::Zero` (the 20480-B
+  residual) installed a `{1,10240}` bf16 TILE shadow; the eager pass only
+  primed the zero and returned false (host memset + `MarkHostWritten`), so
+  eager ended with the consumer-shaped shadow (`{2,5120}` from kRmsNorm's
+  staging) and capture ended with the memset-shaped one. The capture consumer
+  then hit EnsureDevice2D's same-numel arm with the never-warmed
+  `{1,10240} -> {2,5120}` reshape (bench trace `arm726 rows=2 cols=5120
+  dev=1x10240 cap=1`) and the ReshapeViewTiled program's `to_device` wrote
+  mid-trace. FIX (W4 doctrine, one install in both passes): the fresh-slot
+  lane now installs the `{1, cols}` persistent shadow in BOTH passes, so the
+  eager consumer runs — and warms — the same-numel reshape and capture
+  replays it as a program-cache hit. Eager installs are bounded to
+  scratch-scale memsets (bytes <= 64 KiB): an unbounded eager install
+  retained the 3-8 MB weights-load slots and OOMed DRAM (a 268 MB
+  `ttnn::where` then missed by ~7 MB); larger slots keep the pre-fix
+  host-fallback priming. RED: doctest `fresh-slot Memset installs the same
+  shadow in both passes` reproduces the exact bench fatal on the old code
+  (`arm726 rows=2 cols=5120 dev=1x10240 cap=1` -> TT_FATAL at
+  fd_mesh_command_queue.cpp:826, /tmp/red-site2.log); GREEN with the fix
+  (arm726 cap=0 warms, cap=1 cache-hit, /tmp/green-site2.log). SUITE 94/94
+  (/tmp/suite-final2.log). The case also exposed a suite-hygiene bug, fixed
+  here: the `kRopeNeox (small)` case leaked `VT_TT_HOST_FREE_DECODE=0` into
+  every later case; it now restores the ambient value.
+- 2026-09-28 DEVICE GATE (27B leg, /tmp/leg-27b-final3.log, c1:
+  /tmp/leg-27b-c1.log): BENCH_EXIT=1 — the site-2 fatal is gone (capture
+  passes the `{1,10240}->{2,5120}` reshape), but two FURTHER blockers, both
+  previously masked because the leg died at site 2 first, now surface in
+  order: (1) at --concurrency 2, `TryReshapeAndCacheDeviceDecode`
+  (tenstorrent_paged.cpp:643) declines `num_slots > 1` ("decode T=1 only for
+  now"), RAC falls to the host path and `EnsureHost(k)` readbacks mid-capture
+  (fd_mesh_command_queue.cpp:873, "Reads are not supported"); (2) at
+  --concurrency 1 the whole decode capture replays cleanly but
+  `end_trace_capture` OOMs: the trace buffer needs 3,153,969,152 B against
+  ~298 MB free (MeshTrace::populate_mesh_buffer). These are new owed sites
+  (multi-slot RAC device path; 27B decode-trace DRAM fit), not regressions
+  of this fix. Issue stays OPEN for them.
+-
+- 2026-09-28 (worktree row/tt-27b-capture-write) BLOCKER A FIXED — multi-slot
+  decode RAC is capture-safe. ROOT CAUSE: the decline `if (num_slots > 1)
+  return false` at tenstorrent_paged.cpp:642 is NOT new — it is byte-identical
+  since 79ff8f310 (2026-08-18, the #1105 whole-graph capture) and was already
+  present at the W3-era c2 bench commit ead93289b. What changed since W3 is
+  the STATE OF k/v AT THE FALLBACK: the capture-warmup redesign's device-
+  residency waves (28c154d78 "grouped-quant activations serve device-resident"
+  and successors; the suite-repair 12d5d75ee) left the rope K/V outputs as
+  device_current TILE shadows at RAC time under capture. In W3 the host
+  fallback's `EnsureHost(k)` (tenstorrent_paged.cpp:890) found host bytes and
+  did a pure host write; at HEAD the same fallback triggered the device
+  readback mid-capture — TT_FATAL "Reads are not supported during trace
+  capture" (fd_mesh_command_queue.cpp:873). The decline that was a benign
+  perf shortcut in August became a capture-fatal in September.
+  FIX (W4 doctrine, warm in eager what capture replays; no capture-active
+  branch): TryReshapeAndCacheDeviceDecode admits the batched decode
+  (num_slots == T, one token per user) and runs the PROVEN single-user
+  sequence once per user — slice this user's [nkv, d] rows from the rope
+  shadow, 1.0-multiply into a fresh native [1,1,nkv,d], ttnn::copy into that
+  user's own single-shard persistent input (K on worker core u, V on C+u),
+  then one paged_fused_update_cache per user against that user's own [1]
+  update_idx and [1, cols] page-table row; the fused op's
+  override_runtime_arguments re-patches the per-user addresses on the shared
+  cached program, the same mechanism the C=1 replay uses. WarmRacIdx
+  allocates and per-step-refreshes the per-user idx tensors outside capture
+  (the per-step update_idx copy is the C=1 lane's on-device plus_one not yet
+  extended here — recorded as owed perf debt). An earlier design that packed
+  all C users into ONE C-shard sharded input was tried and REJECTED on two
+  tt-metal facts measured in this change: ttnn::copy demands LOGICAL shape
+  equality while a C-shard dest needs padded height C*np (unrepresentable
+  from logical C*nkv, tt_metal::Shape carries no padding), and concat does
+  not preserve per-input padding into its output storage.
+  RED: new doctest `kTENSTORRENT batched decode RAC is capture-safe
+  (num_slots=2)` reproduces the exact leg fatal with the decline restored
+  (/tmp/red-multislot.log, TT_FATAL fd_mesh_command_queue.cpp:873); GREEN
+  with the fix: capture + replay complete and BOTH users' KV verified
+  token-exact (128/128 K, 128/128 V) in the paged-KV device shadow via the
+  new ReadPagedKvShadowForTest hook (/tmp/green-multislot.log). The leg also
+  needed no residency change — the fix is confined to tenstorrent_paged.cpp.
+- 2026-09-28 BLOCKER B ANALYSIS — 27B whole-graph decode trace does not fit;
+  recommendation is REGION-SCOPED capture. The numbers: (1) DEMAND: the
+  captured decode graph records 1,037 tt-metal op entries (the step-decompose
+  legE in-capture census, identical across all four captures —
+  docs/bench-evidence/tt-27b-step-decompose-20260926.md) and end_trace_capture
+  asks for one 3,153,969,152 B DRAM buffer (/tmp/leg-27b-c1.log) — 3.04 MB
+  PER RECORDED COMMAND. The MeshTrace buffer is the replay staging DRAM for
+  the whole recorded command stream; at 1,037 commands the per-command region
+  (launch descriptors, CB/semaphore state, runtime-arg and buffer-descriptor
+  records) dominates, NOT tensor staging (every in-region tensor is
+  persistent/preallocated by the warmup discipline; the 27B DRAM ledger shows
+  the slot census flat) and NOT program binaries (cached outside the trace).
+  (2) SUPPLY: 298,568,896 B free, largest block 280,928,384 B — and ~948
+  MiB of that pressure is the tt-metal#57970-class GDN scatter retention
+  (docs/bench-evidence/tt-ttm-retention-rootcause-20260925.md), i.e. at most
+  ~1.9 GiB recoverable for a 2-request run, taking free to ~2.2 GiB. Even a
+  FULL upstream retention fix leaves 3.15 GB > 2.2 GB: NO whole-graph scope
+  fits the 27B decode step at this command density, and the gap is
+  structural (3,037 more commands than a ~90-command budget of 280 MB).
+  (3) PRECEDENT: region-scoped capture FITS and SERVES — the GDN-region row
+  captured and replayed a multi-layer region inside this same step
+  (.agents/specs/tenstorrent-gdn-region-replay.md,
+  docs/bench-evidence/tt-gdn-region-replay-20260927.md), and the chunked E=1
+  arm already targets a 50 MiB trace region (tenstorrent_capture.cpp's
+  LastTraceBytes discipline). A region of ≤ ~90 commands fits today's free;
+  ≤ ~16 commands fits a 50 MiB region.
+  RECOMMENDATION: region-scoped capture for the 27B decode — capture the
+  recurring per-layer compute region(s) (the kernel-dense GDN/attention/GEMM
+  chain), replay per layer from a host loop, and keep the heterogeneous
+  preamble/RAC/sampling ops eager. Whole-graph stays the right shape for the
+  small models it already serves; forcing 27B into it is bounded by tt-metal's
+  per-command trace cost, which no vllm.cpp-side discipline can shrink. Do
+  not attempt this inside this issue — it is a fresh row (trace-budget,
+  region segmentation, per-region state binding) with its own spec.
+- 2026-09-28 DEVICE GATE (c2 leg, /tmp/leg-27b-c2b.log = monitor
+  bench-c2b): BENCH_EXIT=1 — BLOCKER A's fatal is GONE (the leg no longer
+  dies at RAC; it served thousands of decode steps across ~16 minutes) but
+  fails FURTHER DOWN at the NEXT site of the same class:
+  PagedAttentionKernel's host fallback read mid-capture (EnsureHost ->
+  DownloadToHost -> to_vector, TT_FATAL "Reads are not supported during
+  trace capture" at fd_mesh_command_queue.cpp:873) during a late re-capture
+  — TryPADecodeDevice's multi-slot condition declines in some late state
+  (width change / boundary re-capture) and the leg falls to the PA host
+  path. This is a NEW owed site (multi-slot PA device path under re-capture),
+  the third of the class after the two capture-write sites and the RAC one.
+  Also fixed en route: the batched user split handles the 27B's real rope
+  shadow geometry — rank-2 token-row [T, nkv*d] (per-head d-ALIGNED column
+  slices, the proven recipe) as well as rank-3 [C, nkv, d].
+  RESIDUAL (recorded, not hidden): the new doctest passes standalone and in
+  12-case subsets, but in the FULL 95-case suite it measures 96/128 K elems
+  exact — user 1's SECOND head (and only that head) reads uninitialized bytes
+  from the paged-KV shadow after the eager pass, deterministically, with the
+  same 32-elem count across eager/capture/replay; the per-user
+  mesh_command_queue().finish() sync did not clear it. The fused-update
+  input dumps verify the device inputs are correct (out_headmax exact for
+  all four per-user copies), so the divergence is inside tt-metal's
+  copy/program-cache interaction for the LAST per-user input in a deep
+  program-cache history. Owed: bisect which preceding case poisons the
+  variant, then either the tt-metal report or a per-user program isolation.
+  SUITE: 94/95 passed (the residual is the only failure; every pre-existing
+  case stays green).
+- 2026-09-28 Blocks A and B status: A = the decline is REMOVED and the
+  batched RAC device path serves (warm in eager, replay in capture —
+  /tmp/red-multislot.log reproduces the old fatal, /tmp/green-multislot.log
+  the green case); the serve leg now reaches the PA site above. B = the
+  trace-budget analysis above stands (region-scoped recommended). Issue
+  stays OPEN for the PA multi-slot site, the batched-lane residual, and the
+  27B decode-trace DRAM fit.
+- 2026-09-28 (worktree row/tt-27b-capture-write) PA MULTI-SLOT SITE FIXED —
+  the same doctrine as Blocker A, third site of the class. ROOT CAUSE
+  (diagnosed live, /tmp/leg-27b-diag.log = monitor bench-c2-diag, the leg
+  re-run with VT_TT_TRACE_DEBUG=1): the final traces before the fatal are
+  "PA q_from_device FAILED: tenstorrent PA: batched (B>1) Q 4D materializa-
+  tion is not capture-safe; the host Q path must serve this step" ->
+  "PA device decode FAILED: vt: tenstorrent: PA Q host path is not capture-
+  safe (from_vector readback)" (tenstorrent_paged.cpp:1280) -> the host PA
+  oracle's EnsureHost readback mid-trace (TT_FATAL fd_mesh_command_queue.cpp:873).
+  The decline was the explicit `if (tt_capture_active() && Bu > 1) throw` in
+  TryPagedAttentionDeviceDecode's identity-Q path — a stale W3-era guard.
+  Its premise ("that program calls to_device — forbidden during trace
+  capture") predates the W4 doctrine: the B>1 arm runs the IDENTICAL
+  multiply(reshape(...)) chain in both passes, so the eager step warms the
+  reshape program for the exact input/output spec and the captured call is a
+  program-cache HIT; a spec the warmup did not warm still fatals loudly at
+  the miss (the W4 divergence detector, not a defect to guard against).
+  FIX: the guard is deleted; both passes run the same chain (one hunk in
+  tenstorrent_paged.cpp, no capture-active branch).
+  RED: new doctest `kTENSTORRENT batched decode PagedAttention is capture-
+  safe (num_reqs=2)` (mirror of the RAC case) fails on HEAD for the right
+  reason (/tmp/red-pa.log): the capture pass takes the decline (trace in
+  /tmp/red-pa2.log shows the exact q_from_device FAILED chain) and the
+  captured+replayed output mismatches the device-path reference 2048/2048 —
+  the test stages generation-B K/V through RAC into the DEVICE paged-KV
+  shadow before the capture, so only a device-served PA can reproduce it.
+  GREEN: /tmp/green-pa.log — capture pass serves (q_from_device OK cap=1),
+  replay-vs-eagerB 0/2048 mismatched elems, both users nonzero.
+  EN ROUTE BUG (own issue ISSUE-LOCAL-01M3KM4R2KQN5WXTM57W8BD849): the new
+  case exposed that the batched RacIdxCache lane lacks the C=1 lane's
+  page-table width-change guard — fixed in the same change (evidence in that
+  issue). The RAC residual flake is NOT this: it still reproduces in the
+  full suite and stays owed.
+  SUITE: 95/96 (/tmp/suite-pa3.log) — every pre-existing case green, the PA
+  case green in-suite (0/2048), the only failure the recorded RAC residual
+  (126/128 K/V, user-1 second head).
+  DEVICE GATE: see the next dated entry (bench-c2c).
+- 2026-09-28 DEVICE GATE (c2 leg, /tmp/leg-27b-c2c.log = monitor bench-c2c,
+  post-fix HEAD): BENCH_EXIT=1, but the PA multi-slot site is GONE — the leg
+  served ~12+ minutes of batched decode through SIX successful boundary
+  re-captures (11:24:06, 11:26:40, 11:29:13, 11:31:46, 11:34:15, 11:36:40)
+  that previously died at the PA host readback. The fatal moved PAST the PA
+  site to the ALREADY-RECORDED Blocker B structural limit: the last
+  end_trace_capture asks for 3,128,655,872 B of DRAM trace staging (2.9 GiB,
+  ~3 MB/recorded command) against 278,858,624 B free / 266,655,872 B largest
+  block — bank_manager.cpp:495 OOM, same whole-graph-does-not-fit conclusion
+  as the Blocker B analysis above (3.15 GB demand, ~2.2 GiB recoverable best
+  case). No TPOT table: the leg died at a re-capture before completing the
+  32-token horizon. Blocker B stays with its fresh trace-budget row; this
+  issue's PA site is closed by the red/green + full-suite evidence above.
+- 2026-09-28 (worktree row/tt-27b-region-capture-spec, dc99071ad + f376b512c):
+  the Blocker-B region-scoped recommendation was implemented as the row's wave
+  1 and the fit wall was MEASURED. The dense decode driver captures ONE REGION
+  PER LAYER under VLLM_CPP_REGION_CAPTURE=1 through the bare GraphBreak seam;
+  the red-first device gate proves a TWO-region capture replays byte-identical
+  to eager across the boundary (region 0 = 2,048 B, region 1 = 3,088,384 B).
+  The 27B c1 leg then died at the 8th segment close: tt-metal mesh_trace.cpp:125,
+  trace buffer 4,226,469,888 B vs allocation high-water 4,229,506,816 B — the
+  64 live regions SUM to the whole graph's ~3.15 GB staging (each region owns
+  its trace staging until release, and all 64 replay every step), so
+  segmentation does not shrink the fit demand. The whole-graph OOM and this
+  collision are one structural limit. Evidence:
+  docs/bench-evidence/tt-region-capture-20260928.md. The 27B decode-trace
+  DRAM-fit site stays OPEN, now with the segmentation result recorded: it
+  closes only on a tt-metal-side change (shared/reusable trace staging, a
+  trace_region_size policy, or a #57970 retention recovery that actually
+  covers 3.15 GB) — escalated with these numbers beside tt-metal#57970. The
+  in-flow RAC C=1 segfault the legs exposed is ISSUE-LOCAL-01M3M0K390EM40W5R9BR5A2KZ7.
+
+- 2026-09-28 UPLOAD-GUARD WAVE (worktree row/tt-27b-region-capture-spec,
+  commits 6473ae731/286947603/39e2ca8ef, evidence
+  docs/bench-evidence/tt-capture-upload-guard-20260928.md): the
+  trace-record-audit inline-upload lever landed and the c1 leg FALSIFIED the
+  attribution. The guard (UploadRows/UploadRowsBf16 refuse capture-scope
+  uploads by name; the AddKernel broadcast operand warms into a hash-keyed
+  cache) is red-first proven (pre-fix: raw TT_FATAL
+  fd_mesh_command_queue.cpp:826, no refusal; post-fix: named refusal, warmed
+  capture 1,024 B). SUITE: 98/526,778 assertions, only the pre-recorded RAC
+  residual flake fails. MONEY LEG: 27B whole-graph c1 BENCH_EXIT=1, no TPOT —
+  but zero capture-scope uploads fired (no [TT-UP] line in the leg) and the
+  end_trace_capture demand is byte-identical 3,153,969,152 B; the warmed
+  MatmulBT region still closes at 3,088,384 B. Model C (inline H2D payload) is
+  falsified: the ~3 MB/command is the quant-matmul program class's own
+  recorded launch stream on this tt-metal pin. This issue's fit site stays
+  OPEN with a sharper next step: dump the tt::LogDispatch command stream for
+  one captured MatmulBTQuantGrouped launch and attribute the ~3 MB of
+  bypass_data (candidates: per-launch relay of kernel-binary pages for
+  programs missing the 1,024 KB prefetch ringbuffer fit at
+  fd_mesh_command_queue.cpp:453, or per-launch config/RTA page writes scaling
+  with the quant program's footprint). The escalation beside tt-metal#57970
+  now carries a reproducing op-scale case: a single warmed MatmulBT region
+  closes at 3,088,384 B on this pin.
+- 2026-09-28 ATTRIBUTION (worktree row/tt-27b-region-capture-spec, evidence
+  docs/bench-evidence/tt-launch-record-attribution-20260928.md): the
+  binary-relay candidate is FALSIFIED and the ~3 MB/command is attributed to
+  tt-metal's per-core launch-record path. Source (pin d20b8e27f29): the
+  1,024 KB prefetch-ringbuffer fit (fd_mesh_command_queue.cpp:453,
+  dispatch_settings.cpp:72) only chooses relay_paged vs relay_ringbuffer —
+  BOTH record the kernel binary BY REFERENCE to the resident DRAM kernels
+  buffer (dispatch.cpp:1942-1971, 2079-2116); binary bytes never enter
+  bypass_data; and load_binaries refuses first-time loads mid-capture by name
+  (mesh_workload.cpp:201-205). Our keep-quant kernel measures .text 51,648 B +
+  .data 3,372 B — 18x under the threshold. Device (region-handoff doctest under
+  gdb breakpins, capture window gated on tt_capture_active()): the 3,088,384 B
+  region record is exactly 284 issue_queue_reserve chunks (~10.9 KB each) of
+  the recorded command stream for the ONE full-grid MatmulBTQuantGrouped
+  program, with ZERO in-capture buffer-data writes (write_to_device_buffer: 0
+  real hits) — so no inline H2D payload, ours or tt-metal's. The record scales
+  with the keepquant program's per-core config/RTA dispatch footprint
+  (RmsNorm-class program: 2,048 B total), i.e. tt-metal-side. Our-side grid
+  shrink only scales the record linearly (halving the grid halves 3.15 GB —
+  still OOM) and is recorded as a bound, not a fix. OPEN NEXT (one step): a
+  logging-enabled pin build (TT_METAL_ENABLE_LOGGING=ON + TT_METAL_LOGGER_
+  LEVEL=TRACE, names verified at tt-logger.hpp:98,182 — current release builds
+  compile LogDispatch out, which is why the cheap logger leg was silent) to
+  name the dominant per-chunk class; the attribution does not depend on it.
+- 2026-09-29 (pin advanced to upstream-live 98134127a7b, pin head 6449cf13f7b,
+  logging build dir build_logging): the logger discriminator CLOSED the open
+  next step, and it flips the locus to OURS. On the new pin the focused leg
+  reads region 0 = 2,048 B; region 1 = 2,965,504 B (/tmp/tregion-newpin.log,
+  1/1 case, 1,032/1,032 assertions) — the record barely moved, so upstream's
+  576+ commits did not touch the per-core record path (create_trace_node /
+  issue_queue_reserve unchanged in dispatch.cpp). The capture window is
+  exactly 254 one-shot command-sequence fetches summing 2,961,024 B and
+  contains 11,040 per-core Unique RTA (UNICAST) writes (40-48 B payload each,
+  page-granular when recorded) plus 110 full-grid CB/DFB config pages. Those
+  per-core RTAs are OUR keepquant program's SetRuntimeArgs stream
+  (tenstorrent_keepquant.cpp:2108-2130): 12 words per core, 10 of them
+  shape-global constants and only r0 = c*tcols / rc = clamp(...) varying —
+  both derivable in-kernel from the core coordinate. CONCRETE FIX (ours):
+  compute r0/rc in the kernel, launch with SetCommonRuntimeArgs only, delete
+  the per-core SetRuntimeArgs stream. Expected record: ~2.97 MB -> the
+  RmsNorm-class floor (~2-16 KB per captured launch, ~200x), which closes the
+  27B whole-graph 3.15 GB trace demand. The earlier "tt-metal-side" locus is
+  thereby refined: tt-metal faithfully records what our program asks it to
+  dispatch per core; the shrink lever is ours.
+- 2026-09-29 (worktree row/tt-27b-region-capture-spec, fix commits a1661114b +
+  the RTA-removal commit): the CONCRETE FIX above LANDED and was measured.
+  The keepquant program now launches SetCommonRuntimeArgs-only (14 words) and
+  derives r0/rc in-kernel from the core coordinate
+  (tenstorrent_keepquant.cpp, CARG_* table; host per-core SetRuntimeArgs loop
+  deleted). Correctness held: E=1 grouped keep-quant capture-x2 byte-identity
+  PASS with a partial last core, host derivation-parity doctest added. 27B c1
+  arbiter (fresh build2, pin 6449cf13f7b): trace demand
+  3,153,969,152 -> 2,925,109,248 B (-228,859,904 B = 1,037 launches x ~920
+  cores x one 256 B recorded RTA page) — the lever is real and SPENT, but
+  BENCH_EXIT=1: the whole-graph trace STILL does not fit; c1 does not serve.
+  Two small-shape A/Bs (region-handoff 2,965,504 B both binaries; grouped
+  capture 48,316,416 B both binaries, dispatch-log Unique-RTA counts
+  identical) show the ~2.9 MB per captured command at those shapes was never
+  the RTA stream: the dominant remaining class is per-launch full-grid
+  program command-sequence payload (CB/DFB config pages per sequence). The
+  KB-floor gate for this vehicle was removed as unreachable by this lever
+  (records: docs/bench-evidence/tt-keepquant-rta-fix-20260929.md §3-§5). The
+  fit wall stands at 2,925,109,248 B; next lever is the per-launch
+  command-sequence payload class (issue stays OPEN, evidence above updated).
+- 2026-09-29 (worktree row/tt-27b-region-capture-spec, this leg): the
+  per-launch "config-page payload" hypothesis is REFUTED, and the record's
+  locus is now program COUNT in the grouped decode chain. The §4 synthetic
+  bisect (docs/bench-evidence/tt-trace-config-page-repro-20260929.md §5,
+  repro_bisect_program_shape.cpp) built the keepquant program's exact shape
+  raw (full-grid CoreRange DM kernel, common RTAs) and swept CB count
+  (4/8), CB page size (4/16 KiB) and kernel binary size (32-256 KiB .rodata
+  tables): every variant records 1,024 B/launch — the packed relay collapses
+  all of them, so non-identical per-core config pages are not the 2.82 MB.
+  The region-handoff doctest on HEAD (now gated at 64 KiB, RED measured
+  2,965,504 B, /tmp/region-red.log) showed the default dispatch there is the
+  W4a GROUPED arm: the region is ceil(N/8)=8 chunks x (~85 eltwise decode
+  programs from DecodeKeepQuantWordsF32 Q6_K + ~8 matmul-chain programs) ≈
+  680 programs x the 4-17 KB per-program floor = 2.97 MB. NEXT LEVER (one
+  step): collapse the decode to ONE custom-kernel program per launch — the
+  bit-exact int8-dot kernel is the existence proof at the floor — and the
+  64 KiB doctest gate is the arbiter. The 27B fit wall keeps its
+  2,925,109,248 B bound; c1/c2 re-measure legs stay blocked behind the
+  fusion. Issue stays OPEN.
+- 2026-09-29 (same leg, suite): the 64 KiB region gate landed RED by design
+  (commit 7a0f1ca3f). Full TT suite (build2, pin 6449cf13f7b): 99 cases,
+  97 passed, 2 failed — the owed 2261 flake (126 == 128) and the new
+  intentional red gate (11527, 2,965,504 > 65,536); 527,817/527,819
+  assertions; all keepquant bit-exact capture-x2 cases green
+  (/tmp/suite-final.log, teardown segfault after the run is pre-existing).
+  27B money legs NOT run: no fix landed this leg, so c1 would reproduce the
+  recorded 2,925,109,248 B / BENCH_EXIT=1 outcome; the legs stay blocked
+  behind the decode-fusion lever.
diff --git a/.agents/issues/BACKEND-TENSTORRENT-QWEN35/ISSUE-LOCAL-01M3KM4R2KQN5WXTM57W8BD849.md b/.agents/issues/BACKEND-TENSTORRENT-QWEN35/ISSUE-LOCAL-01M3KM4R2KQN5WXTM57W8BD849.md
new file mode 100644
index 000000000..e63071313
--- /dev/null
+++ b/.agents/issues/BACKEND-TENSTORRENT-QWEN35/ISSUE-LOCAL-01M3KM4R2KQN5WXTM57W8BD849.md
@@ -0,0 +1,35 @@
+ID: ISSUE-LOCAL-01M3KM4R2KQN5WXTM57W8BD849
+Title: RacIdxCache batched lane mishandles a page-table width change
+Row: BACKEND-TENSTORRENT-QWEN35
+State: OPEN
+Kind: bug
+GitHub: -
+Mirror: PENDING
+Availability: FULL
+Created: 2026-09-28
+Updated: 2026-09-28
+Closed: -
+
+## Problem
+
+WarmRacIdx keys RacIdxCache by (num_slots, block_size) but not page-table width. The C=1 lane reallocates on block_table_cols != e.pt_width (retire + realloc, the #1105 discipline); the batched (num_slots>1) lane added in e39f2cf3f has no such guard: its refresh branch indexes batched_pt_host with the CALLER's block_table_cols against a vector sized at allocation width (OOB read) and copy_to_device's a [1, new_cols] host tensor into a [1, old_cols] device tensor — TT_FATAL 'Host tensor has different shape' (tensor_apis.cpp:161). Exposed by the new batched-PA capture doctest (cols=2) running after the batched-RAC doctest (cols=1) in the full suite.
+
+## Resolution
+
+- 2026-09-28 (worktree row/tt-27b-capture-write) FIXED. The batched lane now
+  mirrors the C=1 lane's `pt_width` discipline: any `block_table_cols !=
+  e.batched_pt_width` on a live entry retires the per-user page tables into
+  `batched_retired_pts` (kept alive — never free a buffer a recorded trace
+  addresses, #1105), reallocates them at the new width, and resets
+  `batched_pt_host`; `batched_pt_width` records the allocation width.
+  `batched_update_idxs` ([1] per user) and the sharded inputs are
+  width-independent and untouched. Evidence: before the fix the new
+  batched-PA capture doctest (page-table cols=2) threw TT_FATAL
+  "Host tensor has different shape" (tensor_apis.cpp:161) when it ran after
+  the batched-RAC doctest (cols=1) in the full suite — both share RacIdxCache
+  key (num_slots=2, block_size=32); /tmp/suite-pa.log. After the fix the full
+  96-case suite is 95/96 with the only failure the pre-existing owed RAC
+  residual (/tmp/suite-pa3.log: 126/128, user-1 second head — the signature
+  recorded at e39f2cf3f, unchanged); the PA case reads replay-vs-eagerB
+  0/2048 mismatched elems in the same full-suite run.
+-
diff --git a/.agents/issues/BACKEND-TENSTORRENT-QWEN35/ISSUE-LOCAL-01M3M0K390EM40W5R9BR5A2KZ7.md b/.agents/issues/BACKEND-TENSTORRENT-QWEN35/ISSUE-LOCAL-01M3M0K390EM40W5R9BR5A2KZ7.md
new file mode 100644
index 000000000..ed683ce09
--- /dev/null
+++ b/.agents/issues/BACKEND-TENSTORRENT-QWEN35/ISSUE-LOCAL-01M3M0K390EM40W5R9BR5A2KZ7.md
@@ -0,0 +1,19 @@
+ID: ISSUE-LOCAL-01M3M0K390EM40W5R9BR5A2KZ7
+Title: 27B c1 decode: RAC C=1 lane routes through unallocated batched tensors — segfault at the first cold decode step
+Row: BACKEND-TENSTORRENT-QWEN35
+State: CLOSED
+Kind: bug
+GitHub: -
+Mirror: PENDING
+Availability: FULL
+Created: 2026-09-28
+Updated: 2026-09-28
+Closed: 2026-09-28
+
+## Problem
+
+At row/tt-27b-region-capture HEAD ec4e8a824, the Qwen3.8-27B-Q4_K_M c1 leg (--concurrency 1) segfaults in ttnn::copy inside ReshapeAndCacheKernel during the COLD eager decode step (capturing=0; /tmp/leg-control-c1.log, /tmp/leg-region-c1-diag.log, 2026-09-28, thalia). Control leg without VLLM_CPP_REGION_CAPTURE crashes identically, so this is pre-existing on the base, not the region arm. Root cause: e39f2cf3f rewrote TryReshapeAndCacheDeviceDecode as one per-user batched loop (rac_entry.batched_in[u], batched_update_idxs[u], batched_page_table[u]) but WarmRacIdx allocates those ONLY for num_slots>1 — for C=1 it allocates the shared sharded_in/sharded_in_v/update_idxs/page_table and its warm gate admits C=1 on `allocated` alone, so the loop indexes empty vectors (empty ttnn::Tensor -> null storage -> ttnn::Tensor::memory_config() segfault). The commit's claim 'The C=1 lane is untouched' is false; no c1 leg ran on this branch since e39f2cf3f (the doctrine legs were c2). Fix: restore the proven C=1 sequence verbatim (build_input over the whole shadow into the shared sharded tensors, one paged_fused_update_cache against the shared update_idxs/page_table) beside the batched loop.
+
+## Resolution
+
+2026-09-28: fixed in the same flow that found it (commit restoring the C=1 lane verbatim beside the batched loop). Red /tmp/leg-control-c1.log + /tmp/leg-region-c1-diag.log (cold-step segfault, capturing=0, with and without VLLM_CPP_REGION_CAPTURE); green /tmp/leg-region-c1-fix.log (cold step + capture pass run, leg proceeds to the unrelated fit wall) and the full TT suite green on the RAC lane apart from the recorded OWED batched flake. Full detail in the issue's Resolution section.
diff --git a/.agents/specs/tt-27b-region-capture.md b/.agents/specs/tt-27b-region-capture.md
new file mode 100644
index 000000000..eeec793a9
--- /dev/null
+++ b/.agents/specs/tt-27b-region-capture.md
@@ -0,0 +1,307 @@
+# Spec: region-scoped decode capture for the Tenstorrent 27B — the whole graph does not fit, so capture what recurs
+
+Row: `tt-27b-region-capture`. State: DRAFT (2026-09-28).
+Issue: `ISSUE-LOCAL-01M3JXEFQKSZP23PP2HWY9G0VQ` (Blocker B owns the
+trace-budget analysis this row implements; the row closes the 27B
+decode-trace DRAM-fit site of that issue).
+Builds on: the capture-safety doctrine landed on
+`row/tt-27b-capture-write` — EnsureDevice2D one-chain (`ab7cdb359`),
+the fresh-slot Memset shadow in both passes bounded to ≤ 64 KiB
+(`965822766`), the batched decode RAC per-user device path
+(`e39f2cf3f`), and the PA stale-guard deletion + page-table width
+guard (`809067ccb`). The OWED RAC flake from `e39f2cf3f` (the doctest
+order-sensitive under full-suite program-cache history) carries forward,
+not into this row's scope.
+Precedent: `.agents/specs/tenstorrent-gdn-region-replay.md` — region-
+scoped capture already serves on the 9B (served == eager byte-identical,
+8 replays > 4 captures, the in-place state commit at
+`src/vt/tenstorrent/tenstorrent_gdn.cpp:933` and `:2441`).
+Git integration: ONE pull request (spec + implementation + gates),
+developer decision recorded 2026-09-28 in
+`.agents/developer-preferences.md` under `## Git integration`.
+
+## Scope
+
+The Qwen3.8-27B dense GDN-hybrid decode step on Tenstorrent (the c2 leg,
+Q4_K_M anchor arm) cannot use whole-graph trace capture: the captured
+graph records 1,037 tt-metal commands and `end_trace_capture` asks for
+one 3,153,969,152 B DRAM staging buffer (3.04 MB per recorded command)
+against 298,568,896 B free; even a full tt-metal#57970 retention
+recovery (~948 MiB back, to ~2.2 GiB) leaves 3.15 GB > 2.2 GB. The
+demand is per-command launch RECORDS — descriptors, CB/semaphore state,
+runtime args — not tensor staging (every in-region tensor is persistent
+under the warmup discipline) and not program binaries (cached outside
+the trace). This row replaces whole-graph capture for the 27B with
+per-layer compute regions captured and replayed per layer from a host
+loop, keeping the heterogeneous preamble, RAC, and sampling eager.
+
+Whole-graph capture stays the shape for every model it already fits. The
+row ships a size/fit predicate (below), not a flip of the default.
+
+## Upstream anchors
+
+- vLLM capture semantics: CUDA-graph capture in vLLM is per-iteration,
+  whole-forward, and falls back to eager when capture does not fit
+  (`vllm/worker/gpu_model_runner.py` — the capture-size/eager fallback
+  path). vLLM never captures the sampler or the per-step host logic;
+  the captured region is the model forward. This row mirrors that
+  polarity one level finer: capture the per-layer forward compute, keep
+  the step plumbing eager. The eager fallback behavior (decline capture,
+  serve eager, name the missing part) is the vLLM behavior this row's
+  predicate inherits.
+- tt-metal `MeshTrace` constraints: the trace buffer is the replay
+  staging DRAM for the whole recorded command stream;
+  `MeshTrace::populate_mesh_buffer` / `bank_manager.cpp:495` OOM at
+  3,128,655,872 B requested vs 278,858,624 B free / 266,655,872 B
+  largest block (the c2c leg, `/tmp/leg-27b-c2c.log`, 2026-09-28), and
+  3,153,969,152 B vs 298,568,896 B free on the c1 leg. Mid-trace host
+  writes and readbacks are fatals (`fd_mesh_command_queue.cpp:826` and
+  `:873`), which is the constraint the landed doctrine already serves.
+- The per-command economics: ~30 ms/command trace execution
+  (`tt-27b-step-decompose-20260926.md`, the legE in-capture census,
+  1,037 entries identical across all four captures) and
+  `tt-capture-economics-20260927.md` (the ~31.1 s caller window is
+  98-99% of real-length TPOT).
+
+## Design
+
+### What is captured: per-layer compute regions
+
+The 27B decode step is a homogeneous inner loop: 64 layers of
+GDN/attention + GEMM + norms between one heterogeneous preamble and the
+RAC/sampling tail. Capture ONE region per layer (or per contiguous block
+of layers, if the boundary discipline below prices blocks cheaper): the
+region contains that layer's kernel-dense compute chain and nothing
+else. The eager preamble (embedding, first norms), the RAC KV write,
+the LM head, and the sampler stay OUTSIDE the trace — the same
+polarity vLLM captures with (model forward only, never the sampler),
+applied per layer.
+
+### Region budget math
+
+- The GDN precedent's fit number: the chunked E=1 arm targets a 50 MiB
+  trace region (`src/vt/tenstorrent/tenstorrent_capture.cpp:90`,
+  `LastTraceBytes()` discipline) and the 9B GDN region captured and
+  replayed inside it.
+- Today's free pool admits ~90 commands: 90 × 3.04 MB ≈ 274 MB against
+  298,568,896 B free.
+- The 50 MiB budget admits ~16 commands: 16 × 3.04 MB ≈ 48.6 MiB.
+- The 27B layer census: 1,037 commands / 64 layers ≈ 16.2 commands per
+  layer. **The derived budget: ≤ 50 MiB per region, i.e. ≤ ~16 commands
+  per captured region** — one layer per region lands on the 50 MiB
+  precedent's budget exactly. A layer block that stays ≤ ~90 commands
+  (~274 MB) is admissible only if the boundary-count risk (below)
+  prices it cheaper; the 50 MiB per-region cap is the spec's default,
+  and `LastTraceBytes()` asserts it at capture time (a region over cap
+  declines capture for that region and serves it eager, named).
+
+### Region boundary and state handoff discipline
+
+Every cross-region value (residual stream, GDN ssm/conv state slots, KV
+pages, rope outputs) lives in the persistent preallocated shadows the
+trace reads and writes IN PLACE — the proven W3/W4 in-place commit
+pattern (the 9B GDN region's `tenstorrent_gdn.cpp:933`/`:2441`
+discipline, extended fleet-wide by `965822766` and `e39f2cf3f`). No
+region boundary installs, frees, or re-shadows a state tensor: a fresh
+device tensor at a commit invalidates the address a captured region
+baked (the #3327-class defect — the region re-reads the freed block and
+re-scatters garbage over the live slot). Warmup runs every region
+eagerly first, so every captured call is a program-cache hit — the W4
+doctrine, one chain in both passes, no capture-active branch. The host
+loop replays region after region, re-patching per-region runtime args
+(the RAC per-user `override_runtime_arguments` mechanism) between
+regions; per-region `LastTraceBytes()` and the boundary counter are the
+fit and handoff instruments.
+
+### Whole-graph stays for small models: the size/fit predicate
+
+Capture scope is decided per model by a fit predicate evaluated before
+the first capture: estimate the whole-graph demand (the recorded command
+count × the measured per-command cost, both from `VT_TT_TRACE_DEBUG`'s
+census and `LastTraceBytes()`) against the free DRAM headroom; when it
+fits (the 9B and the small classic lane), capture whole-graph as today.
+When it does not, take the region-scoped path. The predicate uses
+measured numbers, not model size alone; a model whose whole graph fits
+never changes behavior. The 27B region-scoped arm is a named decline of
+whole-graph with a message that says why, per the refused-arm rule.
+
+## Risks
+
+- **Region count × per-region overhead vs the ~31.1 s window.** 64
+  regions × ~16 commands × ~30 ms/command ≈ 30.7 s/step of trace
+  execution — the same command stream the whole graph replays, so the
+  per-command term is unchanged; the NEW cost is the per-replay fixed
+  overhead (host loop dispatch, runtime-arg re-patch, boundary
+  bookkeeping). At a generous 5-10 ms per region boundary that is
+  0.32-0.64 s/step, 1-2% of the 31.1 s window — noise against the win:
+  the whole-graph arm SERVES NOTHING today (BENCH_EXIT=1 at the trace
+  OOM), so region-scoped's ~31 s/step is not a regression against 31 s,
+  it is the difference between a served 27B TT arm and none. The
+  remaining lever on the 31.1 s itself is the dispatch-cost row's work,
+  per-command, not this row's.
+- **Boundary re-capture storms.** A region that declines at replay
+  (shape change, page-table width change) falls to eager and may
+  re-capture; a per-step storm would multiply the fixed cost 64x. The
+  c2c leg survived SIX whole-graph boundary re-captures serving ~12
+  minutes of batched decode (11:24:06 → 11:36:40) — re-capture is
+  bounded in practice, and the landed width-change guards
+  (`809067ccb`, the batched-lane guard) removed the known decline
+  causes. The gate records `boundary` counts per step; a boundary rate
+  above the c2c precedent fails the gate.
+- **State-slot binding across regions (the #3327 class).** The whole-
+  graph row's founding defect (a replay reading a freed block) recurs at
+  region scale if ANY region's captured trace references a non-
+  persistent buffer. Mitigation is the standing doctrine plus a red-first
+  test per state class (below), not hope: every cross-region slot must
+  be a persistent shadow bound before capture, and the test suite
+  mutates each binding to prove the tests would catch a fresh-tensor
+  commit.
+- **The region-scoped arm becomes a parallel path.** Route through the
+  existing capture machinery (`tenstorrent_capture.cpp`, the
+  `LastTraceBytes` discipline, the RAC re-patch seam) — no new capture
+  implementation beside it. A genuinely unreachable upstream behavior
+  needs one exact tracked exception, not a hand-written twin.
+
+## Tests
+
+Red-first per region handoff (each test fails on the pre-row tree for
+the right reason, mirroring the fix-train's doctest pattern):
+
+1. **Per-layer region capture**: the 27B-shaped decode graph captures N
+   per-layer regions, each `LastTraceBytes()` ≤ 50 MiB, and replays
+   token-exact against the eager reference.
+2. **State handoff across boundaries**: for each cross-region state
+   class (residual, GDN ssm/conv, KV pages), a multi-region replay whose
+   correctness requires the in-place commit — the capture installs a
+   fresh tensor in the mutated variant and the test must fail there
+   (the reviewer's mutation target).
+3. **The fit predicate**: a model whose whole graph fits takes the
+   whole-graph path (byte-identical behavior to today); one that does
+   not takes the region path; neither changes the other's served stream.
+4. **The predicate decline names the missing part**: an over-budget
+   region declines with a message that names the command count and the
+   budget.
+5. **Device gate — the 27B c2 leg**: `BENCH_EXIT=0` on the anchor arm
+   (`--concurrency 2`, both `VT_TT_KEEPQUANT_INT8DOT=0` and `=1`), the
+   served stream token-exact vs eager, and the TPOT table recorded per
+   leg (the economics curve's recipe, one fresh process per leg). The
+   leg's replay-vs-eager token identity is the KV/token-exactness gate.
+
+## Gates
+
+1. Every test above red-before on the pre-row tree, green after.
+2. The device suite at the standing bar (94/95 + the OWED RAC flake
+   recorded, not silently absorbed; this row must not depend on the
+   flaky ordering).
+3. The 27B c2 leg `BENCH_EXIT=0` with the TPOT table and the boundary
+   count at or under the c2c precedent (six re-captures per ~12 min).
+4. Served == eager token-exactness on the anchor arm, both INT8DOT
+   settings.
+5. Standard gates: `agent-preflight.sh --staged`, `check-commit-style`,
+   `check-commit-trailers`, `check-agent-record` on the changed files.
+
+## Evidence plan
+
+Dated `docs/bench-evidence/tt-region-capture-<date>.md`: the per-region
+`LastTraceBytes()` census, the boundary counts, the TPOT table over the
+same legs as `tt-capture-economics-20260927.md` (4/16/64, one fresh
+process per leg, the recorded anchor recipe), the fit predicate's
+decisions per model, and the red/green test logs. The bench-evidence
+file is this row's single measurement record; no number lives in two
+files.
+
+## Stop conditions
+
+- A per-layer region cannot be captured under 50 MiB (the layer's
+  command count is structurally over budget) AND a block-of-layers
+  region cannot hold the boundary discipline → stop, record the census,
+  the arm stays masked on the tt-metal lane (per-command trace cost is
+  not shrinkable from vllm.cpp — escalate with the numbers beside
+  tt-metal#57970).
+- A cross-region state slot cannot bind persistently in-region → stop,
+  record; the 9B precedent says this should not happen, but a
+  design-level impossibility stops the row, not a workaround.
+- Region-scoped replay serves tokens that diverge coherently from eager
+  on the anchor → stop, the adjudication row decides which side matches
+  the vLLM oracle.
+
+## Owed
+
+- The RAC doctest flake from `e39f2cf3f` (order-sensitive under
+  full-suite program-cache history; 126/128 K elems, user-1 second
+  head): NOT this row's fix, but this row's suite must stay green with
+  it recorded; its bisect (which case poisons the variant) stays a
+  separate unit.
+- INT8DOT re-measure downstream: the INT8DOT lever's throughput verdict
+  is re-measured on the served region-scoped arm once it lands; this
+  row records only correctness on both INT8DOT settings.
+- The sampler-bracket interaction: the dispatch-cost row
+  (`.agents/specs/tenstorrent-dispatch-cost.md`,
+  `row/TT-DISPATCH-COST`) brackets inside the ~31.1 s caller window
+  assuming a single whole-graph trace execution; region scope splits
+  that window into 64 per-region executions. Its baseline
+  (31,100-31,172 ms, flat) must be re-recorded on the region-scoped arm
+  before its lever verdict binds; sequence the bracketing measurement
+  BEFORE this row lands, or re-baseline it after — one of the two,
+  recorded in both specs.
+
+## Git integration
+
+One pull request: spec, implementation, tests, evidence, and this row's
+records land together (developer decision, 2026-09-28, recorded in
+`.agents/developer-preferences.md`). Branch `row/tt-27b-region-capture`.
+
+## Now
+
+2026-09-28 (wave 2, same worktree, commits `6473ae731`, `286947603`,
+`39e2ca8ef`): the trace-record audit's inline-upload lever was implemented and
+the leg FALSIFIED the attribution. Landed: the capture-scope upload guard
+(`UploadRows`/`UploadRowsBf16` refuse under capture by name after the
+`VT_TT_TRACE_DEBUG` route print; `AddKernel`'s broadcast operand warms into a
+hash-keyed cache with a named capture-scope miss refusal), and the red-first
+device test pinning the contract (unwarmed capture-scope upload refuses;
+warmed capture records 1,024 B). SUITE: 98 cases / 526,778 assertions,
+526,777 green, the only failure the pre-recorded RAC residual flake. MONEY
+LEG (27B whole-graph c1, evidence
+`docs/bench-evidence/tt-capture-upload-guard-20260928.md`): BENCH_EXIT=1, no
+TPOT — but ZERO `[TT-UP]` uploads under capture, and the demand is
+byte-identical 3,153,969,152 B. With the warmed warmed-MatmulBT region still
+closing at 3,088,384 B, the audit's inline-payload model C is FALSIFIED: the
+~3 MB/command is the quant-matmul program class's own recorded launch stream,
+not a capture-scope upload. The guard stays as hardening; the fit site is
+tt-metal-side and now has a precise next step (dump the `tt::LogDispatch`
+command stream for one captured quant-matmul launch and attribute the ~3 MB
+of `bypass_data`). Owed unchanged: the RAC doctest flake, the INT8DOT
+re-measure, the sampler-bracket re-baseline, the model-scale region
+served-replay doctest, the 6 `EnsureHostBytes DURING CAPTURE` readbacks.
+
+2026-09-28 (wave 1, worktree `row/tt-27b-region-capture-spec`, commits
+`dc99071ad`, `f376b512c`, and the RAC C=1 fix + census beneath): the region
+machinery LANDED and the fit wall was MEASURED. Landed: the per-region
+trace-staging census (`BreakableGraph::region_bytes()`, probe-fed from the
+Tenstorrent registrar; `VT_REGION_CENSUS` prints per segment so a mid-scope
+death still leaves its record), the pure `WholeGraphTraceFits` predicate, the
+dense decode driver's region arm (`VLLM_CPP_REGION_CAPTURE=1`, one region per
+layer via the bare `GraphBreak()`, per-region 50 MiB assertion with a named
+sticky decline), and the red-first device handoff gate (TWO regions on the real
+trace backend, replay byte-identical to eager; region 0 = 2,048 B, region 1 =
+3,088,384 B). An in-flow bug blocked every leg: e39f2cf3f's batched RAC rewrite
+had routed C=1 through batched tensors WarmRacIdx never allocates - the first
+cold decode step segfaulted on any c1 leg
+(ISSUE-LOCAL-01M3M0K390EM40W5R9BR5A2KZ7, fixed same-flow, red
+`/tmp/leg-control-c1.log`, green `/tmp/leg-region-c1-fix.log`).
+MEASURED (evidence `docs/bench-evidence/tt-region-capture-20260928.md`): the
+region arm dies at the 8th segment close on tt-metal `mesh_trace.cpp:125` -
+trace buffer 4,226,469,888 B vs allocation high-water 4,229,506,816 B. All 64
+live regions sum to the whole graph's ~3.15 GB staging demand; segmentation
+does not shrink the fit. This is the spec's stop condition firing on the fit
+axis: the census is recorded, the arm stays masked (env-gated, default-off,
+named decline in the tree), and the numbers escalate beside tt-metal#57970. NO
+TPOT table - no leg completes a step horizon. NEXT: the stop-condition
+adjudication - either a tt-metal-side change (shared/reusable staging, a
+non-zero trace_region_size policy, or #57970 retention recovery large enough
+for 3.15 GB) reopens the row, or the 27B TT arm's next lever is the
+dispatch-cost row (per-command ~30 ms term) on the eager arm while this stays
+masked. Owed unchanged: the RAC doctest flake, the INT8DOT re-measure, the
+sampler-bracket re-baseline, plus the model-scale region served-replay
+doctest.
diff --git a/.gitignore b/.gitignore
index 2a8e09e2f..e8e124927 100644
--- a/.gitignore
+++ b/.gitignore
@@ -44,3 +44,4 @@ generated/
 # (`651994a03`) re-added all three from another worktree. The ignore has to be
 # tracked to hold. Write a PR body OUTSIDE the repository.
 .prbody/
+build2/
diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md
index 9931d7517..d51c8cc29 100644
--- a/docs/ENVIRONMENT.md
+++ b/docs/ENVIRONMENT.md
@@ -102,6 +102,7 @@ These change how the engine runs and have no CLI flag (or complement one).
 | `VT_CPU_SPIN_ROUNDS` | `4096` on aarch64, `256` elsewhere | How many relax rounds a CPU-threadpool waiter spins before yielding its core, in `Threadpool::Barrier` and `Threadpool::PollForWork`. A waiter that never yields costs a full scheduler timeslice per dispatch as soon as the pool is wider than the cores available to it, and a stock run reaches that because the pool defaults to hardware concurrency while the process has other runnable threads. `0` restores the never-yield spin for a same-binary A/B. Neither setting changes any computed value: the yield is a scheduling hint only |
 | `VLLM_PREFIX_CACHING_HASH_SEED` | `0` (fixed) | Seed for the prefix-cache block hash, mirroring vLLM's `PYTHONHASHSEED`. `random` makes block hashes non-deterministic across processes, which takes any persisted or shared KV cache to a 0% hit rate. Keep it fixed if you rely on cross-process prefix reuse |
 | `VLLM_KV_EVENTS_USE_INT_BLOCK_HASHES` | `1` (on) | Whether published KV-cache events carry block hashes as an int (the low 64 bits of the sha256 digest) rather than the raw 32 bytes, mirroring vLLM's env of the same name and its default. Set `0` to publish the raw bytes. Only affects the KV-cache event payload (`--kv-events-config`); it does not change the internal block hashing or the cache itself |
+| `VLLM_CPP_REGION_CAPTURE` | `0` (off) | Tenstorrent backend: capture the 27B decode graph as one region per layer instead of one whole graph (`1` enables). Each region is bounded to 50 MiB of trace staging and an over-budget region declines to eager by name; the whole-graph arm is still attempted first when the fit predicate passes. The default path is unchanged — the arm is insurance for the whole-graph trace-fit wall and for devices where whole-graph capture cannot fit |
 | `VLLM_PLUGINS` | unset (load all registered) | Comma-separated allowlist of general plugins to load in `LoadGeneralPlugins()`, mirroring vLLM's `VLLM_PLUGINS`. Unset loads every registered plugin; an empty string loads none; a list loads only the named plugins. A plugin that throws is logged and skipped (the load never aborts the engine). See [.agents/specs/plugin-system.md](../.agents/specs/plugin-system.md) |
 | `VT_LMCACHE_HOST` | `127.0.0.1` | Default LMCache server host for the `lm://` connector. The `kv_connector_extra_config.host` key overrides it. See [KV offload](KV-OFFLOAD.md) |
 | `VT_LMCACHE_PORT` | `65432` | Default LMCache server port. The `kv_connector_extra_config.port` key overrides it |
diff --git a/docs/bench-evidence/tt-capture-upload-guard-20260928.md b/docs/bench-evidence/tt-capture-upload-guard-20260928.md
new file mode 100644
index 000000000..e248dcb57
--- /dev/null
+++ b/docs/bench-evidence/tt-capture-upload-guard-20260928.md
@@ -0,0 +1,84 @@
+# tt capture-scope upload guard — the leg that falsifies the inline-upload attribution (2026-09-28)
+
+Worktree `row/tt-27b-region-capture-spec`, fixes `286947603` (red test) +
+`39e2ca8ef` (guard + broadcast-operand cache), audit input
+`docs/bench-evidence/tt-trace-record-audit-20260928.md`. Legs:
+`~/.local/logs/maki/CeRgUcXiFK5bYWGSPS4sy/monitor-1790617827-8e22/stdout.log`
+(the full suite), `.../monitor-1790618190-57bb/{stdout,stderr}.log` (the 27B
+whole-graph c1 leg), and the in-tree region-handoff test run.
+
+## What landed
+
+1. **The guard.** `UploadRows` and `UploadRowsBf16`
+   (`src/vt/tenstorrent/tenstorrent_residency.cpp`) refuse any H2D upload with
+   `tt_capture_active()` set, by name, after the required
+   `VT_TT_TRACE_DEBUG` print on the route. `AddKernel`'s broadcast operand
+   (`src/vt/tenstorrent/tenstorrent_ops.cpp`) moves behind a cache keyed by
+   host pointer, geometry, and an FNV-1a hash of the operand's values: the
+   eager pass uploads once, the capture pass serves the resident copy, and a
+   capture-scope miss refuses by name (no silent inline, no value staleness).
+2. **The red-first test** ("kTENSTORRENT capture-scope upload refuses and the
+   warmed capture records the 2 KB floor"): red on `286947603`'s parent — the
+   unwarmed capture-scope upload fired and died on tt-metal's own
+   `TT_FATAL fd_mesh_command_queue.cpp:826 !trace_id_.has_value()` (no named
+   refusal) after printing `[TT-UP] UploadRowsBf16 from_span WRITE during
+   capture rows=1024 cols=1508`. Green after the fix: named refusal + the
+   warmed capture records **1,024 B**.
+
+## The money leg and what it actually showed
+
+27B whole-graph c1 (`VT_TT_TRACE_DEBUG=1 VT_TT_KEEPQUANT_INT8DOT=0`,
+`--num-prompts 2 --input-len 128 --output-len 32 --concurrency 1`):
+**BENCH_EXIT=1. No TPOT table — the leg died at the same capture end.**
+The fatal is byte-identical to the pre-row one: `end_trace_capture` asks for
+**3,153,969,152 B** against 298,568,896 B free
+(`bank_manager.cpp` OOM, `assert.hpp:104`).
+
+But the census around it is decisive:
+
+- **Zero** `[TT-UP]` lines in the whole leg: no upload route (staging,
+  broadcast-Add, rope cache, ids) attempted an H2D write under capture. The
+  doctrine's eager pass already covers every site — the guard never had to
+  fire.
+- **Zero** `[TT-KQ]` keep-quant word-shadow refusals: that route was already
+  refusing/covered.
+- The 6 `EnsureHostBytes DURING CAPTURE` readbacks fired as before (known
+  sync hazard, zero trace bytes, still owed).
+
+So on this pin, **the ~3.15 GB demand persists with zero capture-scope
+uploads — the audit's model C (inline H2D payload) is falsified for the
+default whole-graph arm.** The discriminating experiment the audit itself
+proposed settles it at op scale: the region-handoff test with the MatmulBT
+fully warmed closes region 1 at exactly **3,088,384 B** — the same close the
+audit attributed to a ~3.086 MB inline upload. That payload is NOT an upload;
+it is the recorded per-program command stream of the MatmulBT program class
+(a warmed RmsNorm region still closes at 2,048 B, so binaries-by-relay holds
+for small programs — the ~3 MB rides with the quant-matmul program's launch
+record, mechanism unattributed at dispatch level).
+
+## Verdict and next hypothesis
+
+- The guard and the warmable broadcast operand are correct hardening and stay
+  (they convert the pre-fix raw TT_FATAL into a named refusal and remove the
+  per-call broadcast re-upload in both passes), but they do not shrink the
+  whole-graph record, because there was nothing left to shrink from the
+  upload side.
+- The 27B decode-trace DRAM-fit site closes only on a tt-metal-side
+  attribution: dump the dispatch command stream
+  (`tt::LogDispatch` trace level) for one captured `MatmulBTQuantGrouped`
+  launch and find the ~3 MB of `bypass_data` words. Next candidates: the
+  program's kernel-binary relay pages being re-relayed per launch for
+  programs that miss the 1,024 KB prefetch ringbuffer
+  (`fd_mesh_command_queue.cpp:453` fit decision), or per-launch config-page
+  writes scaling with the quant-program's CB/RTA footprint.
+- Region arm unchanged: `VLLM_CPP_REGION_CAPTURE=1` stays env-gated; the
+  64 live regions summing to the same ~3.15 GB (previous evidence) is now
+  doubly explained — the demand is per-program, so segmentation cannot help
+  either.
+
+## Suite
+
+98 cases, 526,778 assertions: **526,777 passed, 1 failed** — the pre-recorded
+batched-RAC residual flake (96/128 K/V, user-1 second head,
+`test_tenstorrent_backend.cpp:2261`), the exact owed failure the issue
+already records. No new failures from the guard.
diff --git a/docs/bench-evidence/tt-keepquant-rta-fix-20260929.md b/docs/bench-evidence/tt-keepquant-rta-fix-20260929.md
new file mode 100644
index 000000000..dd8583443
--- /dev/null
+++ b/docs/bench-evidence/tt-keepquant-rta-fix-20260929.md
@@ -0,0 +1,90 @@
+# the per-core RTA fix: measured, and what it did NOT close (2026-09-29)
+
+Worktree `row/tt-27b-region-capture-spec` at the fix commit. Follows
+[tt-launch-record-attribution-20260928.md](tt-launch-record-attribution-20260928.md),
+whose verdict named our keepquant program's per-core `SetRuntimeArgs`
+(`tenstorrent_keepquant.cpp:2108-2130`) as the ~2.9 MB per captured launch.
+
+## 1. The fix (landed)
+
+`src/vt/tenstorrent/tenstorrent_keepquant.cpp`:
+
+- The kernel (`kernel_main`) reads ALL runtime words from
+  `SetCommonRuntimeArgs` (14 words: the 3 bank bases + M/K/N/nb/wpb/act_f32/
+  mtile/qb_pad/enc/tcols + grid_x) and derives the per-core slice in-kernel:
+  `c = get_relative_logical_y() * grid_x + get_relative_logical_x()`,
+  `row0 = c * tcols`, `rowc = row0 >= N ? 0 : min(tcols, N - row0)` — the
+  exact guard the deleted host loop applied, including the fully-idle tail.
+  Grid is part of the workload key, so `grid_x` is shape-global per program.
+- The host per-core `SetRuntimeArgs` loop (12 words × grid_cores per call) is
+  DELETED; every word moves to the common-args vector, set once on a
+  workload miss and updated in place on a hit. This call site is the only
+  caller of the program; nothing else needs per-core args on it.
+
+Correctness: the keepquant capture-x2 byte-identity suite stays green on the
+fix (E=1 grouped keep-quant capture, full test suite below); the shape
+carries a partial last core, so the in-kernel clamp is device-proven. A new
+host doctest pins the in-kernel derivation to the deleted loop's values for
+every core across partial/idle tail shapes.
+
+## 2. Measured on the 27B whole-graph capture (the arbiter)
+
+c1 leg (`Qwen3.8-27B-Q4_K_M`, 2x128/32, `VT_TT_TRACE_DEBUG=1`), fresh
+build2 against the new pin `6449cf13f7b`:
+
+| | trace demand at end_mesh_capture |
+|---|---|
+| pre-fix (attribution doc) | 3,153,969,152 B |
+| post-fix (this leg, /tmp/money-c1.log) | 2,925,109,248 B |
+
+The fix removed **228,859,904 B** ≈ 1,037 captured launches × ~920 grid
+cores × one recorded 256 B RTA page each — exactly the per-core
+`SetRuntimeArgs` stream the fix deleted. `BENCH_EXIT=1`: the capture still
+fatals `Out of Memory: Not enough space to allocate 2925109248 B` (bank
+manager, 8 banks). **The whole-graph trace still does not fit DRAM; c1 does
+not serve.** The remaining ~2.93 GB is NOT per-core RTAs.
+
+## 3. The attribution doc's per-launch magnitude was wrong; its direction was right
+
+Two controlled A/Bs on small keepquant captures (same command, red vs green
+binary, new-pin libs):
+
+- region-handoff doctest: region 1 = **2,965,504 B on BOTH binaries** —
+  byte-identical. Dispatch-log counts (TT_METAL_LOGGER_LEVEL=TRACE,
+  build_logging): 22,079 Unique-RTA lines on both.
+- E=1 grouped keep-quant capture: device trace demand
+  **48,316,416 B on BOTH binaries**, capture-x2 byte-identity PASS on both.
+
+So at these shapes the per-core RTA stream contributed ~0 to the recorded
+region — the 2.9 MB per captured command is dominated by the ~250 remaining
+"one-shot program command sequence" fetches per launch (full-grid CB/DFB
+config pages and per-sequence chunks), which are per-launch, not per-core.
+The 27B A/B is the honest measurement: −228.9 MB real, wall standing.
+
+## 4. OPEN NEXT (updated)
+
+The per-core RTA lever is SPENT (landed, correct, ~7.3% of the demand). The
+dominant remaining class is per-launch program command-sequence payload:
+~1,037 launches × ~2.7 MB, i.e. tt-metal records each launch's full-grid
+CB/DFB configuration per sequence. Candidate levers, in traceable order:
+
+1. Count the remaining capture-window classes with the existing
+   logging-enabled discriminator on a 27B capture attempt (the 254-fetch
+   census of the attribution doc, rerun post-fix) — name the per-sequence
+   payload composition before touching anything.
+2. tt-metal-side: whether `create_trace_node` can dedupe/re-reference
+   unchanged full-grid config pages across replays of the same program
+   (upstream question; pin-local experiment first).
+3. Our-side: fewer full-grid CB/config-bearing programs per launch (merge
+   programs), or capture at coarser launch granularity — our-side grid
+   shrink only scales linearly and stays OOM (recorded as a bound, not a
+   fix).
+
+## 5. Test reconciliation
+
+The region-handoff KB-bound gate added red-first for this fix measured
+2,965,504 B before AND after (§3), so the KB floor is not reachable by
+removing per-core RTAs and the gate was removed rather than left red. The
+keepquant capture-x2 byte-identity gates and the derivation-parity doctest
+stand. The 27B trace-fit assertion remains the bench leg (the only vehicle
+at that scale), still failing at 2,925,109,248 B — the row's fit wall.
diff --git a/docs/bench-evidence/tt-launch-record-attribution-20260928.md b/docs/bench-evidence/tt-launch-record-attribution-20260928.md
new file mode 100644
index 000000000..a89768c85
--- /dev/null
+++ b/docs/bench-evidence/tt-launch-record-attribution-20260928.md
@@ -0,0 +1,165 @@
+# tt launch-record attribution — what the ~3.04 MB per captured command is (2026-09-28)
+
+**UPDATE 2026-09-29 (advanced pin, logging-enabled discriminator).** The pin
+advanced from `9161e8fdb27+4` (head `d20b8e27f29`) to upstream-live
+`98134127a7b` + the same local series rebased (pin head `6449cf13f7b`, branch
+`vllm-cpp-pin/20260923-adv`), with `TT_METAL_ENABLE_LOGGING=ON` in a dedicated
+`build_logging` build dir (env `TT_METAL_LOGGER_LEVEL=TRACE`,
+`TT_METAL_LOGGER_TYPES=Dispatch`, names per
+`build_logging/include/tt-logger/tt-logger.hpp:98,182`). Upstream did NOT
+change the per-core trace record between the pins (576+ commits:
+`create_trace_node`/`issue_queue_reserve` in
+`tt_metal/impl/program/dispatch.cpp` untouched; the only nearby changes are
+sub-device setup caching `262a365421f` and trace-allocation-tracker fixes
+`c05eff45369`, neither of which touches the per-core record). The focused leg
+on the new pin reads `region 0 = 2,048 B; region 1 = 2,965,504 B`
+(`/tmp/tregion-newpin.log`, case 1/1 passed, 1,032/1,032 assertions) — the
+pin move shaved ~4% off the record; the mechanism stands. The logger
+discriminator closes the open question: the capture window is exactly 254
+`Writing Program Command Sequence` one-shot fetches summing 2,961,024 B
+(= region 1 + region 0 + trace header pages, `/tmp/burst.log` extraction of
+`/tmp/tregion-newpin.log`), and the window contains **11,040 per-core Unique
+RTA (UNICAST) writes** at 40–48 B payload each plus 110 full-grid CB/DFB
+config pages — i.e. the record is OUR keepquant program's per-core
+`SetRuntimeArgs` stream
+(`src/vt/tenstorrent/tenstorrent_keepquant.cpp:2108-2130`: 12 words per core,
+of which 10 are shape-global constants and only `r0`/`rc` vary, and those two
+are `c*tcols` and its clamp — computable in-kernel from the core coordinate).
+Verdict flips to **OURS** (see §3 below for the tt-metal-side framing this
+update refines): replacing the per-core `SetRuntimeArgs` with in-kernel
+`r0/rc` derivation + `SetCommonRuntimeArgs`-only launches drops the record to
+the RmsNorm-class floor (~2–16 KB per captured launch, ~200×), which closes
+the whole 27B decode-trace DRAM-fit site.
+
+---
+
+Worktree `row/tt-27b-region-capture-spec` at `5f7ac1fa7` (probes were
+uncommitted scratch, reverted before the commit). Fixes the open candidate from
+[tt-capture-upload-guard-20260928.md](tt-capture-upload-guard-20260928.md):
+the 27B whole-graph trace demand is byte-identical 3,153,969,152 B ≈ 1,037 ×
+~3.04 MB, capture-scope uploads are eliminated (zero fires), and the leading
+hypothesis was per-launch kernel-binary relay for programs that miss the
+1,024 KB prefetch ringbuffer. Device evidence: focused legs under
+`flock /home/lu_zero/gpu.lock` on thalia (P150), the region-handoff doctest
+(`tests/vt/test_tenstorrent_backend.cpp:11376`, reproduces census `region 0 =
+2,048 B; region 1 = 3,088,384 B`), run under `gdb -batch` with breakpoints on
+the tt-metal pin's (`~/Sources/tt/tt-metal`, rev `d20b8e27f29`, the exact tree
+the release libs were built from) dispatch internals. Logs: `/tmp/tregion2.log`
+(census reproduction), `/tmp/tt-pcs-census3.log` (per-sequence census),
+`/tmp/tt-cw2.log` / `tt-iq3.log` (per-chunk command-queue census),
+`/tmp/tt-iq4.log` (chunk backtraces), `/tmp/tt-wb3.log` / `tt-iqc.log`
+(capture-window-gated probes).
+
+## 1. The binary-relay hypothesis is REFUTED at source
+
+- The prefetcher "ringbuffer" is 1,024 KB
+  (`tt_metal/impl/dispatch/util/dispatch_settings.cpp:72`, shrunk to 67 KB only
+  when two CQs share a dispatch engine, `:113`); the fit decision is
+  `max_program_kernels_sizeB <= ringbuffer_size()` per MeshWorkload
+  (`tt_metal/distributed/fd_mesh_command_queue.cpp:453`,
+  `tt_metal/distributed/mesh_workload.cpp:272-275`).
+- Either way, the recorded command stream carries the binary BY REFERENCE, not
+  by value: `add_prefetch_relay_paged` sub-commands pointing at the program's
+  DRAM `kernels_buffer` pages when the program does not fit
+  (`tt_metal/impl/program/dispatch.cpp:1942-1961`, multicast
+  `add_prefetch_relay_paged_packed` `:2079-2116`), or
+  `add_prefetch_relay_ringbuffer` when it does (`:1961-1971`). The relay reads
+  device DRAM at replay; the bytes never enter the host-recorded `bypass_data`
+  (`fd_mesh_command_queue.cpp:1660-1665`). Data point in-tree: region 0's
+  complete RmsNorm record is 2,048 B — no room for any compiled binary, so
+  binaries ride by reference for cached programs of any size.
+- tt-metal additionally REFUSES first-time binary loads during capture by name:
+  `MeshWorkloadImpl::load_binaries` TT_FATALs "Cannot load new binaries during
+  trace capture … Warm up before capturing a trace"
+  (`tt_metal/distributed/mesh_workload.cpp:201-205`). Our warm pass covers
+  this; the focused legs never hit it.
+- OUR kernel sizes: the compiled keepquant device kernel (both cache variants,
+  `~/.cache/tt-metal-cache/192705464604149581/kernels/Kernel_Source_Code/
+  {6195984949586344520,16729977526042552009}/brisc/brisc.elf`) measures
+  `.text` 51,648 B + `.data` 3,372 B (`size(1)`) — **18× under the 1,024 KB
+  ringbuffer threshold**. (The unstripped ELF artifact is 1,124,984 B, but that
+  size is `.rela`/symbol tables, not the transfer payload class the relay
+  commands reference.) The "one of our keep-quant binaries exceeds the
+  ringbuffer so the full binary is relayed into the trace per launch"
+  mechanism does not exist on this pin.
+
+## 2. Device attribution of region 1's 3,088,384 B
+
+Method notes: the release libs have `TT_METAL_ENABLE_LOGGING=OFF`
+(`build*/CMakeCache.txt`), so `TT_METAL_LOGGER_LEVEL=TRACE` yields zero
+`tt::LogDispatch` output (verified: the logger leg produced only the doctest
+lines) — the dispatch-log route needs a pin rebuild and was not taken.
+Instead the legs breakpinned the pin's dispatch internals directly:
+
+- **The record is pure command-stream chunks, not data writes.** Gating
+  `SystemMemoryManager::issue_queue_reserve` on our `tt_capture_active()` flag
+  (`/tmp/tt-iqc.log`): the capture window contains **284 reserve chunks and
+  nothing else**; gating `buffer_dispatch::write_to_device_buffer` the same way
+  (`/tmp/tt-wb3.log`): **zero real in-capture buffer-data writes** (the only 2
+  hits register-size 0). So no H2D payload — ours or tt-metal's — is recorded
+  inline in the MatmulBT region, consistent with the zero `[TT-UP]`/`[TT-KQ]`
+  census of the guard leg.
+- **The 3,088,384 B is those 284 chunks.** `issue_queue_reserve` is exactly how
+  the bypass buffer grows in trace recording
+  (`tt_metal/impl/dispatch/system_memory_manager.cpp:519-524`); contiguous runs
+  of the leg's chunk list sum to exactly 3,088,384 B (and to the 2,048 B of
+  region 0 as a single chunk).
+- **The class is the program's own recorded launch stream, scaled per core.**
+  Region 1 records ONE `MatmulBTQuantGrouped` launch whose program spans the
+  FULL worker grid (one `CreateKernelFromString` over
+  `CoreRange{0,0}-{grid.x-1,grid.y-1}`, full-grid CBs, per-core
+  `SetRuntimeArgs` — `src/vt/tenstorrent/tenstorrent_keepquant.cpp:2017-2136`),
+  while region 0's RmsNorm program (ttnn, few cores) records 2,048 B total. The
+  pin's per-program record path writes per-sequence chunks through
+  `write_data_to_cq` (`tt_metal/impl/program/dispatch.cpp:3443-3512`), and the
+  trace node snapshots per-core-range RTA/CB/DFB config payloads
+  (`create_trace_node`, `dispatch.cpp:3520-3651`) — the MB scale rides with the
+  keepquant program's per-core configuration/RTA footprint (284 chunks ≈ 2 ×
+  ~142 grid cores at ~10.9 KB average), not with headers, relay payloads, or
+  any tensor upload.
+
+## 3. Verdict
+
+| Class | Bytes (region 1) |
+|---|---|
+| Command headers + configs for a small program (region 0 floor) | ≤ 2,048 B |
+| Inline H2D payload (audit model C) | **0** — zero in-capture buffer writes |
+| Kernel-binary relay payload | **0** — relay is by reference; kernel is 18× under the ringbuffer threshold anyway |
+| The keepquant program's own recorded per-core launch stream | **~3,086,336 B** (284 chunks, ~10.9 KB each) |
+
+**Fix locus: tt-metal-side.** The record scales with per-core dispatch writes
+for a full-grid custom-kernel program; the shrink levers are (a) tt-metal
+recording per-core config/RTA pages once per program (or by reference) instead
+of per launch, or (b) tt-metal's prefetcher-cache path covering the custom
+kernel's config stream. Expected win if the record drops to the RmsNorm-class
+floor: region record 3,088,384 → ~2-16 KB per command, whole-graph
+3,153,969,152 → tens of MB — the entire 27B decode-trace DRAM-fit site closes.
+
+**Our-side partial lever (not a fix):** shrinking the keepquant grid (fewer,
+fatter cores) scales the record roughly linearly with core count; halving the
+grid halves a 3.15 GB demand — still an OOM, so it only bounds the mechanism,
+it does not close the site. Refuse to ship a "fix" that only re-shapes the
+program.
+
+## Open discriminator (one step)
+
+Which per-sequence class dominates the 284 chunks (per-core config-buffer page
+vs per-core RTA write vs packed-binary sub-commands) needs one leg against a
+logging-enabled pin build (`TT_METAL_ENABLE_LOGGING=ON` +
+`TT_METAL_LOGGER_LEVEL=TRACE`, env names verified at
+`build_Release/include/tt-logger/tt-logger.hpp:98,182`) or an equivalent
+`issue_queue_reserve` caller-tag patch. The class attribution above (per-core
+launch stream, tt-metal-side) does not depend on it.
+
+## What falsified what
+
+- "program does not fit the 1,024 KB ringbuffer → full binary relayed through
+  the trace per launch": **falsified** — relay is by reference in both arms
+  (`dispatch.cpp:1942-1971`), and the keepquant kernel is 55 KB against a
+  1,024 KB threshold.
+- "config/RTA page writes scaling with tensor geometry": **refined** — the
+  scaling is with the program's CORE-GRID footprint (per-core dispatch
+  writes), not tensor geometry; the payload is the recorded command stream
+  itself.
+- Audit model C (inline H2D payload): **falsified on-device** — zero
+  in-capture buffer writes despite the exact 3,088,384 B close.
diff --git a/docs/bench-evidence/tt-region-capture-20260928.md b/docs/bench-evidence/tt-region-capture-20260928.md
new file mode 100644
index 000000000..b594307f3
--- /dev/null
+++ b/docs/bench-evidence/tt-region-capture-20260928.md
@@ -0,0 +1,93 @@
+# tt-27b-region-capture — wave-1 evidence (2026-09-28)
+
+Worktree `row/tt-27b-region-capture-spec`, HEAD ec4e8a824 + the commits below.
+Host thalia (P150), all device work under `/home/lu_zero/gpu.lock`, device reset
+before every leg, `~/Sources/tt/env-tt-common.sh`, tt-metal pin
+`~/Sources/tt/tt-metal-pin` (read-only). One fresh process per leg.
+
+## The region handoff (spec tests 1+2, op scale)
+
+- RED: the pre-row tree has no per-region census —
+  `tests/vt/test_tenstorrent_backend.cpp`'s
+  `kTENSTORRENT region replay: state handoff across a region boundary` does not
+  compile there (no `vt::BreakableGraph::region_bytes()`, no
+  `vt::GraphRegionBytesProbe`), and with `VLLM_CPP_REGION_CAPTURE=1` set the
+  whole-graph arm dies at the trace OOM (every leg below, EXIT=139/1).
+- GREEN (focused, `/tmp/tregion.log`, re-run `/tmp/tregion2.log`): TWO regions on the real tt-metal trace
+  backend — region 0 `RmsNorm` writes the PERSISTENT norm buffer in place,
+  `vt::GraphBreak()` (bare form, no eager call) closes segment 1 and opens
+  segment 2, region 1 `MatmulBT` reads that same buffer. Per-region census:
+  **region 0 = 2,048 B, region 1 = 3,088,384 B**, both ≤ 50 MiB. Each replay is
+  BYTE-IDENTICAL to the eager reference (memcmp, 4,096 B f32). The reviewer's
+  mutation target is the #3327-class defect: give region 0 a FRESH output
+  buffer and region 1's baked address reads freed storage — the replay stops
+  being the eager bytes.
+- Host-side census + predicate (`/tmp` run of `test_breakable_graph -tc="tt-27b*"`):
+  12/12 assertions — `WholeGraphTraceFits` declines only on a measured
+  over-budget estimate (1,037 × 3,187,104 B > 298,568,896 B free = FALSE;
+  zeroed fields = TRUE), and every segment records its staging delta
+  (4 segments from 3 breaks on the recording backend).
+
+## The in-flow bug: RAC C=1 lane (ISSUE-LOCAL-01M3M0K390EM40W5R9BR5A2KZ7)
+
+- RED (`/tmp/leg-control-c1.log`, `/tmp/leg-region-c1-diag.log`): the c1 leg
+  segfaults in `ttnn::copy` inside `ReshapeAndCacheKernel` at the FIRST COLD
+  eager decode step (capturing=0) — identical WITH and WITHOUT
+  `VLLM_CPP_REGION_CAPTURE`, so pre-existing on the base, not the region arm.
+- Cause: e39f2cf3f's batched-lane rewrite routed C=1 through the batched
+  arrays WarmRacIdx never allocates for one user.
+- GREEN (`/tmp/leg-region-c1-fix.log`): C=1 restored verbatim
+  (shared `sharded_in`/`update_idxs`/`page_table`); the cold step and the
+  capture pass now run — the leg reaches deep into per-segment capture (8
+  `BeginCapture`s) before hitting the fit wall below.
+
+## The 27B c1 leg — the fit wall, measured
+
+`/tmp/leg-region-c1-fix.log` (2026-09-28 14:52, fresh reset):
+
+- Per-segment capture proceeds segment by segment; the 8th segment close dies
+  in tt-metal `populate_mesh_buffer` (`mesh_trace.cpp:125`):
+  `Trace buffer at address 4,226,469,888 overlaps with DRAM activity during
+  trace capture. Allocation high water mark: 4,229,506,816`.
+- **The verdict: segmentation does not shrink the total staging demand.** Each
+  live region owns its trace staging until released, and all 64 regions must
+  stay live to replay every step, so the 64 per-layer regions sum to the same
+  ~3.15 GB (1,037 commands × ~3.04 MB/command) the whole graph asked for —
+  against the same free DRAM. The whole-graph OOM and this region-arm
+  collision are ONE structural limit wearing two fatals.
+- `BENCH_EXIT`: 1 (engine fatal). **No TPOT table: the leg cannot complete a
+  step horizon.** The honest baseline stands as the spec records it: the
+  whole-graph arm serves NOTHING (BENCH_EXIT=1 at its own trace OOM) and the
+  region arm serves nothing YET — the 27B TT arm stays masked.
+- Per the spec's stop conditions: the census is recorded, the per-command
+  trace cost is not shrinkable from vllm.cpp, and the numbers go beside
+  tt-metal#57970. The arm remains env-gated (`VLLM_CPP_REGION_CAPTURE`) and
+  default-off; the over-cap decline (named, sticky, per size) is in the tree.
+
+## What did NOT get measured
+
+- The TPOT table (no completed leg), the boundary-re-capture rate, the c2 leg,
+  the INT8DOT correctness leg — all blocked by the fit wall above.
+- The in-tree model-scale served-replay doctest for the region arm — owed to
+  the wave that lands a serving arm.
+
+## Suite
+
+Full TT suite (`ctest -R tenstorrent`, fresh binaries, `/tmp/suite-region2.log`,
+2026-09-28): **526,771 / 526,773 assertions green**; the single failing case is
+the RECORDED OWED RAC doctest flake (126/128 K elems, user-1 second head — the
+order-sensitive-under-full-suite-program-cache residual e39f2cf3f recorded; this
+row does not depend on the flaky ordering). The region handoff case passes
+standalone AND in-suite. The first suite run (`/tmp/suite-region.log`,
+526,771/526,773) failed only this test's own census assertion under ~500 prior
+cases' allocator history (stale staging level) — hardened to assert region 1's
+self-bounded delta; the handoff byte-exactness was green in both runs.
+
+## Gate summary
+
+- Focused handoff doctest: GREEN (standalone + in-suite).
+- Full TT suite: 526,771/526,773, the one failure the recorded OWED flake.
+- 27B c1 leg: BENCH_EXIT=1 (fit wall, recorded above); no TPOT table — the
+  spec's stop condition fired, honestly.
+- Commit gates: check-commit-style, check-commit-trailers, check-agent-record
+  on the changed files — run at landing.
diff --git a/docs/bench-evidence/tt-trace-config-page-repro-20260929.md b/docs/bench-evidence/tt-trace-config-page-repro-20260929.md
new file mode 100644
index 000000000..1b545d754
--- /dev/null
+++ b/docs/bench-evidence/tt-trace-config-page-repro-20260929.md
@@ -0,0 +1,135 @@
+# trace config-page repro: the full-grid × per-launch scaling hypothesis is REFUTED (2026-09-29)
+
+Follows [tt-keepquant-rta-fix-20260929.md](tt-keepquant-rta-fix-20260929.md),
+which left the 27B whole-graph trace at 2,925,109,248 B (1,037 launches,
+~2.82 MB/launch post per-core-RTA fix) with the remaining dominant class
+named "per-launch CB/DFB config pages", suspected to scale as config-page
+count × grid. This note records the standalone reproducer built to test that
+scaling on stock ops — and the refutation.
+
+## 1. The reproducer
+
+`repro_trace_config_pages.cpp` in this directory. Standalone C++ harness
+against the pin's libs (no vllm.cpp), ~85 lines. One trivial op
+(`ttnn::multiply`, bf16, TILE), captured inside a trace at two grid extents:
+
+- case 1: [32,32] interleaved (1-core extent),
+- case 2: [32×110, 64] HEIGHT_SHARDED on L1 across the FULL 11×10 = 110-core
+  compute grid of the P150 (shard grid verified in-process: `memory layout=2
+  buffer=1 shard grid cores=110`),
+- case 3: 8 launches of case 2 inside one trace.
+
+Trace bytes read from `MeshDevice::get_trace_buffers_size()` after
+`end_trace_capture` and before `release_trace` (live total,
+`tt_metal/distributed/mesh_trace.cpp:62` adds `padded_size` on commit and
+`trace_buffer.cpp:24` subtracts on release, so the reading is exactly the one
+trace's padded size; `MeshTraceBuffer::desc->total_trace_size` is not
+reachable through any installed header).
+
+Program cache enabled; each variant warmed before capture (capture refuses
+new binaries, `mesh_workload.cpp:201-205`). Run recipe:
+`flock /home/lu_zero/gpu.lock`, `reset` + 15 s, `TT_METAL_HOME` /
+`LD_LIBRARY_PATH` at `~/Sources/tt/tt-metal-pin/build_release_script`
+(pin `6449cf13f7b` ≈ upstream `98134127a7b` + local series), aarch64
+clang/gnu-16. Log: `/tmp/ttrace-repro3.log`, `EXIT=0`.
+
+## 2. Measured
+
+| op | grid extent | trace bytes |
+|---|---|---|
+| ttnn::multiply [32,32] bf16, 1 launch | 1 core (interleaved) | 17,408 |
+| ttnn::multiply [3520,64] bf16, 1 launch | 110 cores (full-grid sharded) | 17,408 |
+| ttnn::multiply [3520,64] bf16, 8 launches | 110 cores | 139,264 (17,408/launch) |
+
+**No grid scaling.** The full-grid program records byte-identical trace size
+to the 1-core program, and per-launch cost is grid-independent: the 8-launch
+trace is exactly 8 × 17,408. Stock tt-metal records a full-grid program's
+command sequence at the RmsNorm-class KB floor (17.4 KB padded, i.e. ~2–16 KB
+unpadded — the same floor the attribution doc measured for our RmsNorm
+record).
+
+## 3. Verdict
+
+The task's stop condition fired: a full-grid trivial op does NOT reproduce
+the MB-per-launch class, so the generic "per-launch config pages × grid"
+framing is wrong as an upstream ask. tt-metal already records a stock
+multi-core program compactly — consistent with the packed relay
+(`add_prefetch_relay_paged_packed`, `dispatch.cpp:2079-2116`) collapsing
+identical per-core config pages. **No issue filed.**
+
+The 27B residual (~2.82 MB × 1,037 launches) must be specific to our
+keepquant program's config structure, not to grid extent per se. The
+attribution doc's own dispatch-log census already points there: the captured
+window contains ~250 one-shot "Writing Program Command Sequence" fetches per
+launch. Candidate discriminators, in order: (a) number of distinct kernel
+config pages per launch (the fused chain instantiates more kernels than one
+ttnn op), (b) number of distinct CB config pages (many circular buffers vs
+the op's two), (c) failure to hit the packed relay because pages are not
+byte-identical across cores (per-core-varying content), (d) program size
+exceeding the packed-path threshold and falling back to per-page relay.
+
+## 4. Re-scope
+
+The upstream ask is dead in its current shape. The next leg is local: bisect
+keepquant's per-launch recorded bytes against a synthetic program that varies
+(a)-(d) one at a time on this same harness (a compile-time N-kernel/N-CB
+program, same 110-core grid). If a synthetic with our config-page count
+reproduces MB/launch while an equal-grid 1-kernel/2-CB program stays at KB,
+the finding is a config-page-count × pages-not-packed gap worth either an
+upstream report (with the synthetic) or a local program-shape fix. Until
+then this row's remaining trace-fit wall keeps its recorded bound.
+
+Reproducer: `repro_trace_config_pages.cpp`; build and run recipe in §1.
+
+## 5. The named bisect: program-shape classes are ALL innocent (2026-09-29, late)
+
+The §4 leg ran. `repro_bisect_program_shape.cpp` in this directory builds raw
+tt-metal programs shaped exactly like the keepquant int8-dot program
+(DataMovement kernel over the FULL 11x10 grid via one `CoreRange`, CBs,
+`SetCommonRuntimeArgs` only, `DM_DEDICATED_NOC`, warm-then-trace) and varies
+one knob at a time. Recipe as §1, binary `/tmp/bisect`, log
+`/tmp/bisect-run.log` + `/tmp/bt32.log`, `EXIT=0`.
+
+| program | trace bytes/launch |
+|---|---|
+| 1 kernel, 4 CBs, 4 KiB pages (our shape, test scale) | 1,024 |
+| same, 16 KiB CB pages | 1,024 |
+| 8 CBs | 1,024 |
+| 32 KiB / 64 KiB / 128 KiB / 256 KiB .rodata table in the kernel | 1,024 (all) |
+
+Three refutations in one sweep: per-core CB descriptor count, CB page size,
+and KERNEL BINARY SIZE (to 256 KiB, far past our ~50 KiB source) each leave
+the record at the 1,024 B floor. The packed relay collapses all of them.
+Non-identical pages are NOT our 2.82 MB/launch.
+
+## 6. Where the 2,965,504 B actually lives
+
+The region-handoff doctest (KB-floor gate, this row) reads region 1 =
+2,965,504 B at [1,512] -> [1,1024] Q6_K on HEAD 600bfd60e
+(/tmp/region-red.log). The default dispatch there is the W4a GROUPED arm
+(VT_TT_KEEPQUANT_INT8DOT unset), so the captured region is not the int8-dot
+program at all: it is the chunked f32-exact chain
+(tenstorrent_keepquant.cpp:1323-1387) — ceil(N/8) = 8 chunks, and EACH chunk
+runs the eltwise Q6_K word decode `DecodeKeepQuantWordsF32`
+(tenstorrent_keepquant.cpp:629-719: per (h,r) ~12 elementwise/slice/concat
+programs, 2 halves = ~85 recorded programs) plus the slice/typecast/
+to_layout/multiply/sum/permute chain (~8 more). ~680 programs x the ~4-17 KB
+per-program command-sequence floor the reproducer measured = the ~2.97 MB
+region. The class is PROGRAM COUNT in our decode chain, not page identity —
+discriminator (a) from §3, at the whole-region scale, and the ~250
+one-shot-fetch census was these ops, not per-core pages.
+
+The KB single-program path already exists in the tree: the int8-dot kernel
+(one program per launch, bit-exact vs the CPU integer vec_dot oracle) records
+at the floor by §5. The next lever is therefore to collapse the decode chain
+to one custom-kernel program (generalize the int8-dot kernel to the default
+path, or fuse the eltwise decode), NOT any dispatch/config change. The
+region doctest's 64 KiB gate (landed RED at 2,965,504 B) is its arbiter.
+
+## 7. Re-scope (second)
+
+The §4 stop condition fired again, one level down: no program-shape knob
+moves the record; the multiplier is how many programs the keepquant grouped
+arm launches per call. Until that fusion lands, the 27B trace-fit wall keeps
+its recorded 2,925,109,248 B bound and the c1/c2 re-measure legs stay blocked
+behind it.
diff --git a/docs/bench-evidence/tt-trace-config-page-repro-20260929/repro_bisect_program_shape.cpp b/docs/bench-evidence/tt-trace-config-page-repro-20260929/repro_bisect_program_shape.cpp
new file mode 100644
index 000000000..a9148e0ec
--- /dev/null
+++ b/docs/bench-evidence/tt-trace-config-page-repro-20260929/repro_bisect_program_shape.cpp
@@ -0,0 +1,138 @@
+// repro_bisect_program_shape.cpp — the N-kernel/N-CB/binary-size bisect the
+// tt-trace-config-page-repro-20260929 note §4 names. Synthetic raw-tt-metal
+// programs shaped like the keepquant int8-dot program (DataMovement kernel over
+// the full grid, CBs, common RTAs), one knob at a time, each warmed then traced:
+//   argv: [cb_count] [kernel_count] [table_kb] [cb_page_kb] [core_span]
+// Prints trace bytes per launch per variant. Registering the realtime program
+// profiler prints per-program dispatch byte classes (BINARY / RTARGS / CB_CONFIG).
+
+#include <tt-metalium/host_api.hpp>
+#include <tt-metalium/mesh_workload.hpp>
+#include <tt-metalium/mesh_command_queue.hpp>
+#include <tt-metalium/device.hpp>
+#include <tt-metalium/bfloat16.hpp>
+#include <tt-metalium/mesh_trace_id.hpp>
+#include <ttnn/device.hpp>
+#include <ttnn/operations/trace.hpp>
+
+#include <atomic>
+#include <cstdio>
+#include <cstdlib>
+#include <string>
+#include <vector>
+
+using namespace tt::tt_metal;
+
+static uint32_t Capture(ttnn::MeshDevice& dev, const std::function<void()>& body) {
+  auto tid = ttnn::operations::trace::begin_trace_capture(&dev, std::nullopt);
+  body();
+  ttnn::operations::trace::end_trace_capture(&dev, tid, std::nullopt);
+  const uint32_t bytes = dev.get_trace_buffers_size();
+  ttnn::operations::trace::execute_trace(&dev, tid, std::nullopt, /*blocking=*/true);
+  ttnn::operations::trace::release_trace(&dev, tid);
+  return bytes;
+}
+
+static std::string KernelSrc(uint32_t table_kb) {
+  // table_kb > 0 pads the binary with a used const table (lands in .rodata,
+  // exactly like the keepquant IQ tables).
+  std::string s = "#include \"api/dataflow/dataflow_api.h\"\n";
+  s += "#include <cstdint>\n";
+  if (table_kb) {
+    const uint32_t n = table_kb * 256u;
+    s += "static const uint32_t kTable[" + std::to_string(n) + "] = {";
+    for (uint32_t i = 0; i < n; ++i)
+      s += (i ? "," : "") + std::to_string(0x01020304u + i * 0x9E3779B9u) + "u";
+    s += "};\n";
+  }
+  s += R"(
+void kernel_main() {
+  constexpr uint32_t out_base = get_compile_time_arg_val(0);
+  constexpr uint32_t page = get_compile_time_arg_val(1);
+  uint32_t a0 = get_common_arg_val<uint32_t>(0);
+  uint32_t a1 = get_common_arg_val<uint32_t>(1);
+  uint64_t src = get_noc_addr(a0);
+  uint32_t local = out_base;
+  for (uint32_t i = 0; i < page; i += 16) {
+    uint32_t w = a0 + a1 + i;
+    volatile uint32_t* p = reinterpret_cast<volatile uint32_t*>(local + i);
+    p[0] = w;
+  }
+  noc_async_write(out_base, src + 64, page);
+  noc_async_write_barrier();
+}
+)";
+  return s;
+}
+
+struct Variant {
+  const char* name;
+  uint32_t cbs;
+  uint32_t kernels;
+  uint32_t table_kb;
+  uint32_t cb_page_kb;
+};
+
+int main(int argc, char** argv) {
+  std::vector<Variant> variants;
+  if (argc >= 5) {
+    variants.push_back({"custom", static_cast<uint32_t>(std::atoi(argv[1])),
+                        static_cast<uint32_t>(std::atoi(argv[2])),
+                        static_cast<uint32_t>(std::atoi(argv[3])),
+                        static_cast<uint32_t>(std::atoi(argv[4]))});
+  } else {
+    variants = {
+        {"tiny-kern, 4cb, small pages", 4, 1, 0, 4},
+        {"tiny-kern, 4cb, 16K pages  ", 4, 1, 0, 16},
+        {"tiny-kern, 8cb             ", 8, 1, 0, 4},
+        {"2 kernels, 4cb             ", 4, 2, 0, 4},
+        {"big-table 32KB             ", 4, 1, 32, 4},
+        {"big-table 64KB             ", 4, 1, 64, 4},
+        {"big-table 128KB            ", 4, 1, 128, 4},
+        {"big-table 256KB            ", 4, 1, 256, 4},
+    };
+  }
+
+  auto dev = ttnn::open_mesh_device(0, DEFAULT_L1_SMALL_SIZE, 1ull << 28);
+  dev->enable_program_cache();
+  const CoreCoord g = dev->compute_with_storage_grid_size();
+  printf("grid: %ux%u\n", (unsigned)g.x, (unsigned)g.y);
+
+  for (const auto& v : variants) {
+    tt::tt_metal::Program program = tt::tt_metal::CreateProgram();
+    const CoreRange full({0, 0}, {g.x - 1, g.y - 1});
+    const uint32_t cb_page = v.cb_page_kb * 1024u;
+    for (uint32_t c = 0; c < v.cbs; ++c) {
+      CircularBufferConfig cfg(v.cbs * cb_page,
+                               {{static_cast<tt::CBIndex>(static_cast<int>(tt::CBIndex::c_0) + static_cast<int>(c)),
+                                 tt::DataFormat::Float32}});
+      cfg.set_page_size(static_cast<tt::CBIndex>(static_cast<int>(tt::CBIndex::c_0) + static_cast<int>(c)), cb_page);
+      CreateCircularBuffer(program, full, cfg);
+    }
+    const std::string src = KernelSrc(v.table_kb);
+    for (uint32_t k = 0; k < v.kernels; ++k) {
+      // Disjoint vertical halves so two kernels never share a core.
+      CoreRange kr({0u, k * 5u}, {g.x - 1, k * 5u + 4u});
+      std::vector<uint32_t> cargs = {1024 * 64 + k * 16, cb_page};
+      CreateKernelFromString(program, src, kr,
+                             DataMovementConfig{.processor = DataMovementProcessor::RISCV_0,
+                                                .noc = NOC::RISCV_0_default,
+                                                .noc_mode = NOC_MODE::DM_DEDICATED_NOC,
+                                                .compile_args = cargs});
+    }
+    SetCommonRuntimeArgs(program, 0, std::vector<uint32_t>{1024 * 1024, 12345});
+    tt::tt_metal::distributed::MeshWorkload workload;
+    workload.add_program(tt::tt_metal::distributed::MeshCoordinateRange(dev->shape()), std::move(program));
+    tt::tt_metal::Program& prog = workload.get_programs().begin()->second;
+    // warm (load binaries) then trace one launch
+    tt::tt_metal::distributed::EnqueueMeshWorkload(dev->mesh_command_queue(), workload, false);
+    dev->mesh_command_queue().finish();
+    const uint32_t bytes = Capture(*dev, [&] {
+      prog.set_runtime_id(1);
+      tt::tt_metal::distributed::EnqueueMeshWorkload(dev->mesh_command_queue(), workload, false);
+    });
+    printf("%s : cbs=%u kern=%u table=%uKB cbpage=%uKB -> trace %u B (%u/launch)\n",
+           v.name, v.cbs, v.kernels, v.table_kb, v.cb_page_kb, bytes, bytes);
+  }
+  return 0;
+}
diff --git a/docs/bench-evidence/tt-trace-config-page-repro-20260929/repro_trace_config_pages.cpp b/docs/bench-evidence/tt-trace-config-page-repro-20260929/repro_trace_config_pages.cpp
new file mode 100644
index 000000000..c1d066942
--- /dev/null
+++ b/docs/bench-evidence/tt-trace-config-page-repro-20260929/repro_trace_config_pages.cpp
@@ -0,0 +1,81 @@
+// repro_trace_config_pages.cpp — minimal reproducer for per-launch CB/DFB
+// config-page scaling of trace-recorded command sequences (tt-metal,
+// Blackhole P150, eager+trace, program cache on).
+//
+// One trivial op (ttnn::multiply, bf16, zero data payload, kernel binary
+// relayed by reference) is captured inside a trace twice: once at a 1-core
+// extent (interleaved [32,32]) and once spanning the FULL device grid
+// (height-sharded L1 over every compute core). The trace descriptor's
+// total_trace_size (MeshTraceBuffer::desc, host-side assembled command
+// stream) is printed for both, plus an N-launch loop to show per-launch
+// linear growth. Expected if config pages are per-launch: full-grid trace
+// >> 1-core trace, and the N-launch trace scales ~N x.
+
+#include <ttnn/device.hpp>
+#include <ttnn/operations/eltwise/binary/binary.hpp>
+#include <ttnn/operations/trace.hpp>
+#include <tt-metalium/buffer.hpp>
+#include <tt-metalium/bfloat16.hpp>
+#include <tt-metalium/mesh_trace_id.hpp>
+
+#include <cstdio>
+#include <vector>
+
+using namespace tt::tt_metal;
+
+static uint32_t Capture(ttnn::MeshDevice& dev, const std::function<void()>& body) {
+  auto tid = ttnn::operations::trace::begin_trace_capture(&dev, std::nullopt);
+  body();
+  ttnn::operations::trace::end_trace_capture(&dev, tid, std::nullopt);
+  const uint32_t bytes = dev.get_trace_buffers_size();
+  ttnn::operations::trace::execute_trace(&dev, tid, std::nullopt, /*blocking=*/true);
+  ttnn::operations::trace::release_trace(&dev, tid);
+  return bytes;
+}
+
+int main() {
+  auto dev = ttnn::open_mesh_device(/*device_id=*/0, /*l1_small_size=*/DEFAULT_L1_SMALL_SIZE,
+                                    /*trace_region_size=*/1ull << 28);
+  dev->enable_program_cache();
+  const CoreCoord g = dev->compute_with_storage_grid_size();
+  const uint32_t cores = g.x * g.y;
+  printf("device compute grid: %ux%u (%u cores)\n", (unsigned)g.x, (unsigned)g.y, cores);
+
+  // Case 1: [32,32] bf16, interleaved -> a 1-core extent.
+  const tt::tt_metal::TensorSpec small_spec(
+      tt::tt_metal::Shape({32, 32}),
+      tt::tt_metal::TensorLayout(tt::tt_metal::DataType::BFLOAT16, tt::tt_metal::PageConfig(tt::tt_metal::Layout::TILE),
+                                 tt::tt_metal::MemoryConfig{}));
+  ttnn::Tensor a1 = ttnn::Tensor::from_vector<float>(std::vector<float>(32 * 32, 1.0f), small_spec, dev.get());
+  ttnn::Tensor b1 = ttnn::Tensor::from_vector<float>(std::vector<float>(32 * 32, 2.0f), small_spec, dev.get());
+  ttnn::multiply(a1, b1);  // warm: binaries must be loaded before capture
+  const uint32_t one_core = Capture(*dev, [&] { ttnn::multiply(a1, b1); });
+
+  // Case 2: same op, height-sharded across the FULL compute grid.
+  const uint32_t rows = cores * 32;
+  const tt::tt_metal::ShardSpec shard(
+      CoreRangeSet(CoreRange({0, 0}, {g.x - 1, g.y - 1})), {32, 64}, ShardOrientation::ROW_MAJOR);
+  const tt::tt_metal::TensorSpec big_spec(
+      tt::tt_metal::Shape({rows, 64}),
+      tt::tt_metal::TensorLayout(tt::tt_metal::DataType::BFLOAT16, tt::tt_metal::PageConfig(tt::tt_metal::Layout::TILE),
+                                 tt::tt_metal::MemoryConfig{tt::tt_metal::TensorMemoryLayout::HEIGHT_SHARDED,
+                                                            tt::tt_metal::BufferType::L1, shard}));
+  ttnn::Tensor a2 = ttnn::Tensor::from_vector<float>(std::vector<float>(rows * 64, 1.0f), big_spec, dev.get());
+  ttnn::Tensor b2 = ttnn::Tensor::from_vector<float>(std::vector<float>(rows * 64, 2.0f), big_spec, dev.get());
+  ttnn::Tensor out2 = ttnn::multiply(a2, b2);  // warm
+  const auto& mc = out2.memory_config();
+  printf("case2 memory layout=%d buffer=%d shard grid cores=%u\n",
+         (int)mc.memory_layout(), (int)mc.buffer_type(),
+         (unsigned)(mc.shard_spec() ? mc.shard_spec()->grid.num_cores() : 0));
+  const uint32_t full_grid = Capture(*dev, [&] { ttnn::multiply(a2, b2); });
+
+  // Case 3: 8 launches of the same full-grid op inside ONE trace.
+  const uint32_t full_grid_x8 =
+      Capture(*dev, [&] { for (int i = 0; i < 8; ++i) ttnn::multiply(a2, b2); });
+
+  printf("trace bytes, 1 launch, 1-core extent   ([32,32] interleaved):   %u\n", one_core);
+  printf("trace bytes, 1 launch, full-grid extent (%u cores sharded):   %u\n", cores, full_grid);
+  printf("trace bytes, 8 launches, full-grid extent:                    %u (%u/launch)\n",
+         full_grid_x8, full_grid_x8 / 8);
+  return 0;
+}
diff --git a/docs/bench-evidence/tt-trace-record-audit-20260928.md b/docs/bench-evidence/tt-trace-record-audit-20260928.md
new file mode 100644
index 000000000..595055db1
--- /dev/null
+++ b/docs/bench-evidence/tt-trace-record-audit-20260928.md
@@ -0,0 +1,148 @@
+# tt trace-record audit — what the 3.15 GB staging demand actually is (2026-09-28)
+
+Audit of the 27B decode trace-record composition. Evidence:
+`/tmp/leg-region-c1-fix.log` (the 64-region c1 leg, 2026-09-28 14:52),
+`docs/bench-evidence/tt-region-capture-20260928.md` at `180befff5` (the 2-region
+focused census), the whole-graph numbers in
+`.agents/issues/BACKEND-TENSTORRENT-QWEN35/ISSUE-LOCAL-01M3JXEFQKSZP23PP2HWY9G0VQ.md:83,129-134`
+(file committed at `ec4e8a824`; the issue directory currently carries
+`ISSUE-LOCAL-01M3918KQ580Z3NHVRNXVF15FZ.md`), and the tt-metal pin
+`~/Sources/tt/tt-metal-pin` (read-only). No device legs were run; source and
+existing logs suffice.
+
+## 1. The arithmetic model comparison
+
+Raw numbers:
+
+- Whole graph (issue :83,:129-132): `end_trace_capture` asks for one
+  **3,153,969,152 B** staging buffer; the captured decode graph records
+  **1,037 tt-metal commands** → mean **3,041,286 B/command**.
+- Focused 2-region census (region-capture evidence, "The region handoff"):
+  region 0 (one `RmsNorm` command) closes at **2,048 B**; region 1 (one
+  `MatmulBT` command) closes at **3,088,384 B**.
+- 64-region leg: 8 segments closed sequentially, the 8th close dies in
+  `populate_mesh_buffer` (`mesh_trace.cpp:118-125`) at allocation high-water
+  4,229,506,816 B. The per-segment close sizes were not yet logged in that leg
+  (the per-segment staging-delta census landed in the test harness at
+  `180befff5`, after the leg), so the leg constrains the total, not the shape.
+
+Three models fitted:
+
+| Model | Prediction | Verdict |
+|---|---|---|
+| A. per-command uniform | every region ≈ 3.04 MB/command | **refuted by region 0**: one full RmsNorm close = 2,048 B, 1,487× below the mean |
+| B. fixed-per-trace/region F + small c | from the focused test F + c ≤ 3,090,432 B; 64 regions ≤ 198 MB. To reach 3.15 GB needs F ≈ 49.3 MB/region — 16× the largest region close ever measured | **refuted at measured magnitudes** |
+| C. per-command heterogeneous: ~2 KB dispatch headers + inline H2D payload wherever a capture-scope upload fires | region 0 = headers only (✓ 2,048 B); region 1 = headers + ~3.086 MB payload; 1,037 × 3,088,384 B = 3.20 GB ≈ the whole-graph 3.15 GB within 1.5% | **supported, high confidence** |
+
+Region 1's 3,088,384 B decomposes exactly as 2,048 B of command headers plus a
+3,086,336 B payload — a byte count consistent with one bf16 tensor of
+~1.54 M elements written host→device during capture (3,088,384 = 2 × 1,544,192;
+bf16 payloads are recorded byte-exact, see §2). The whole-graph mean
+(3,041,286 B) sits 1.5% below that single-command measurement, i.e. the 3.15 GB
+is ~99.9% per-command inline payload, not fixed overhead. If the 64-region
+fixed-per-trace hypothesis were right, the whole-graph number would have to be
+~64 × 3.09 MB ≈ 198 MB — the leg's 4.23 GB high-water at segment 8 already
+exceeds 21 regions × 49 MB, and the focused test measured a complete region
+close (staging level before/after) at 3.09 MB, not 49 MB.
+
+**The data supports per-command-linear with a ~3 MB dominant term. Confidence:
+high** on the model choice (two independent single-command measurements bracket
+the whole-graph mean); the exact payload identity per command in the 27B graph
+is unmeasured (§4's attribution leg closes that).
+
+## 2. What tt-metal puts in the trace buffer per command
+
+Read path, pin `~/Sources/tt/tt-metal-pin`:
+
+- The trace buffer's content is the host-recorded dispatch command stream:
+  `tt_metal/distributed/mesh_trace.cpp:56-59` sizes it from
+  `MeshTraceDescriptor::total_trace_size`, and `:155-172` writes
+  `mesh_trace_data.data` (uint32 command words) into the DRAM buffer. The
+  fatal the leg hit is the top-down DRAM overlap check at
+  `mesh_trace.cpp:108-152` (message at `:118-125`).
+- Recording is sysmem bypass mode: `FDMeshCommandQueue::record_begin` sets
+  `set_bypass_mode(true)` at `tt_metal/distributed/fd_mesh_command_queue.cpp:1350`;
+  everything the queue issues during capture is appended to
+  `bypass_data` and reduced into `ordered_trace_data` at
+  `fd_mesh_command_queue.cpp:1660-1665` (per-device `max_trace_size`).
+- Per program launch in bypass: `fd_mesh_command_queue.cpp:449-462` creates a
+  trace node; `program_dispatch::create_trace_node`
+  (`tt_metal/impl/program/dispatch.cpp:3520`) snapshots RTAs, circular-buffer
+  configs, dataflow-buffer configs, and cross-node config pages — KBs.
+- **Kernel binaries are NOT embedded per command.** The binary command
+  sequence emits relay commands that reference the program's persistent
+  storage: `add_prefetch_relay_paged` pointing at the program's
+  `kernels_buffer` DRAM pages when uncached
+  (`tt_metal/impl/program/dispatch.cpp:1942-1961`), or
+  `add_prefetch_relay_ringbuffer` pointing at the prefetcher ring buffer when
+  the program fits the cache (`:1961-1971`). The cache is the 1,024 KB
+  prefetch ringbuffer (`tt_metal/impl/dispatch/util/dispatch_settings.cpp:72`;
+  shrunk to 67 KB only when two CQs share one dispatch engine, `:113`), and the
+  fit decision is `max_program_kernels_sizeB <= prefetcher_cache_sizeB` at
+  `fd_mesh_command_queue.cpp:453`. **Data confirms this**: region 0's complete
+  2,048 B close cannot carry any compiled binary, so binaries ride by
+  reference.
+- What IS copied inline: any H2D write issued during capture. Bypass mode
+  records the write command with its payload so the replay re-writes the same
+  bytes. This is the only MB-scale per-command mechanism in the record path.
+- Fixed per-trace cost: `record_begin`/`record_end` wrappers, the
+  `exec_buf_end` epilogue (`fd_mesh_command_queue.cpp:1661`), completion and
+  event-reset bookkeeping (`:1691-1705`) — bytes to KB, consistent with region
+  0's 2,048 B.
+
+So tt-metal's minimal command costs ~2 KB (headers + go signal + snapshot
+configs); the fixed per-trace cost is the same order. **A 3 MB command is not
+tt-metal overhead — it is a captured inline write.**
+
+## 3. Our capture path — staging findings
+
+- `EnsureDevice2D` (`src/vt/tenstorrent/tenstorrent_residency.cpp:337`)
+  lazily uploads via `from_vector` whenever a tensor is not yet device-current.
+  If the first touch happens inside a capture scope, the whole tensor is
+  recorded inline into the trace. There is no refusal or debug print on this
+  route (the `tt_capture_active()` guards there only divert reshape paths).
+- Broadcast `AddKernel` builds a replicated `[rows, d]` host tensor and
+  `from_vector`s it unconditionally — with an explicit
+  "[TT-UP] AddKernel from_vector WRITE during capture" debug line
+  (`src/vt/tenstorrent/tenstorrent_ops.cpp:113-133`). A rows×d×4 B payload is
+  inlined per captured Add.
+- `WarmDecodeIds` refreshes the ids device buffer with a `copy_to_device`
+  per step, including the capture step
+  (`src/vt/tenstorrent/tenstorrent_capture.cpp:314-317`) — small (n × 4 B).
+- Six `EnsureHostBytes DURING CAPTURE` fired in the c1 leg
+  (log `/tmp/leg-region-c1-fix.log:1006-1011`; print at
+  `tenstorrent_residency.cpp:1114`). These are blocking device→host readbacks
+  inside capture — a capture-scope sync hazard, but they add no trace bytes.
+- Region-1 composition explained: the focused test's `MatmulBT` command
+  recorded 2,048 B of headers plus one ~3.086 MB inline H2D payload —
+  byte-count consistent with the matmul's bf16 weight (~1.54 M elems) being
+  uploaded during capture via the lazy staging route instead of warmed before
+  it.
+
+## 4. Conclusion and the lever
+
+**Per-command dominates; fewer/bigger traces is the wrong lever.**
+Quantified best cases:
+
+- If the ~3 MB/command term is an inline capture-scope upload, eliminating it
+  drops every command to region 0's measured floor: 1,037 × ~2 KB ≈
+  **2.1 MB total** — three orders of magnitude under the 3.15 GB wall, and the
+  whole-graph capture fits 2.2 GB trivially. This is the actionable lever:
+  attribute then hoist the upload (warm every tensor before
+  `TraceBeginCapture`; refuse/flag any `from_vector`/`EnsureDevice2D` upload
+  with `tt_capture_active()`, the way `tenstorrent_ops.cpp:130` already
+  prints for AddKernel's broadcast route).
+- The fewer-traces alternative (capture the biggest 8 layers whole) does not
+  survive the arithmetic: even taking the measured per-command mean at face
+  value, 8 layer-regions ≈ 130 commands × 3.04 MB ≈ 395 MB — it only buys a
+  ~5.6× demand cut by dropping 87% of the commands, and it is dominated by the
+  upload fix, which buys ~1,500×.
+
+Discriminating leg (cheap, no new mechanism): rerun the focused 2-region case
+with `VT_TT_TRACE_DEBUG=1` plus a print added on `EnsureDevice2D`'s upload
+route; region 1's payload attribution line either names the weight upload or
+falsifies §3's reading. The committed per-segment staging-delta census
+(`180befff5`) gives the per-region denominators for the same leg.
+
+No tt-metal record-size ask is indicated by this data: the record path is
+already minimal for warmed programs (region 0 = 2,048 B proves it).
diff --git a/include/vt/breakable_graph.h b/include/vt/breakable_graph.h
index 11ccb7bdb..eab9e12ef 100644
--- a/include/vt/breakable_graph.h
+++ b/include/vt/breakable_graph.h
@@ -195,6 +195,34 @@ struct GraphBreakStats {
 GraphBreakStats GetGraphBreakStats();
 void ResetGraphBreakStats();
 
+// ---------------------------------------------------------------------------
+// The per-region trace-staging byte probe (tt-27b-region-capture).
+// ---------------------------------------------------------------------------
+// The seam is backend-agnostic and cannot name tt-metal's trace buffer; the
+// Tenstorrent registrar installs `LastTraceBytesForTest` here, and every
+// segment close then records what it contributed to the staging (see
+// `BreakableGraph::region_bytes`). Uninstalled (nullptr, the default) records
+// nothing: the census is opt-in per backend, never a cross-backend cost.
+using GraphRegionBytesProbe = int64_t (*)();
+void SetGraphRegionBytesProbe(GraphRegionBytesProbe probe);
+
+// The WHOLE-GRAPH FIT predicate (tt-27b-region-capture, spec `## Design`).
+// Pure and injectable so a unit case can decide both arms without a device.
+// vLLM's capture-size list walks capture sizes DOWNWARD past a size that does
+// not fit (gpu_model_runner.py's `may_replay_capture`/eager fallback): the
+// same polarity — over budget declines capture, names nothing here, the
+// caller names it. Zeroed fields (a backend with no census, a lane with no
+// measurement) answer TRUE: no measurement is never evidence of no fit.
+struct WholeGraphFitEstimate {
+  int64_t recorded_command_count = 0;  // the VT_TT_TRACE_DEBUG census
+  int64_t per_command_bytes = 0;       // measured staging bytes per command
+  int64_t free_bytes = 0;              // device free DRAM at decision time
+};
+inline bool WholeGraphTraceFits(const WholeGraphFitEstimate& e) {
+  if (e.per_command_bytes <= 0 || e.recorded_command_count <= 0) return true;
+  return e.recorded_command_count * e.per_command_bytes <= e.free_bytes;
+}
+
 // The one kill switch, read ONCE per process into a function-local static.
 // Today six drivers each read `VLLM_CPP_CUDAGRAPH` for themselves and the three
 // single-shape drivers invented their own switch instead, so there is no one
@@ -275,6 +303,16 @@ class BreakableGraph {
   // marker in `qwen3_5.cpp`, which is a bench-only build.
   void* segment(size_t i) const { return i < segments_.size() ? segments_[i] : nullptr; }
   size_t break_count() const { return break_fns_.size(); }
+  // THE PER-REGION TRACE-STAGING CENSUS (tt-27b-region-capture). One entry per
+  // segment, in capture order: the bytes the backend's trace staging held at
+  // that segment's close MINUS the bytes held at its open — i.e. what THIS
+  // segment alone contributed to the trace buffer. Filled only when a probe is
+  // installed (`SetGraphRegionBytesProbe`); empty (and never consulted) on a
+  // backend that registers none, so the CUDA lane records no census and pays
+  // nothing. A driver that captures regions asserts each entry against its
+  // per-region budget and declines by name when one is over — the whole graph
+  // fit question, answered per region instead of once for 1,037 commands.
+  const std::vector<int64_t>& region_bytes() const { return region_bytes_; }
   bool captured() const { return !segments_.empty(); }
   int64_t replay_count() const { return replays_; }
 
@@ -323,6 +361,7 @@ class BreakableGraph {
   Backend* backend_ = nullptr;
   std::vector<void*> segments_;
   std::vector<std::function<void()>> break_fns_;
+  std::vector<int64_t> region_bytes_;
   int64_t replays_ = 0;
   bool capture_failed_ = false;
   std::exception_ptr capture_error_;
@@ -471,6 +510,9 @@ class GraphCaptureScope {
     Event* e;
   };
   std::vector<ForkedQueue> forks_;
+  // The trace-staging byte level when THIS segment opened, so the close can
+  // record what the segment contributed (see `BreakableGraph::region_bytes`).
+  int64_t region_bytes_base_ = 0;
   // Joins every outstanding fork onto the capture queue and clears the set.
   // Called by `EndSegment` BEFORE `Backend::EndCaptureGraph`.
   void JoinOutstandingForks();
diff --git a/scripts/env-doc-allowlist.txt b/scripts/env-doc-allowlist.txt
index 53a04bd1d..ab7f20510 100644
--- a/scripts/env-doc-allowlist.txt
+++ b/scripts/env-doc-allowlist.txt
@@ -267,3 +267,4 @@ VT_V4_W32_WARPS
 VT_KEV_LAYER_DUMP
 VT_VK_DISABLE
 VT_VK_DISABLE_PAGED_ATTN
+VT_REGION_CENSUS
diff --git a/src/vllm/model_executor/models/qwen3_5.cpp b/src/vllm/model_executor/models/qwen3_5.cpp
index 0f5e187c3..4c9f02d2a 100644
--- a/src/vllm/model_executor/models/qwen3_5.cpp
+++ b/src/vllm/model_executor/models/qwen3_5.cpp
@@ -51,7 +51,9 @@
 #include <map>
 #include <memory>
 #include <mutex>
+#include <numeric>  // tt-27b-region-capture: std::accumulate over the region census
 #include <string>
+#include <string_view>  // tt-27b-region-capture: VLLM_CPP_REGION_CAPTURE parse
 #include <unordered_map>
 #include <utility>
 #include <optional>
@@ -10200,6 +10202,15 @@ static DBuf DenseForwardLayers(Dev d, const Tensor& hidden_in,
     // DFlash DF-AUX-TAPS: capture (hidden+res) at configured boundaries. Inert
     // (no-op) when aux_out is null — every non-DFlash caller.
     MaybeCaptureAuxTap(d, l, aux_layer_ids, aux_out, hidden.t(), res.t(), T, H);
+    // tt-27b-region-capture: ONE REGION PER LAYER. The bare break splits the
+    // kPiecewise scope into a new segment with NO eager call and NO
+    // destination — the region boundary is a pure capture split, and the
+    // handoff is the in-place one: hidden/res are pool-backed buffers whose
+    // captured addresses the #2274 pinning holds for the graph's life, and
+    // the GDN ssm/conv + KV state slots are persistent shadows committed IN
+    // PLACE (tenstorrent_gdn.cpp's W3 discipline). Inert (a counter tick)
+    // in every kFull scope and every eager call — byte-identical to today.
+    vt::GraphBreak();
     // VT_DUMP_ACT (issue #41, ROCm 0.8B forward-divergence fix spike W1; keyed
     // and completed for #2590): dump the residual stream after each layer.
     //
@@ -12090,6 +12101,39 @@ ForwardLogits Qwen3_5DecodeGraph::Step(
   return fl;
 }
 
+// ─── tt-27b-region-capture: the region-scoped decode-capture arm ─────────────
+// The 27B decode graph does not fit ONE whole-graph trace: end_trace_capture
+// asks for one ~3.15 GB staging buffer against ~298 MB free (the spec's
+// `## Scope` census, 1,037 recorded commands), and the whole-graph arm serves
+// nothing. Region scope splits the same command stream into ONE REGION PER
+// LAYER: the decode driver opens its capture kPiecewise and DenseForwardLayers
+// emits an in-place boundary after each layer (`vt::GraphBreak()`, the bare
+// form — no eager call, no destination; the layer outputs flow device-side
+// through the SAME persistent buffers a whole-graph capture bakes, which is
+// exactly the in-place handoff discipline the #3327 class demands — no region
+// boundary installs, frees, or re-shadows a state tensor). The replay is the
+// container's host loop: segment, (no-op break), segment, ... — the per-region
+// runtime-arg re-patch the RAC per-user mechanism already serves, because the
+// RAC/rope hooks read the SAME persistent device inputs every segment bakes.
+// Sizing: 1,037 commands / 64 layers ≈ 16.2 commands per layer ≈ 48.6 MiB at
+// the measured 3.04 MB per command — the GDN precedent's 50 MiB region budget
+// (tenstorrent_capture.cpp:90), asserted per region from the probe-fed census
+// (`BreakableGraph::region_bytes()`), with an over-cap region declining the
+// capture BY NAME (below).
+// OFF by default in this slice (`VLLM_CPP_REGION_CAPTURE=1` opts in): the fit
+// predicate's automatic model-by-model wiring (`vt::WholeGraphTraceFits`) is
+// the next wave — wiring it now would re-route the 9B whole-graph arm the
+// census cannot yet price per model.
+static bool RegionCaptureRequested() {
+  static const bool v = [] {
+    const char* e = std::getenv("VLLM_CPP_REGION_CAPTURE");
+    return e != nullptr && e[0] != '\0' && std::string_view(e) != "0";
+  }();
+  return v;
+}
+// The GDN precedent's fit number (tenstorrent_capture.cpp:90).
+constexpr int64_t kRegionCaptureBudgetBytes = 50 * 1024 * 1024;
+
 // ─── Qwen3_5DenseDecodeGraph (27B dense decode CUDA-graph driver) ────────────
 // The 27B DENSE sibling of Qwen3_5DecodeGraph. Same cold→warm→replay state
 // machine, same padded-batch capture set (kDecodeGraphSizes), same persistent
@@ -12167,6 +12211,11 @@ struct Qwen3_5DenseDecodeGraph::Impl {
     vt::BreakableGraph graph;
     int fa_cols = -1;                 // captured block-table column count
     bool warm = false;
+    // tt-27b-region-capture: a named per-region over-budget DECLINE (see the
+    // census below) is sticky for this size — an over-budget layer's command
+    // stream does not shrink between steps, so re-capturing every step would
+    // be the boundary storm the spec's risk names. The slot serves EAGER.
+    bool region_declined = false;
     int64_t replays = 0;
     // R2: the cur_pos the device held after this slot's last seeding step or
     // replay (WarmDecodePos continuation predicate, qwen3.cpp #2469).
@@ -12742,7 +12791,7 @@ ForwardLogits Qwen3_5DenseDecodeGraph::Step(
 
   // Warm: the pool + residency were warmed for this size by the previous (eager)
   // step. CAPTURE the dense layer region once, instantiate the graph, launch it.
-  if (s.warm) {
+  if (s.warm && !s.region_declined) {
     // #1380: THE POOL MUST BE ABLE TO SERVE THE WHOLE CAPTURED FORWARD, not one
     // block of one tensor. This used to alloc-and-free a single [S, vocab] f32
     // block, on the reasoning that the capture RETAINS its logits while the
@@ -12894,10 +12943,21 @@ ForwardLogits Qwen3_5DenseDecodeGraph::Step(
     // nothing), cap_end the scope DESTRUCTION (EndCaptureGraph: the tt-metal
     // trace finalize + trace-buffer build — the capture-invocation cost).
     StepPhaseClk::time_point sph_tc1{};
+    // tt-27b-region-capture: the mode IS the fit decision. Region scope
+    // (env opt-in this slice; the automatic `vt::WholeGraphTraceFits`
+    // wiring is the next wave) splits the same command stream one layer per
+    // region — the whole-graph staging demand (~3.15 GB for 1,037 recorded
+    // commands) never accrues, because each region's trace buffer lands
+    // inside the 50 MiB budget the census below asserts. Every GraphBreak
+    // in the forward is INERT in the kFull arm, so the default shape is
+    // byte-identical to the one the comment above records.
+    const bool region_scope = RegionCaptureRequested();
     {
       const StepPhaseClk::time_point sph_tc0 =
           sph.on ? StepPhaseClk::now() : sph.t0;
-      vt::GraphCaptureScope scope(b, impl_->queue, s.graph, vt::GraphCaptureMode::kFull);
+      vt::GraphCaptureScope scope(b, impl_->queue, s.graph,
+          region_scope ? vt::GraphCaptureMode::kPiecewise
+                       : vt::GraphCaptureMode::kFull);
       if (sph.on) sph.cap_begin_ms = StepPhaseMsOf(sph_tc0, StepPhaseClk::now());
       sph_tc1 = StepPhaseClk::now();
       if (d.q.device.type == vt::DeviceType::kTENSTORRENT) {
diff --git a/src/vt/breakable_graph.cpp b/src/vt/breakable_graph.cpp
index 71b0520b5..c9cacf67a 100644
--- a/src/vt/breakable_graph.cpp
+++ b/src/vt/breakable_graph.cpp
@@ -123,12 +123,27 @@ void CopyOutput(Backend& b, Queue& q, std::map<std::string, Tensor>& dst,
 
 BreakableGraph::~BreakableGraph() { Reset(); }
 
+// tt-27b-region-capture: the installed byte probe. Function-local static so TU
+// order never decides who owns it; the Tenstorrent registrar installs
+// `LastTraceBytesForTest`, everyone else leaves the null that records nothing.
+GraphRegionBytesProbe& RegionBytesProbe() {
+  static GraphRegionBytesProbe probe = nullptr;
+  return probe;
+}
+
+void SetGraphRegionBytesProbe(GraphRegionBytesProbe probe) {
+  RegionBytesProbe() = probe;
+}
+
 void BreakableGraph::Reset() {
   if (backend_ != nullptr) {
     for (void* g : segments_) backend_->DestroyGraph(g);
   }
   segments_.clear();
   break_fns_.clear();
+  // The census describes the released capture, the same rule `replays_` and the
+  // failure accessors below follow.
+  region_bytes_.clear();
   // The replay count describes the graph that was just released. Leaving it
   // behind makes the next capture report replays it never ran, and G3's whole
   // job is to be the number nobody has to trust twice.
@@ -252,6 +267,9 @@ GraphCaptureScope::~GraphCaptureScope() {
 
 void GraphCaptureScope::BeginSegment() {
   if (!active_ || segment_open_) return;
+  // tt-27b-region-capture: the byte level BEFORE this segment records, so the
+  // close computes what THIS region contributed alone.
+  if (GraphRegionBytesProbe probe = RegionBytesProbe()) region_bytes_base_ = probe();
   b_->BeginCapture(*q_);
   segment_open_ = true;
 }
@@ -266,6 +284,17 @@ void GraphCaptureScope::EndSegment() {
   segment_open_ = false;  // cleared FIRST: a throwing end must not be retried
   void* seg = b_->EndCaptureGraph(*q_);
   g_->segments_.push_back(seg);
+  // tt-27b-region-capture: the per-region census entry (see region_bytes()).
+  // VT_REGION_CENSUS prints per segment, because a capture that DIES
+  // mid-scope (the 27B fit collision does) must still leave its per-region
+  // record: the post-scope summary only exists for a capture that finished.
+  if (GraphRegionBytesProbe probe = RegionBytesProbe()) {
+    const int64_t bytes = probe() - region_bytes_base_;
+    g_->region_bytes_.push_back(bytes);
+    if (std::getenv("VT_REGION_CENSUS") != nullptr)
+      std::fprintf(stderr, "[REGION-CAPTURE] segment %zu: %lld B staging\n",
+                   g_->region_bytes_.size() - 1, static_cast<long long>(bytes));
+  }
   g_segments.fetch_add(1, std::memory_order_relaxed);
 }
 
diff --git a/src/vt/tenstorrent/tenstorrent_backend.cpp b/src/vt/tenstorrent/tenstorrent_backend.cpp
index fd061b40f..4fdc0abd2 100644
--- a/src/vt/tenstorrent/tenstorrent_backend.cpp
+++ b/src/vt/tenstorrent/tenstorrent_backend.cpp
@@ -26,6 +26,7 @@
 //     reference tier stays gated off for this device — no free correctness
 //     net for an unregistered op.
 #include "vt/backend.h"
+#include "vt/breakable_graph.h"
 #include "vt/tenstorrent/tenstorrent_device.h"
 
 #include <cstdlib>
@@ -129,6 +130,11 @@ struct Registrar {
     // static init (unspecified TU order rules out trusting another TU's
     // initializer to have probed already).
     if (!DeviceAvailable()) return;
+    // tt-27b-region-capture: install the trace-staging byte probe so every
+    // capture segment records what it contributed (`BreakableGraph::
+    // region_bytes()`). The 50 MiB per-region budget is asserted by the driver
+    // that captures regions, against these numbers.
+    vt::SetGraphRegionBytesProbe(&LastTraceBytesForTest);
     static TenstorrentBackend backend;
     RegisterBackend(DeviceType::kTENSTORRENT, &backend);
   }
diff --git a/src/vt/tenstorrent/tenstorrent_device.h b/src/vt/tenstorrent/tenstorrent_device.h
index e68442972..5f6600abb 100644
--- a/src/vt/tenstorrent/tenstorrent_device.h
+++ b/src/vt/tenstorrent/tenstorrent_device.h
@@ -310,9 +310,17 @@ void WarmPagedKvShadow(void* k_cache_data, void* v_cache_data,
                       int64_t num_blocks, int64_t block_size,
                       int64_t num_kv_heads, int64_t head_size,
                       int64_t used_blocks);
+// TEST-ONLY (the DeviceShadowExact pattern): read the paged-KV DEVICE shadow
+// for this cache buffer back to host (row-major [nb, bs, nkv, d] floats), so
+// a focused case can verify what the captured RAC replay wrote without going
+// through the stale host master.
+bool ReadPagedKvShadowForTest(const void* k_cache_data, float* dst, int64_t n);
 #else
 inline void WarmPagedKvShadow(void*, void*, int64_t, int64_t, int64_t, int64_t,
                               int64_t) {}
+inline bool ReadPagedKvShadowForTest(const void*, float*, int64_t) {
+  return false;
+}
 #endif
 
 // GDN conv-state shadow serveability (decode side): true when the transposed
diff --git a/src/vt/tenstorrent/tenstorrent_internal.h b/src/vt/tenstorrent/tenstorrent_internal.h
index 257bf39a2..f7d9b4fa3 100644
--- a/src/vt/tenstorrent/tenstorrent_internal.h
+++ b/src/vt/tenstorrent/tenstorrent_internal.h
@@ -34,6 +34,7 @@
 #include <map>
 #include <mutex>
 #include <optional>
+#include <unordered_map>
 #include <thread>
 #include <tuple>
 #include <utility>
@@ -116,6 +117,9 @@ std::tuple<ttnn::Tensor, std::optional<ttnn::Tensor>> chunk_gated_delta_rule(
     const std::optional<ttnn::Tensor>& initial_state = std::nullopt,
     bool output_final_state = false, uint32_t chunk_size = 64,
     bool use_qk_l2norm = false, bool output_head_major = false,
+    // The advanced pin (98134127a7b, chunk_gated_delta_rule.hpp:41) adds the
+    // multicast flag at this position; mirror 1:1.
+    bool use_mcast = true,
     const std::optional<ttnn::MemoryConfig>& memory_config = std::nullopt,
     const std::optional<DeviceComputeKernelConfig>& compute_kernel_config = std::nullopt,
     const std::optional<ttnn::Tensor>& eye = std::nullopt,
diff --git a/src/vt/tenstorrent/tenstorrent_keepquant.cpp b/src/vt/tenstorrent/tenstorrent_keepquant.cpp
index 508ec5a7e..7c26d9907 100644
--- a/src/vt/tenstorrent/tenstorrent_keepquant.cpp
+++ b/src/vt/tenstorrent/tenstorrent_keepquant.cpp
@@ -1559,19 +1559,30 @@ constexpr const char* kKeepQuantInt8DotKernelSrc = R"TTKQ(
 #include "api/dataflow/dataflow_api.h"
 #include "keepquant_kernel_code.h"
 
-// Per-core runtime args (the SetRuntimeArgs stream).
-constexpr uint32_t ARG_M = 0;
-constexpr uint32_t ARG_K = 1;
-constexpr uint32_t ARG_N = 2;
-constexpr uint32_t ARG_NB = 3;        // weight blocks per row (K / elems)
-constexpr uint32_t ARG_WPB = 4;       // staged i32 words per block
-constexpr uint32_t ARG_ACT_F32 = 5;   // 1: f32 activation bytes, 0: bf16
-constexpr uint32_t ARG_ROW0 = 6;      // first weight column of this core (4*group0)
-constexpr uint32_t ARG_ROWC = 7;      // real columns this core dots (0: idle)
-constexpr uint32_t ARG_MTILE = 8;     // activation rows per quantize tile
-constexpr uint32_t ARG_QB_PAD = 9;    // 16B-aligned activation-quant row bytes
-constexpr uint32_t ARG_ENC = 10;      // 0/1/2/3/4 = Q4_K/Q5_K/Q6_K/Q8_0/IQ3_XXS
-constexpr uint32_t ARG_TCOLS = 11;    // padded tile width (uniform): groups_per_core*4
+// Shape-global runtime args (the SetCommonRuntimeArgs stream). The 27B
+// trace-fit fix (tt-launch-record-attribution-20260928): every word here is
+// identical across cores, so the program records ONE common-args page per
+// captured launch instead of one per-core Unique RTA UNICAST page per core —
+// 11,040 recorded pages ≈ 2.97 MB per launch is what blew the 27B whole-graph
+// trace out of DRAM. Indices 0-2 are the three bank bases the original
+// common-args set carried; the shape words follow. The only per-core-varying
+// words (row0 = c*tcols, rowc = its clamp) are derived in-kernel from the
+// core coordinate.
+constexpr uint32_t CARG_W_ADDR = 0;
+constexpr uint32_t CARG_A_ADDR = 1;
+constexpr uint32_t CARG_O_ADDR = 2;
+constexpr uint32_t CARG_M = 3;
+constexpr uint32_t CARG_K = 4;
+constexpr uint32_t CARG_N = 5;
+constexpr uint32_t CARG_NB = 6;        // weight blocks per row (K / elems)
+constexpr uint32_t CARG_WPB = 7;       // staged i32 words per block
+constexpr uint32_t CARG_ACT_F32 = 8;   // 1: f32 activation bytes, 0: bf16
+constexpr uint32_t CARG_MTILE = 9;     // activation rows per quantize tile
+constexpr uint32_t CARG_QB_PAD = 10;   // 16B-aligned activation-quant row bytes
+constexpr uint32_t CARG_ENC = 11;      // 0..13, the enc-select dispatch below
+constexpr uint32_t CARG_TCOLS = 12;    // padded tile width: groups_per_core*4
+constexpr uint32_t CARG_GRID_X = 13;   // core-grid width: c = y*grid_x + x
+constexpr uint32_t kNumCommonArgs = 14;
 
 // CB scratch (self-cycled: reserve -> use -> push -> pop; no consumer core).
 constexpr uint32_t CB_F32 = 0;  // one activation row widened to f32 (K*4 B)
@@ -1597,18 +1608,27 @@ void kernel_main() {
   const auto acc_o =
       TensorAccessor(args_o, get_common_arg_val<uint32_t>(2));
 
-  const uint32_t M = get_arg_val<uint32_t>(ARG_M);
-  const uint32_t K = get_arg_val<uint32_t>(ARG_K);
-  const uint32_t N = get_arg_val<uint32_t>(ARG_N);
-  const uint32_t nb = get_arg_val<uint32_t>(ARG_NB);
-  const uint32_t wpb = get_arg_val<uint32_t>(ARG_WPB);
-  const uint32_t act_f32 = get_arg_val<uint32_t>(ARG_ACT_F32);
-  const uint32_t row0 = get_arg_val<uint32_t>(ARG_ROW0);
-  const uint32_t rowc = get_arg_val<uint32_t>(ARG_ROWC);
-  const uint32_t mtile = get_arg_val<uint32_t>(ARG_MTILE);
-  const uint32_t qb_pad = get_arg_val<uint32_t>(ARG_QB_PAD);
-  const uint32_t enc = get_arg_val<uint32_t>(ARG_ENC);
-  const uint32_t tcols = get_arg_val<uint32_t>(ARG_TCOLS);  // padded tile width
+  const uint32_t M = get_common_arg_val<uint32_t>(CARG_M);
+  const uint32_t K = get_common_arg_val<uint32_t>(CARG_K);
+  const uint32_t N = get_common_arg_val<uint32_t>(CARG_N);
+  const uint32_t nb = get_common_arg_val<uint32_t>(CARG_NB);
+  const uint32_t wpb = get_common_arg_val<uint32_t>(CARG_WPB);
+  const uint32_t act_f32 = get_common_arg_val<uint32_t>(CARG_ACT_F32);
+  const uint32_t mtile = get_common_arg_val<uint32_t>(CARG_MTILE);
+  const uint32_t qb_pad = get_common_arg_val<uint32_t>(CARG_QB_PAD);
+  const uint32_t enc = get_common_arg_val<uint32_t>(CARG_ENC);
+  const uint32_t tcols = get_common_arg_val<uint32_t>(CARG_TCOLS);  // padded tile width
+  const uint32_t grid_x = get_common_arg_val<uint32_t>(CARG_GRID_X);
+  // The per-core slice, derived from the core coordinate instead of a
+  // per-core SetRuntimeArgs word (the 27B trace-fit fix): c enumerates the
+  // grid row-major, exactly the host's {c % grid.x, c / grid.x} mapping.
+  const uint32_t c =
+      static_cast<uint32_t>(get_relative_logical_y()) * grid_x +
+      static_cast<uint32_t>(get_relative_logical_x());
+  const uint32_t row0 = c * tcols;
+  const uint32_t rowc =
+      row0 >= N ? 0u
+                : ((tcols < N - row0) ? tcols : (N - row0));
   if (rowc == 0 || M == 0) return;
 
   const uint32_t word_bytes = wpb * 4;
@@ -1988,6 +2008,18 @@ void MatmulBTQuantInt8DotKernel(Queue& q, Tensor& out, const Tensor& a,
       std::to_string(grid.y);
 
   std::lock_guard<std::mutex> workload_guard(Int8DotWorkloadMutex());
+  // The ONE common-args vector the program runs (see the kernel's CARG_*
+  // table): the three bank bases followed by the shape-global words. Built per
+  // call, set once on a miss and updated in place on a hit.
+  const auto common_args = [&] {
+    return std::vector<uint32_t>{
+        static_cast<uint32_t>(words.mesh_buffer().address()),
+        static_cast<uint32_t>(dev_a.mesh_buffer().address()),
+        static_cast<uint32_t>(dev_out.mesh_buffer().address()),
+        static_cast<uint32_t>(M), static_cast<uint32_t>(K),
+        static_cast<uint32_t>(N), static_cast<uint32_t>(nb), wpb, act_f32,
+        mtile, qb_pad, enc_sel, tcols, static_cast<uint32_t>(grid.x)};
+  };
   auto& workload_cache = Int8DotWorkloadCache();
   auto workload_it = workload_cache.find(workload_key);
   const bool workload_miss = workload_it == workload_cache.end();
@@ -2070,12 +2102,12 @@ void MatmulBTQuantInt8DotKernel(Queue& q, Tensor& out, const Tensor& a,
             .opt_level = tt::tt_metal::KernelBuildOptLevel::O2,
             .compiler_include_paths = {KeepQuantKernelIncludeDir()}});
     // The ONE legal initial common-args set (kernel.cpp:786: common runtime
-    // args can only be set once; later calls update them in place).
-    tt::tt_metal::SetCommonRuntimeArgs(
-        program, kernel,
-        {static_cast<uint32_t>(words.mesh_buffer().address()),
-         static_cast<uint32_t>(dev_a.mesh_buffer().address()),
-         static_cast<uint32_t>(dev_out.mesh_buffer().address())});
+    // args can only be set once; later calls update them in place). ALL the
+    // runtime words live here since the 27B trace-fit fix: per-core
+    // SetRuntimeArgs recorded one Unique RTA UNICAST page per core (~2.97 MB
+    // per captured launch), and the two words that varied per core (row0,
+    // rowc) are derived in-kernel from the core coordinate.
+    tt::tt_metal::SetCommonRuntimeArgs(program, kernel, common_args());
 
     tt::tt_metal::distributed::MeshWorkload workload;
     workload.add_program(
@@ -2092,42 +2124,26 @@ void MatmulBTQuantInt8DotKernel(Queue& q, Tensor& out, const Tensor& a,
   // reaches its cached program the same way (workload.get_programs(),
   // device_operation.hpp:184). Per-call runtime args on it (the ttnn
   // override_runtime_arguments contract): the common args carry this call's
-  // buffer addresses; the per-core args the shape, encoding and column-slice
-  // words. Dispatch commands regenerate from these on every enqueue
-  // (mesh_workload.cpp:210), so a re-enqueued workload always runs this
-  // call's values.
+  // buffer addresses plus the shape/encoding words. Dispatch commands
+  // regenerate from these on every enqueue (mesh_workload.cpp:210), so a
+  // re-enqueued workload always runs this call's values. NO per-core
+  // SetRuntimeArgs remains on this program: the 27B trace-fit fix derives
+  // row0/rowc in-kernel from the core coordinate (grid is part of the
+  // workload key, so grid_x is shape-global per program), and every other
+  // word is identical across cores. This call is the only call site of the
+  // program — nothing else needs per-core args on it.
   tt::tt_metal::Program& program =
       workload_it->second.workload.get_programs().begin()->second;
   if (!workload_miss) {
     // Update the common args IN PLACE on a reused program (kernel.cpp:786
-    // forbids a second set). The three words are this call's bank bases: the
-    // words shadow, the activation, the out page — the GetCommonRuntimeArgs
-    // pattern (ttnn unary_program_factory.cpp:647-652). A miss just set them
-    // with this call's addresses.
-    auto& common_args =
+    // forbids a second set) — the GetCommonRuntimeArgs pattern (ttnn
+    // unary_program_factory.cpp:647-652). A miss just set them with this
+    // call's values.
+    auto& crta =
         tt::tt_metal::GetCommonRuntimeArgs(program, workload_it->second.kernel);
-    common_args[0] = static_cast<uint32_t>(words.mesh_buffer().address());
-    common_args[1] = static_cast<uint32_t>(dev_a.mesh_buffer().address());
-    common_args[2] = static_cast<uint32_t>(dev_out.mesh_buffer().address());
-  }
-
-  std::vector<tt::tt_metal::CoreCoord> core_coords;
-  std::vector<std::vector<uint32_t>> per_core;
-  core_coords.reserve(grid_cores);
-  per_core.reserve(grid_cores);
-  for (uint32_t c = 0; c < grid_cores; ++c) {
-    const uint32_t r0 = c * tcols;
-    const uint32_t rc =
-        r0 >= static_cast<uint32_t>(N)
-            ? 0u
-            : std::min(tcols, static_cast<uint32_t>(N) - r0);
-    core_coords.push_back(tt::tt_metal::CoreCoord{c % grid.x, c / grid.x});
-    per_core.push_back({static_cast<uint32_t>(M), static_cast<uint32_t>(K),
-                        static_cast<uint32_t>(N), static_cast<uint32_t>(nb),
-                        wpb, act_f32, r0, rc, mtile, qb_pad, enc_sel, tcols});
+    const std::vector<uint32_t> next = common_args();
+    for (uint32_t i = 0; i < next.size(); ++i) crta[i] = next[i];
   }
-  tt::tt_metal::SetRuntimeArgs(program, workload_it->second.kernel,
-                               core_coords, per_core);
   // A fresh runtime id before EVERY enqueue, hit or miss — what the ttnn
   // dispatch does unconditionally (device_operation.hpp:181-186).
   program.set_runtime_id(static_cast<uint64_t>(
diff --git a/src/vt/tenstorrent/tenstorrent_ops.cpp b/src/vt/tenstorrent/tenstorrent_ops.cpp
index b014fa824..522e8d6cd 100644
--- a/src/vt/tenstorrent/tenstorrent_ops.cpp
+++ b/src/vt/tenstorrent/tenstorrent_ops.cpp
@@ -104,6 +104,56 @@ void MatmulBTKernel(Queue&, Tensor& out, const Tensor& a, const Tensor& b) {
 // [rows, d] tile rather than relying on ttnn's own broadcast rules — keeps
 // this kernel's behavior pinned to the CPU reference rather than to
 // whatever ttnn::add happens to support today.
+// tt-27b-region-capture: the broadcast `b` upload must be warmable. The eager
+// pass uploads the replicated [rows, d] tensor once and caches it keyed by the
+// host pointer, the geometry, and a hash of `b`'s d values, so the capture
+// pass finds it resident (a changed host value hashes differently and is
+// re-uploaded in the next eager pass — a capture-scope miss is a warm hole and
+// is refused by name). Before this cache the upload ran unconditionally in
+// BOTH passes: every captured broadcast Add inlined a rows*d*4 B payload into
+// the trace (the inline-capture-scope-upload class the
+// tt-trace-record-audit-20260928 audit pinned at ~3 MB per command).
+ttnn::Tensor BroadcastOperandDevice(const Tensor& b, uint32_t rows, uint32_t d,
+                                    MeshDevice& device) {
+  struct Entry {
+    uint32_t rows, d;
+    uint64_t hash;
+    ttnn::Tensor dev;
+  };
+  static std::unordered_map<const void*, Entry> cache;
+  static std::mutex mutex;
+  uint64_t h = 1469598103934665603ull;
+  for (uint32_t i = 0; i < d; ++i) {
+    float v = LoadElemF32(b, i);
+    uint32_t bits;
+    std::memcpy(&bits, &v, sizeof(bits));
+    h = (h ^ bits) * 1099511628211ull;
+  }
+  if (tt_capture_active()) {
+    std::lock_guard<std::mutex> g(mutex);
+    auto it = cache.find(b.data);
+    if (it != cache.end() && it->second.rows == rows && it->second.d == d &&
+        it->second.hash == h)
+      return it->second.dev;
+    VT_CHECK(false,
+             "tenstorrent: AddKernel broadcast replicated-tensor upload "
+             "refused inside an open trace capture — no resident copy for "
+             "this operand at TraceBeginCapture; warm the add in the eager "
+             "pass (tt-27b-region-capture)");
+  }
+  std::vector<float> replicated(static_cast<size_t>(rows) * d);
+  for (uint32_t r = 0; r < rows; ++r)
+    for (uint32_t c = 0; c < d; ++c)
+      replicated[static_cast<size_t>(r) * d + c] = LoadElemF32(b, c);
+  ttnn::Tensor dev =
+      ttnn::Tensor::from_vector<float>(replicated, TileSpecOf(rows, d), &device);
+  {
+    std::lock_guard<std::mutex> g(mutex);
+    cache[b.data] = Entry{rows, d, h, dev};
+  }
+  return dev;
+}
+
 void AddKernel(Queue&, Tensor& out, const Tensor& a, const Tensor& b) {
   TT_OP_TRACE("Add");
   VT_CHECK(a.rank == 2 && out.rank == 2, "tenstorrent kAdd: `a`/`out` must be rank-2 in W0");
@@ -125,13 +175,9 @@ void AddKernel(Queue&, Tensor& out, const Tensor& a, const Tensor& b) {
   ttnn::Tensor dev_b;
   if (bcast) {
     EnsureHost(b);
-    std::vector<float> replicated(static_cast<size_t>(rows) * d);
-    for (uint32_t r = 0; r < rows; ++r)
-      for (uint32_t c = 0; c < d; ++c)
-        replicated[static_cast<size_t>(r) * d + c] = LoadElemF32(b, c);
     if (std::getenv("VT_TT_TRACE_DEBUG") != nullptr && tt_capture_active())
-      std::fprintf(stderr, "[TT-UP] AddKernel from_vector WRITE during capture\n");
-    dev_b = ttnn::Tensor::from_vector<float>(replicated, TileSpecOf(rows, d), &device);
+      std::fprintf(stderr, "[TT-UP] AddKernel broadcast operand served under capture\n");
+    dev_b = BroadcastOperandDevice(b, rows, d, device);
   } else {
     dev_b = EnsureDevice2D(b, device);
   }
diff --git a/src/vt/tenstorrent/tenstorrent_paged.cpp b/src/vt/tenstorrent/tenstorrent_paged.cpp
index 86b9ffffd..4adac4bc0 100644
--- a/src/vt/tenstorrent/tenstorrent_paged.cpp
+++ b/src/vt/tenstorrent/tenstorrent_paged.cpp
@@ -608,6 +608,31 @@ struct RacIdxEntry {
   // read and may hold garbage — no zeros tail, no concat, no allocation.
   ttnn::Tensor sharded_in;   // K input (height-sharded)
   ttnn::Tensor sharded_in_v;  // V input (separate — K and V must NOT share the same buffer)
+  // Batched lane (num_slots > 1): the PROVEN single-user tensors, one set per
+  // user. sharded_*[u] sits on its own core (K on worker u, V on worker C+u);
+  // update_idxs[u] is [1], page_table[u] is [1, cols] — each fused-update
+  // call then runs the C=1 shapes, with the op's override_runtime_arguments
+  // re-patching the per-user addresses on the shared cached program. The
+  // per-user idx content is refreshed OUTSIDE capture every step (no on-device
+  // plus_one for this lane yet — recorded as owed: fold the [C] plus_one'd
+  // cur_pos into the per-user reads).
+  std::vector<ttnn::Tensor> batched_in;
+  std::vector<ttnn::Tensor> batched_in_v;
+  std::vector<ttnn::Tensor> batched_update_idxs;
+  std::vector<ttnn::Tensor> batched_page_table;
+  std::vector<int32_t> batched_idx_host;   // last content copied per user
+  std::vector<int32_t> batched_pt_host;    // last page-table row copied per user
+  int64_t batched_pt_width = 0;  // columns the batched page tables were built with
+  // A page-table WIDTH change (block boundary growth, or the shrink when the
+  // longest request finishes) retires the per-user tables here and
+  // reallocates — the C=1 lane's pt_width discipline. A stale-width device
+  // tensor would TT_FATAL the refresh copy_to_device (shape mismatch), and an
+  // old-width batched_pt_host makes the change-detection loop read out of
+  // bounds. The retired tensors stay alive: never free a buffer a recorded
+  // trace addresses (#1105).
+  std::vector<ttnn::Tensor> batched_retired_pts;
+  bool batched_alloc = false;
+  bool batched_in_is_alloc = false;
   uint32_t nkv = 0;
   uint32_t d = 0;
   bool allocated = false;        // ttnn::Tensor::is_allocated() crashes on default-constructed tensors in this build
@@ -639,7 +664,9 @@ bool TryReshapeAndCacheDeviceDecode(const Tensor& k, const Tensor& v,
   const int64_t num_slots = slot_mapping.shape[0];
   if (T < 1 || num_slots < 1) return false;
   if ((d % 32u) != 0u || (bs % 32u) != 0u) return false;
-  if (num_slots > 1) return false;  // decode T=1 only for now
+  // Decode: one token per user (token i is user i). Prefill (T > num_slots)
+  // keeps the host path.
+  if (T != num_slots) return false;
 
   // k/v must carry CURRENT device shadows ([T*nkv, d] TILE bf16 from rope).
   std::optional<ttnn::Tensor> k_dev, v_dev;
@@ -678,14 +705,23 @@ bool TryReshapeAndCacheDeviceDecode(const Tensor& k, const Tensor& v,
     }
   }
 
-  // Paged-KV shadows must exist and cover the target block.
-  const int64_t slot = slot_mapping.Ptr<int64_t>()[0];
+  // Walk ALL users: batched decode carries one slot per user. The paged-KV
+  // shadow must cover the DEEPEST target block; per-user padding slots
+  // (slot < 0) are skipped by paged_update_cache itself (update_idx == -1).
+  const int64_t* slots_ptr = slot_mapping.Ptr<int64_t>();
+  int64_t max_block = -1;
+  bool any_valid = false;
+  for (int64_t u = 0; u < num_slots; ++u) {
+    const int64_t su = slots_ptr[u];
+    if (su < 0) continue;
+    max_block = std::max(max_block, su / bs);
+    any_valid = true;
+  }
   if (std::getenv("VT_TT_TRACE_DEBUG") != nullptr)
-    std::fprintf(stderr, "[TT-TRACE] RAC slot=%lld cap=%d\n",
-                 (long long)slot, (int)tt_capture_active());
-  if (slot < 0) return true;  // nothing to write; treat as handled
-  const uint32_t block = static_cast<uint32_t>(slot / bs);
-  const uint32_t offset = static_cast<uint32_t>(slot % bs);
+    std::fprintf(stderr, "[TT-TRACE] RAC slot0=%lld nslots=%lld max_block=%lld cap=%d\n",
+                 (long long)slots_ptr[0], (long long)num_slots,
+                 (long long)max_block, (int)tt_capture_active());
+  if (!any_valid) return true;  // nothing to write; treat as handled
 
   std::optional<ttnn::Tensor> kc_dev, vc_dev;
   {
@@ -696,9 +732,9 @@ bool TryReshapeAndCacheDeviceDecode(const Tensor& k, const Tensor& v,
       std::fprintf(stderr, "[TT-TRACE] RAC paged-kv shadow k=%d v=%d k_nb=%u\n",
                    skc->device.has_value(), svc->device.has_value(), skc->nb);
     if (!skc->device.has_value() || !svc->device.has_value()) return false;
-    if (skc->nb <= block || skc->nkv != static_cast<uint32_t>(nkv) ||
+    if (skc->nb <= max_block || skc->nkv != static_cast<uint32_t>(nkv) ||
         skc->bs != static_cast<uint32_t>(bs) || skc->d != static_cast<uint32_t>(d)) return false;
-    if (svc->nb <= block || svc->nkv != static_cast<uint32_t>(nkv) ||
+    if (svc->nb <= max_block || svc->nkv != static_cast<uint32_t>(nkv) ||
         svc->bs != static_cast<uint32_t>(bs) || svc->d != static_cast<uint32_t>(d)) return false;
     kc_dev = skc->device;
     vc_dev = svc->device;
@@ -717,14 +753,21 @@ bool TryReshapeAndCacheDeviceDecode(const Tensor& k, const Tensor& v,
                    (long long)num_slots, (long long)bs);
     // WarmRacIdx (driver Refresh slot) refreshes update_idxs/page_table content
     // every step via copy_to_device; here we just verify the tensors exist.
-    if (it == RacIdxCache().end() || !it->second.allocated) {
+    // C=1 verifies the shared pair; the batched lane verifies its per-user set.
+    const bool warmed = num_slots == 1
+                            ? (it != RacIdxCache().end() && it->second.allocated)
+                            : (it != RacIdxCache().end() &&
+                               it->second.batched_alloc &&
+                               it->second.batched_in_is_alloc &&
+                               it->second.batched_in.size() ==
+                                   static_cast<size_t>(num_slots));
+    if (!warmed) {
       VT_CHECK(!tt_capture_active(),
                "tenstorrent: RAC idx tensors not warmed — call WarmRacIdx "
                "outside capture (driver Refresh slot) first");
       return false;
     }
-    // sharded_in must exist (WarmRacIdx needs the paged-KV shadow geometry).
-    if (!it->second.sharded_in_is_alloc) {
+    if (num_slots == 1 && !it->second.sharded_in_is_alloc) {
       VT_CHECK(!tt_capture_active(),
                "tenstorrent: RAC sharded input not warmed — call WarmRacIdx "
                "outside capture after WarmPagedKvShadow");
@@ -827,34 +870,129 @@ bool TryReshapeAndCacheDeviceDecode(const Tensor& k, const Tensor& v,
     return sharded_dst;
   };
   // V first, then K
-  // Debug: dump v_dev properties before sharding
-  ttnn::Tensor v_in = build_input(*v_dev, rac_entry.sharded_in_v, rac_entry);
-  ttnn::Tensor k_in = build_input(*k_dev, rac_entry.sharded_in, rac_entry);
-  // Debug: check v_in for all heads
-  // num_kv_heads_override pins the kernel's head loop to nkv rows: the input
-  // shard is tile-padded (nkv_pad rows) but only the first nkv rows hold data
-  // (upstream decode pattern, test_paged_cache_flexible_geometry.py).
-  // Use paged_fused_update_cache (single call for K+V) instead of two separate
-  // paged_update_cache calls. The fused op has override_runtime_arguments
-  // (the non-fused doesn't), so it works correctly with program cache enabled.
-  // The second separate call would reuse the first's cached program with the
-  // first's buffer addresses (program cache collision).
-  auto [new_kc, new_vc] = ttnn::experimental::paged_fused_update_cache(
-      *kc_dev, k_in, *vc_dev, v_in,
-      /*update_idxs=*/{}, rac_entry.update_idxs,
-      /*share_cache=*/false, rac_entry.page_table,
-      /*batch_offset=*/0, /*compute_kernel_config=*/std::nullopt,
-      /*mesh_coords=*/std::nullopt);
+  // Batched users (num_slots > 1): run the PROVEN single-user sequence once
+  // per user — slice this user's [nkv, d] rows out of the rope shadow,
+  // materialize the fresh native [1,1,nkv,d] via the 1.0 multiply, ttnn::copy
+  // into that user's own single-shard persistent input, then one
+  // paged_fused_update_cache call against that user's own [1] update_idx and
+  // [1, cols] page-table row. Same shapes as the C=1 lane, so call u > 0
+  // hits the warmed program and the fused op's override_runtime_arguments
+  // re-patches the per-user buffer addresses — the exact mechanism the
+  // C=1 replay already relies on. slice/multiply/copy are the warmed,
+  // capture-safe in-region ops.
+  auto user_slice = [&](const ttnn::Tensor& src, uint32_t u,
+                        const RacIdxEntry& entry) -> ttnn::Tensor {
+      // Materialize a FRESH native copy of the whole shadow first (the 1.0
+      // multiply), THEN slice the user's rows out of it. Slicing the
+      // committed rope shadow directly mis-served the second user's last
+      // head in-suite (the unaligned TILE row slice reused a slice program
+      // first compiled against a different offset); from a fresh native
+      // tensor the slice is exact in both regimes.
+      const auto ls = src.logical_shape();
+      ttnn::Tensor fresh = ttnn::multiply(src, 1.0f);
+      // Rank-2 token-row layout [T, nkv*d] (rope's flat commit on the 27B):
+      // user u is ROW u; each head is a d-ALIGNED column range, so slice
+      // per (u, h) column range — the proven per-head recipe — and concat
+      // the heads back into [nkv, d] storage.
+      if (ls.rank() == 2 && ls[0] == num_slots && ls[1] == entry.nkv * entry.d) {
+        std::vector<ttnn::Tensor> heads;
+        heads.reserve(entry.nkv);
+        for (uint32_t h = 0; h < entry.nkv; ++h) {
+          heads.push_back(ttnn::slice(
+              fresh,
+              ttsl::SmallVector<uint32_t>{u, h * entry.d},
+              ttsl::SmallVector<uint32_t>{u + 1u, (h + 1u) * entry.d},
+              ttsl::SmallVector<uint32_t>{1u, 1u}));
+        }
+        return ttnn::concat(heads, /*dim=*/0);
+      }
+      return ls.rank() == 3
+                 ? ttnn::slice(fresh,
+                               ttsl::SmallVector<uint32_t>{u, 0u, 0u},
+                               ttsl::SmallVector<uint32_t>{u + 1u, entry.nkv,
+                                                           entry.d},
+                               ttsl::SmallVector<uint32_t>{1u, 1u, 1u})
+                 : ttnn::slice(fresh,
+                               ttsl::SmallVector<uint32_t>{u * entry.nkv, 0u},
+                               ttsl::SmallVector<uint32_t>{
+                                   (u + 1u) * entry.nkv, entry.d},
+                               ttsl::SmallVector<uint32_t>{1u, 1u});
+  };
+  const uint32_t C = static_cast<uint32_t>(num_slots);
+  // Per-user temporaries (the fresh shadow copy, the slice, the native4
+  // multiply output) are FRESH ALLOCATIONS freed at each iteration's end.
+  // Freeing one while its enqueued copy is still in flight lets the next
+  // iteration's allocation recycle the storage under the deferred reader —
+  // the exact class the retention root-cause names — so hold them until the
+  // end of the function.
+  std::vector<ttnn::Tensor> keepalive;
+  keepalive.reserve(static_cast<size_t>(C) * 3);
+  std::optional<ttnn::Tensor> new_kc, new_vc;
+  if (num_slots == 1) {
+    // ISSUE-LOCAL-01M3M0K390EM40W5R9BR5A2KZ7: THE C=1 LANE, verbatim the
+    // pre-e39f2cf3f form (git `show e39f2cf3f~1` carries it). The batched
+    // rewrite routed ONE user through the batched arrays WarmRacIdx never
+    // allocates for C=1 — the loop indexed empty vectors and the first cold
+    // decode step segfaulted on ttnn::copy into an empty tensor. C=1 feeds
+    // the WHOLE rope shadow (no per-user slice: there is one user) into the
+    // SHARED sharded inputs, and one fused update against the SHARED
+    // (plus_one'd on-device, R2) update_idxs and page-table. V first, then K.
+    // num_kv_heads_override pins the kernel's head loop to nkv rows: the
+    // input shard is tile-padded (nkv_pad rows) but only the first nkv rows
+    // hold data (upstream decode pattern,
+    // test_paged_cache_flexible_geometry.py). The fused op has
+    // override_runtime_arguments (the non-fused doesn't), so it works with
+    // the program cache enabled; two separate calls would reuse the first's
+    // cached program with the first's buffer addresses.
+    ttnn::Tensor v_in = build_input(*v_dev, rac_entry.sharded_in_v, rac_entry);
+    ttnn::Tensor k_in = build_input(*k_dev, rac_entry.sharded_in, rac_entry);
+    auto [nkc, nvc] = ttnn::experimental::paged_fused_update_cache(
+        *kc_dev, k_in, *vc_dev, v_in,
+        /*update_idxs=*/{}, rac_entry.update_idxs,
+        /*share_cache=*/false, rac_entry.page_table,
+        /*batch_offset=*/0, /*compute_kernel_config=*/std::nullopt,
+        /*mesh_coords=*/std::nullopt);
+    new_kc = std::move(nkc);
+    new_vc = std::move(nvc);
+  } else {
+  for (uint32_t u = 0; u < C; ++u) {
+    ttnn::Tensor v_src = user_slice(*v_dev, u, rac_entry);
+    ttnn::Tensor k_src = user_slice(*k_dev, u, rac_entry);
+    ttnn::Tensor v_in = build_input(v_src, rac_entry.batched_in_v[u], rac_entry);
+    ttnn::Tensor k_in = build_input(k_src, rac_entry.batched_in[u], rac_entry);
+    keepalive.push_back(v_src);
+    keepalive.push_back(k_src);
+    keepalive.push_back(v_in);
+    keepalive.push_back(k_in);
+    // num_kv_heads_override pins the kernel's head loop to nkv rows (the
+    // input shard is tile-padded); see the C=1 lane below for the fused-op /
+    // program-cache rationale.
+    auto [ukc, uvc] = ttnn::experimental::paged_fused_update_cache(
+        *kc_dev, k_in, *vc_dev, v_in,
+        /*update_idxs=*/{}, rac_entry.batched_update_idxs[u],
+        /*share_cache=*/false, rac_entry.batched_page_table[u],
+        /*batch_offset=*/0, /*compute_kernel_config=*/std::nullopt,
+        /*mesh_coords=*/std::nullopt);
+    new_kc = std::move(ukc);
+    new_vc = std::move(uvc);
+    // Eager pass only: sync between users, so user u+1's fresh allocations
+    // cannot recycle user u's in-flight temporaries (the deferred-reader
+    // recycle class the retention root-cause names) and each per-user program
+    // variant is patched against quiesced state. A finish inside a live trace
+    // is illegal; the capture pass records the identical op sequence against
+    // trace-tracked allocations instead.
+    if (!tt_capture_active()) SharedMeshDevice().mesh_command_queue().finish();
+  }
+  }  // num_slots == 1 (shared C=1 tensors) / the per-user batched lane
   {
     std::lock_guard<std::mutex> g(PagedKvMutex());
-    PagedKvShadows()[reinterpret_cast<uintptr_t>(k_cache.data)].device = std::move(new_kc);
+    PagedKvShadows()[reinterpret_cast<uintptr_t>(k_cache.data)].device = std::move(*new_kc);
     PagedKvShadows()[reinterpret_cast<uintptr_t>(k_cache.data)].device_current = true;
-    PagedKvShadows()[reinterpret_cast<uintptr_t>(v_cache.data)].device = std::move(new_vc);
+    PagedKvShadows()[reinterpret_cast<uintptr_t>(v_cache.data)].device = std::move(*new_vc);
     PagedKvShadows()[reinterpret_cast<uintptr_t>(v_cache.data)].device_current = true;
   }
   if (std::getenv("VT_TT_TRACE_DEBUG") != nullptr)
-    std::fprintf(stderr, "[TT-TRACE] RAC device update (copy+paged_update_cache)\n");
-  (void)offset; (void)block;
+    std::fprintf(stderr, "[TT-TRACE] RAC device update batched C=%u (per-user copy+paged_update_cache)\n", C);
   return true;
 }
 
@@ -1124,15 +1262,18 @@ bool TryPagedAttentionDeviceDecode(Tensor& out, const Tensor& query, const Tenso
           // path). No metadata view of the [rows, D] buffer represents
           // the per-batch head tiling at B > 1, so materialize the
           // correct [1, B, H, D] TILE tensor through the free reshape's
-          // device program. That program calls to_device — forbidden
-          // during trace capture — so a captured batched step takes the
-          // host Q path (the contract the old fatal enforced, loudly).
+          // device program. W4 doctrine: BOTH passes run that same
+          // multiply(reshape(...)) chain — the eager step warms the
+          // reshape program for this exact input/output spec, so the
+          // captured call is a program-cache HIT, not a mid-trace
+          // to_device. The former capture-active B>1 decline here sent
+          // every captured batched step to the host Q arm, whose refusal
+          // cascaded into PagedAttentionKernel's host oracle and its
+          // EnsureHost(k_cache) readback mid-trace — the 27B c2 leg fatal
+          // (ISSUE-LOCAL-01M3JXEFQKSZP23PP2HWY9G0VQ). A reshape spec the
+          // warmup did not warm still fatals loudly at the program-cache
+          // miss, which is the W4 divergence detector, not a defect.
           const auto ps2d = dev_q_2d.padded_shape();
-          (void)ps2d;
-          if (tt_capture_active() && Bu > 1)
-            throw std::runtime_error(
-                "tenstorrent PA: batched (B>1) Q 4D materialization is "
-                "not capture-safe; the host Q path must serve this step");
           if (Bu == 1) {
             dev_q = ttnn::multiply(
                 ttnn::experimental::view(
@@ -1987,6 +2128,16 @@ void WarmPagedKvShadow(void* k_cache_data, void* v_cache_data,
   warm_one(v_cache_data);
 }
 
+bool ReadPagedKvShadowForTest(const void* k_cache_data, float* dst, int64_t n) {
+  std::lock_guard<std::mutex> g(PagedKvMutex());
+  auto it = PagedKvShadows().find(reinterpret_cast<uintptr_t>(k_cache_data));
+  if (it == PagedKvShadows().end() || !it->second.device.has_value()) return false;
+  auto host = it->second.device->to_vector<float>();
+  const int64_t m = std::min<int64_t>(n, static_cast<int64_t>(host.size()));
+  for (int64_t i = 0; i < m; ++i) dst[i] = host[static_cast<size_t>(i)];
+  return true;
+}
+
 void WarmRacIdx(const void* /*slot_mapping_owner*/, const int64_t* slots,
                 int64_t num_slots, int64_t block_size,
                 const int32_t* block_table, int64_t block_table_cols,
@@ -2050,6 +2201,13 @@ void WarmRacIdx(const void* /*slot_mapping_owner*/, const int64_t* slots,
   const auto key = std::make_pair(num_slots, block_size);
   std::lock_guard<std::mutex> g(RacIdxMutex());
   RacIdxEntry& e = RacIdxCache()[key];
+  // The SHARED [C] update_idxs / page-table tensors serve the C=1 lane only
+  // (its captured paged_update_cache replays against these stable addresses,
+  // and its update_idxs is plus_one'd on-device). The batched lane keeps its
+  // own per-user tensors below and must NOT allocate the shared ones — a
+  // standalone update_idxs allocated after a capture would trip the #1105
+  // frozen-index refusal. The vectors ptv/idxv above still feed it.
+  if (num_slots == 1) {
   // idx/page-table tensors are allocated ONCE per key and their CONTENT is
   // refreshed in place each step (copy_to_device, outside capture). The
   // captured paged_update_cache replays against the stable address and reads
@@ -2113,6 +2271,135 @@ void WarmRacIdx(const void* /*slot_mapping_owner*/, const int64_t* slots,
     }
   }
   e.pt_host = ptv;
+  }  // num_slots == 1 (shared C=1 tensors)
+  // Batched lane (num_slots > 1): per-user persistent tensors, allocated
+  // once, CONTENT refreshed outside capture every step. The per-user
+  // update_idx changes every step (it is the decode position), so unlike the
+  // C=1 lane there is no on-device plus_one yet — the refresh is the same
+  // copy_to_device discipline the page-table refresh already pays (the
+  // on-device advance for this lane is recorded as owed in the entry).
+  if (num_slots > 1) {
+    const uint32_t C = static_cast<uint32_t>(num_slots);
+    // ANY width change retires + reallocates the per-user page tables (same
+    // discipline as the C=1 lane above): the refresh copy_to_device would
+    // TT_FATAL on a shape mismatch, and batched_pt_host sized at the old
+    // width makes the change-detection loop read out of bounds. The driver
+    // resets + re-captures on any column-count change, so the new address is
+    // what the next capture records.
+    bool realloc_pts = false;
+    if (e.batched_alloc && block_table_cols != e.batched_pt_width) {
+      for (auto& pt : e.batched_page_table)
+        e.batched_retired_pts.push_back(std::move(pt));
+      e.batched_page_table.assign(C, ttnn::Tensor());
+      e.batched_pt_host.clear();
+      realloc_pts = true;
+    }
+    if (!e.batched_alloc) {
+      MeshDevice& device2 = device;
+      e.batched_in.assign(C, ttnn::Tensor());
+      e.batched_in_v.assign(C, ttnn::Tensor());
+      e.batched_update_idxs.assign(C, ttnn::Tensor());
+      e.batched_page_table.assign(C, ttnn::Tensor());
+      for (uint32_t u = 0; u < C; ++u) {
+        e.batched_update_idxs[u] = ttnn::Tensor::from_vector<int32_t>(
+            {idxv[static_cast<size_t>(u)]},
+            SpecOf(tt::tt_metal::Shape({1u}), ttnn::DataType::INT32,
+                              ttnn::Layout::ROW_MAJOR),
+            &device2);
+        e.batched_page_table[u] = ttnn::Tensor::from_vector<int32_t>(
+            std::vector<int32_t>(
+                ptv.begin() + static_cast<long>(u) * block_table_cols,
+                ptv.begin() + static_cast<long>(u + 1) * block_table_cols),
+            SpecOf(tt::tt_metal::Shape({1u, static_cast<uint32_t>(block_table_cols)}),
+                   ttnn::DataType::INT32, ttnn::Layout::ROW_MAJOR),
+            &device2);
+      }
+      e.batched_idx_host.assign(idxv.begin(), idxv.end());
+      e.batched_pt_host = ptv;
+      e.batched_pt_width = block_table_cols;
+      e.batched_alloc = true;
+    } else if (realloc_pts) {
+      // Width change on a live entry: rebuild ONLY the per-user page tables
+      // (update_idxs is [1] per user and width-independent; the sharded
+      // inputs must NOT be touched — batched_in_is_alloc still marks them).
+      for (uint32_t u = 0; u < C; ++u) {
+        e.batched_page_table[u] = ttnn::Tensor::from_vector<int32_t>(
+            std::vector<int32_t>(
+                ptv.begin() + static_cast<long>(u) * block_table_cols,
+                ptv.begin() + static_cast<long>(u + 1) * block_table_cols),
+            SpecOf(tt::tt_metal::Shape({1u, static_cast<uint32_t>(block_table_cols)}),
+                   ttnn::DataType::INT32, ttnn::Layout::ROW_MAJOR),
+            &device);
+      }
+      e.batched_pt_host = ptv;
+      e.batched_pt_width = block_table_cols;
+    } else {
+      for (uint32_t u = 0; u < C; ++u) {
+        const int32_t idx_u = idxv[static_cast<size_t>(u)];
+        if (e.batched_idx_host[static_cast<size_t>(u)] != idx_u) {
+          ttnn::Tensor ih = ttnn::Tensor::from_vector<int32_t>(
+              {idx_u}, SpecOf(tt::tt_metal::Shape({1u}), ttnn::DataType::INT32,
+                              ttnn::Layout::ROW_MAJOR));
+          ttnn::copy_to_device(ih, e.batched_update_idxs[u]);
+          e.batched_idx_host[static_cast<size_t>(u)] = idx_u;
+        }
+        bool row_changed = false;
+        for (int64_t c = 0; c < block_table_cols; ++c) {
+          if (e.batched_pt_host[static_cast<size_t>(u) * block_table_cols + c] !=
+              ptv[static_cast<size_t>(u) * block_table_cols + c]) {
+            row_changed = true;
+            break;
+          }
+        }
+        if (row_changed) {
+          ttnn::Tensor ph = ttnn::Tensor::from_vector<int32_t>(
+              std::vector<int32_t>(
+                  ptv.begin() + static_cast<long>(u) * block_table_cols,
+                  ptv.begin() + static_cast<long>(u + 1) * block_table_cols),
+              SpecOf(tt::tt_metal::Shape({1u, static_cast<uint32_t>(block_table_cols)}),
+                     ttnn::DataType::INT32, ttnn::Layout::ROW_MAJOR));
+          ttnn::copy_to_device(ph, e.batched_page_table[u]);
+          for (int64_t c = 0; c < block_table_cols; ++c)
+            e.batched_pt_host[static_cast<size_t>(u) * block_table_cols + c] =
+                ptv[static_cast<size_t>(u) * block_table_cols + c];
+        }
+      }
+    }
+    // Per-user single-shard inputs: K for user u on worker core (u, 0), V on
+    // (C+u, 0) — never the same core, or the second copy's program-cache hit
+    // would write V into K's buffer (the C=1 lane's rule, one pair per user).
+    if (!e.batched_in_is_alloc) {
+      std::lock_guard<std::mutex> pg(PagedKvMutex());
+      for (auto& [ptr, shadow] : PagedKvShadows()) {
+        if (!(shadow.nkv > 0 && shadow.d > 0 && (shadow.d % 32u) == 0u)) continue;
+        const uint32_t np = std::max(32u, ((shadow.nkv + 31u) / 32u) * 32u);
+        const auto grid = device.compute_with_storage_grid_size();
+        if (2u * C > grid.x) break;  // not enough worker cores: leave unallocated
+        for (uint32_t u = 0; u < C; ++u) {
+          auto make = [&](uint32_t cx) {
+            tt::tt_metal::CoreRangeSet cs({tt::tt_metal::CoreRange(
+                tt::tt_metal::CoreCoord(cx, 0), tt::tt_metal::CoreCoord(cx, 0))});
+            tt::tt_metal::ShardSpec ss(cs, {np, shadow.d},
+                                       tt::tt_metal::ShardOrientation::ROW_MAJOR);
+            tt::tt_metal::MemoryConfig sm(
+                tt::tt_metal::TensorMemoryLayout::HEIGHT_SHARDED,
+                tt::tt_metal::BufferType::L1, ss);
+            return ttnn::create_device_tensor(
+                tt::tt_metal::TensorSpec(
+                    tt::tt_metal::Shape({1u, 1u, shadow.nkv, shadow.d}),
+                    tt::tt_metal::TensorLayout(
+                        ttnn::DataType::BFLOAT16,
+                        tt::tt_metal::PageConfig(ttnn::Layout::TILE), sm)),
+                &device);
+          };
+          e.batched_in[u] = make(u);
+          e.batched_in_v[u] = make(C + u);
+        }
+        e.batched_in_is_alloc = true;
+        break;
+      }
+    }
+  }
   // Build the persistent sharded RAC input ONCE from the first available
   // paged-KV shadow's geometry (same nkv/d as the cache): logical
   // [1,1,nkv,d], padded [1,1,nkv_pad,d], HEIGHT_SHARDED L1, shard
diff --git a/src/vt/tenstorrent/tenstorrent_residency.cpp b/src/vt/tenstorrent/tenstorrent_residency.cpp
index 2dc0d2379..3c5cee5bc 100644
--- a/src/vt/tenstorrent/tenstorrent_residency.cpp
+++ b/src/vt/tenstorrent/tenstorrent_residency.cpp
@@ -328,9 +328,22 @@ ttnn::Tensor UploadRows(const float* data, uint32_t rows, uint32_t cols, MeshDev
   if (std::getenv("VT_TT_TRACE_DEBUG") != nullptr && tt_capture_active())
     std::fprintf(stderr, "[TT-UP] UploadRows ptr=%p rows=%u cols=%u\n",
                  static_cast<const void*>(data), rows, cols);
+  // tt-27b-region-capture: an H2D write issued while a trace capture is open
+  // is recorded INLINE into the trace buffer — the 27B whole-graph
+  // 3,153,969,152 B demand was ~1,037 such payloads, not tt-metal record
+  // overhead (docs/bench-evidence/tt-trace-record-audit-20260928.md). The
+  // eager pass warms every upload; an upload that still fires under capture is
+  // a warm hole and is refused by name.
+  if (tt_capture_active()) {
+    if (std::getenv("VT_TT_TRACE_DEBUG") != nullptr)
+      std::fprintf(stderr, "[TT-UP] UploadRows from_vector WRITE during capture\n");
+    VT_CHECK(false,
+             "tenstorrent: UploadRows f32 H2D upload refused inside an open "
+             "trace capture — the payload would be recorded inline into the "
+             "trace; warm the tensor in the eager pass before "
+             "TraceBeginCapture (tt-27b-region-capture)");
+  }
   std::vector<float> host(data, data + static_cast<size_t>(rows) * cols);
-  if (std::getenv("VT_TT_TRACE_DEBUG") != nullptr && tt_capture_active())
-    std::fprintf(stderr, "[TT-UP] UploadRows from_vector WRITE during capture\n");
   return ttnn::Tensor::from_vector<float>(host, TileSpecOf(rows, cols), &device);
 }
 
@@ -410,6 +423,15 @@ ttnn::Tensor UploadRowsBf16(const Tensor& t, uint32_t rows, uint32_t cols,
   if (std::getenv("VT_TT_TRACE_DEBUG") != nullptr && tt_capture_active())
     std::fprintf(stderr, "[TT-UP] UploadRowsBf16 from_span WRITE during capture ptr=%p rows=%u cols=%u\n",
                  (const void*)t.data, rows, cols);
+  // tt-27b-region-capture: same refusal as UploadRows — every arm below
+  // (from_span, persistent enqueue_write, allocating arm) is an H2D write that
+  // a capture would record inline.
+  if (tt_capture_active())
+    VT_CHECK(false,
+             "tenstorrent: UploadRowsBf16 bf16 H2D upload refused inside an "
+             "open trace capture — the payload would be recorded inline into "
+             "the trace; warm the tensor in the eager pass before "
+             "TraceBeginCapture (tt-27b-region-capture)");
   const size_t n = static_cast<size_t>(rows) * static_cast<size_t>(cols);
   // The bytes at t.Ptr are the window's own bf16 bits (bfloat16 is a 2-byte
   // class wrapping the same uint16 pattern).
@@ -707,10 +729,15 @@ ttnn::Tensor EnsureDevice2D(const Tensor& t, MeshDevice& device) {
         const auto ls = s->device->logical_shape();
         if (ls.rank() == 2 && ls[0] == rows && ls[1] == cols)
           return *s->device;
-        if (tt_capture_active()) {
-          s->device = CaptureSafeReshape(*s->device, ttnn::Shape({rows, cols}));
-          return *s->device;
-        }
+        // ONE chain in BOTH passes (W4 doctrine; ISSUE-LOCAL-01M3JXEFQKSZP2
+        // 3PP2HWY9G0VQ). The tt_capture_active() arm this replaced ran a bare
+        // ttnn::reshape on the TILED shadow, a program the eager pass never
+        // warmed (its row-major reshape is a free view), so the first decode
+        // capture on a rank-3-committed slot (CommitDeviceLogical2D, e.g. the
+        // 27B attention output) created ReshapeViewTiledProgramFactory's
+        // program mid-capture and died on its to_device write. The chain
+        // below is warmed by the eager pass at this exact spec, so under
+        // capture it is a program-cache hit with no writes.
         ttnn::Tensor reshaped = ttnn::to_layout(
             ttnn::reshape(ttnn::to_layout(*s->device, ttnn::Layout::ROW_MAJOR),
                           ttnn::Shape({rows, cols})),
@@ -722,6 +749,11 @@ ttnn::Tensor EnsureDevice2D(const Tensor& t, MeshDevice& device) {
           static_cast<uint64_t>(s->dev_rows) * static_cast<uint64_t>(s->dev_cols);
       const uint64_t want = static_cast<uint64_t>(rows) * static_cast<uint64_t>(cols);
       if (have == want) {
+        if (std::getenv("VT_TT_TRACE_DEBUG") != nullptr)
+          std::fprintf(stderr,
+                       "[TT-RESHAPE] arm726 rows=%u cols=%u dev=%ux%u cap=%d\n",
+                       rows, cols, s->dev_rows, s->dev_cols,
+                       (int)tt_capture_active());
         if (tt_capture_active()) {
           ttnn::Tensor reshaped =
               CaptureSafeReshape(*s->device, ttnn::Shape({rows, cols}));
@@ -856,6 +888,26 @@ bool DeviceShadowExact(const Tensor& t, uint32_t rows, uint32_t cols) {
          s->dev_rows == rows && s->dev_cols == cols;
 }
 
+// ISSUE-LOCAL-01M3JXEFQKSZP23PP2HWY9G0VQ test surfaces (external linkage, the
+// DeviceShadowExact pattern — the test TU stays free of ttnn headers).
+// CommitRank3DeviceLogicalForTest reproduces the slot state a decode op
+// leaves after CommitDeviceLogical2D on a rank-3 device result: the shadow's
+// logical shape stays [b, h, d] while the slot records the flat 2D geometry
+// [b*h, d]. The next EnsureDevice2D at [b*h, d] takes the exact-rows/cols arm
+// whose logical shape mismatches — the arm that must run ONE chain in both
+// passes. EnsureDevice2DForTest is that call.
+void CommitRank3DeviceLogicalForTest(Tensor& out, uint32_t b, uint32_t h, uint32_t d) {
+  VT_CHECK(out.IsContiguous(), "CommitRank3DeviceLogicalForTest expects contiguous out");
+  VT_CHECK(out.Numel() == static_cast<int64_t>(b) * h * d, "numel mismatch");
+  MeshDevice& device = SharedMeshDevice();
+  ttnn::Tensor dev = EnsureDevice2D(out, device);
+  ttnn::Tensor r3 = ttnn::reshape(dev, ttnn::Shape({b, h, d}));
+  CommitDeviceLogical2D(out, std::move(r3), b * h, d);
+}
+
+void EnsureDevice2DForTest(Tensor& t) { (void)EnsureDevice2D(t, SharedMeshDevice()); }
+
+
 // to host (the residency win). Host is marked stale until EnsureHost.
 // Device tensor is stored as logical [rows, cols] TILE (may differ from out's
 // rank-3 view as long as numel matches) so a later Reshape+EnsureDevice2D hits.
@@ -1311,40 +1363,21 @@ bool MemsetDeviceIfCapture(void* p, int value, size_t bytes) {
     // bf16-only, the same polarity as the W7 reservation arm: the geometry
     // is derived from the registered byte size, which is dtype-unambiguous
     // only for 2-byte elements.
-    // CAPTURE-ONLY: an eager fresh-slot zero keeps the host fallback. The
-    // byte size does not name a dtype (the f32 KV masters share these pool
-    // blocks), so serving one eagerly would install a wrongly-typed shadow;
-    // inside the capture the write is banned and the buffer is scratch whose
-    // every consumer reads on device, which is what makes the guess safe.
-    // The capture-time zero still finds its tensor: the cold step's
-    // EnsureDevice2D restage primed the zero at this exact spec
-    // (ZeroCachePrime) and the copy program is warm from the eager copy
-    // lane — ZeroCacheGet refuses a capture-time miss by design.
-    if (!tt_capture_active()) {
-      // Prime the zero-cache AND warm the copy program for this spec during
-      // eager warmup: the capture-time lane runs ttnn::copy(zero_src, *fresh)
-      // whose CopyDeviceOperation hash is shape-specific, so a copy never
-      // executed during warmup is not in the program cache and trace capture
-      // fatals on the missing binary. Also prime ZeroCacheGet for the
-      // [1, cols] bf16 TILE spec — a fresh-slot memset whose geometry never
-      // staged (the 27B bench: a 20480-B res.Zero → [1,10240] bf16 TILE)
-      // would miss mid-capture.
-      if (bytes > 0 && (bytes % 2) == 0) {
-        uint32_t cols = static_cast<uint32_t>(bytes / 2);
-        MeshDevice& md = SharedMeshDevice();
-        md.enable_program_cache();
-        auto shape = ttnn::Shape({1u, cols});
-        ZeroCachePrime(shape, ttnn::DataType::BFLOAT16,
-                       ttnn::Layout::TILE, md);
-        ttnn::Tensor zero_src = ZeroCacheGet(
-            shape, ttnn::DataType::BFLOAT16, ttnn::Layout::TILE, md);
-        ttnn::Tensor tmp = ttnn::empty(shape, ttnn::DataType::BFLOAT16,
-                                       ttnn::Layout::TILE, &md,
-                                       ttnn::MemoryConfig{});
-        ttnn::copy(zero_src, tmp);
-      }
-      return false;
-    }
+    // PASS-UNIFIED (site 2 of this issue): the install below
+    // used to be gated on capture, with eager priming the zero and returning
+    // false (host memset + MarkHostWritten). That left the slot in a
+    // different state per pass: eager ended with no device shadow (the
+    // consumer's stage defined its geometry), capture ended with a [1, cols]
+    // bf16 shadow. The 27B bench then hit the same-numel arm of
+    // EnsureDevice2D mid-capture with a reshape spec the eager pass never ran
+    // ([1,10240] -> [2,5120], the residual kRmsNorm consumes) —
+    // ReshapeViewTiledProgramFactory created its program mid-trace and died
+    // on its to_device write. W4 doctrine: warm in eager what capture
+    // replays — one install in BOTH passes, so the eager consumer runs the
+    // same-numel reshape and the capture replays it as a program-cache hit.
+    // The dtype guess is the one the capture lane already made; zeros are
+    // zeros in every dtype, and the eager lane additionally memsets the host
+    // bytes so the host contract (Memset leaves zeros at p) still holds.
     MeshDevice& device_fresh = SharedMeshDevice();
     device_fresh.enable_program_cache();
     uint32_t cols = 0;
@@ -1369,38 +1402,77 @@ bool MemsetDeviceIfCapture(void* p, int value, size_t bytes) {
         fresh = *s->persistent;
       }
     }
-    if (!fresh.has_value()) {
-      fresh = ttnn::empty(ttnn::Shape({1u, cols}), ttnn::DataType::BFLOAT16,
-                          ttnn::Layout::TILE, &device_fresh,
-                          ttnn::MemoryConfig{});
-      std::lock_guard<std::mutex> g(SlotMutex());
-      if (BufferSlot* s = FindSlot(p)) {
-        // W5 semantics: the allocation becomes the slot's persistent buffer,
-        // so the next zero reuses the same device address.
-        s->persistent = fresh;
-        s->persist_rows = 1;
-        s->persist_cols = cols;
+    bool capture = tt_capture_active();
+    // EAGER retention bound: installing a persistent {1, cols} shadow for
+    // EVERY fresh-slot memset eagerly retained multi-MB buffers at weights
+    // load (the 27B warmup memsets 3-8 MB slots) and OOMed DRAM — the device
+    // runs within ~70 MB of full, and a retained 120 KB class × 48 slots ate
+    // the margin a 268 MB ttnn::where needed. Inside the capture the install
+    // is the point (the write is banned, the buffer is scratch); eagerly it
+    // only needs to WARM what capture replays, so keep the install to the
+    // captured residual/scratch class (the 27B site-2 buffer is 20480 B) and
+    // let larger slots keep the pre-fix host fallback.
+    bool scratch_scale = bytes <= (size_t{1} << 16);
+    if (capture || scratch_scale) {
+      if (!fresh.has_value()) {
+        fresh = ttnn::empty(ttnn::Shape({1u, cols}), ttnn::DataType::BFLOAT16,
+                            ttnn::Layout::TILE, &device_fresh,
+                            ttnn::MemoryConfig{});
+        std::lock_guard<std::mutex> g(SlotMutex());
+        if (BufferSlot* s = FindSlot(p)) {
+          // W5 semantics: the allocation becomes the slot's persistent
+          // buffer, so the next zero reuses the same device address.
+          s->persistent = fresh;
+          s->persist_rows = 1;
+          s->persist_cols = cols;
+        }
+      }
+      ttnn::Tensor zero_src = ZeroCacheGet(*fresh, device_fresh);
+      if (std::getenv("VT_TT_TRACE_DEBUG") != nullptr)
+        std::fprintf(stderr,
+                     "[TT-TRACE] device zero-fill (fresh slot %p cols=%u cap=%d)\n",
+                     p, cols, (int)capture);
+      ttnn::Tensor z = ttnn::copy(zero_src, *fresh);
+      (void)z;
+      {
+        std::lock_guard<std::mutex> g(SlotMutex());
+        BufferSlot* s = FindSlot(p);
+        if (s == nullptr) return false;
+        s->device = *fresh;
+        s->dev_rows = 1;
+        s->dev_cols = cols;
+        s->device_current = true;
+        s->host_current = false;
+        s->conv_transposed = false;
+        s->device_reserved = false;  // real zeros installed — reservation spent
+      }
+      if (!capture) {
+        // Eager host contract: Memset leaves zeros at p for direct host
+        // readers, and host_current=true records that honestly (the capture
+        // lane keeps host_current=false — inside the trace nothing reads the
+        // host, and the device lane is the truth).
+        std::memset(p, 0, bytes);
+        std::lock_guard<std::mutex> g(SlotMutex());
+        if (BufferSlot* s = FindSlot(p)) s->host_current = true;
       }
+      return true;
     }
-    ttnn::Tensor zero_src = ZeroCacheGet(*fresh, device_fresh);
-    if (std::getenv("VT_TT_TRACE_DEBUG") != nullptr)
-      std::fprintf(stderr, "[TT-TRACE] device zero-fill (fresh slot %p cols=%u)\n",
-                   p, cols);
-    ttnn::Tensor z = ttnn::copy(zero_src, *fresh);
-    (void)z;
-    {
-      std::lock_guard<std::mutex> g(SlotMutex());
-      BufferSlot* s = FindSlot(p);
-      if (s == nullptr) return false;
-      s->device = *fresh;
-      s->dev_rows = 1;
-      s->dev_cols = cols;
-      s->device_current = true;
-      s->host_current = false;
-      s->conv_transposed = false;
-      s->device_reserved = false;  // real zeros installed — reservation spent
+    // Large eager fresh-slot memset: prime the zero-cache spec and warm the
+    // copy program with a throwaway buffer (no retention), then keep the
+    // host fallback exactly as before this issue's fix.
+    if (bytes > 0 && (bytes % 2) == 0) {
+      auto shape = ttnn::Shape({1u, cols});
+      ZeroCachePrime(shape, ttnn::DataType::BFLOAT16, ttnn::Layout::TILE,
+                     device_fresh);
+      ttnn::Tensor zero_src = ZeroCacheGet(shape, ttnn::DataType::BFLOAT16,
+                                           ttnn::Layout::TILE, device_fresh);
+      ttnn::Tensor tmp = ttnn::empty(shape, ttnn::DataType::BFLOAT16,
+                                     ttnn::Layout::TILE, &device_fresh,
+                                     ttnn::MemoryConfig{});
+      ttnn::copy(zero_src, tmp);
     }
-    return true;
+    std::memset(p, 0, bytes);
+    return false;
   }
   MeshDevice& device = SharedMeshDevice();
   const ttnn::Tensor& shadow = *dev;
diff --git a/tests/vt/test_breakable_graph.cpp b/tests/vt/test_breakable_graph.cpp
index fb8779c8f..de1005060 100644
--- a/tests/vt/test_breakable_graph.cpp
+++ b/tests/vt/test_breakable_graph.cpp
@@ -1225,7 +1225,7 @@ TEST_CASE("BreakableGraph: an outstanding fork is joined BEFORE the segment clos
     vt::ResetGraphBreakStats();
     BreakableGraph g;
     {
-      GraphCaptureScope scope(b, q, g, GraphCaptureMode::kPiecewise);
+      GraphCaptureScope scope(b, q, g, vt::GraphCaptureMode::kPiecewise);
       REQUIRE(scope.active());
       // The model forks and tells the scope, exactly as `laguna.cpp` does.
       vt::GraphNoteFork(aux, done);
@@ -1261,7 +1261,7 @@ TEST_CASE("BreakableGraph: an outstanding fork is joined BEFORE the segment clos
     vt::ResetGraphBreakStats();
     BreakableGraph g;
     {
-      GraphCaptureScope scope(b, q, g, GraphCaptureMode::kPiecewise);
+      GraphCaptureScope scope(b, q, g, vt::GraphCaptureMode::kPiecewise);
       REQUIRE(scope.active());
       vt::GraphNoteFork(aux, done);
       CHECK(scope.outstanding_forks() == 1);
@@ -1363,3 +1363,63 @@ TEST_CASE("BreakableGraph: an outstanding fork is joined BEFORE the segment clos
     CHECK(s.forks_auto_joined == 0);
   }
 }
+
+// ─── tt-27b-region-capture ──────────────────────────────────────────────────
+// The whole-graph fit predicate and the per-region trace-staging census.
+// Host-side: the predicate is pure and the census runs on the recording
+// backend through an injected probe — the same mechanism the Tenstorrent
+// registrar installs `LastTraceBytesForTest` through. The device-side
+// handoff gate (replay-vs-eager byte-exactness across a region boundary on
+// the real trace backend) lives in tests/vt/test_tenstorrent_backend.cpp.
+TEST_CASE("tt-27b-region-capture: WholeGraphTraceFits declines only on a measured over-budget estimate") {
+  // No measurement is never evidence of no fit: zeroed fields answer TRUE so
+  // a backend with no census cannot flip a served arm by silence.
+  CHECK(vt::WholeGraphTraceFits(vt::WholeGraphFitEstimate{}));
+  CHECK(vt::WholeGraphTraceFits({0, 3'187'104, 298'568'896}));
+  CHECK(vt::WholeGraphTraceFits({1'037, 0, 0}));
+  // The recorded 27B census: 1,037 commands x 3.04 MB/command = ~3.15 GB
+  // against 298,568,896 B free — the decline the row exists for.
+  CHECK_FALSE(vt::WholeGraphTraceFits(
+      {1'037, 3'187'104, 298'568'896}));
+  // The same command stream that fits a 4 GiB headroom.
+  CHECK(vt::WholeGraphTraceFits({1'037, 3'187'104, 4'000'000'000}));
+}
+
+TEST_CASE("tt-27b-region-capture: every segment records its trace-staging bytes") {
+  RequireCaptureLane();
+  RecordingCaptureBackend b;
+  vt::Queue q = b.CreateQueue();
+  vt::ResetGraphBreakStats();
+  // Fake the trace-staging byte level the way tt-metal's
+  // get_trace_buffers_size advances across segments: a monotone counter in a
+  // function static, because a GraphRegionBytesProbe is a plain function
+  // pointer and cannot capture.
+  vt::SetGraphRegionBytesProbe([]() -> int64_t {
+    static int64_t n = 0;
+    return ++n * 1'000;
+  });
+
+  BreakableGraph g;
+  // One slot PER break point — the aliasing refusal refuses a shared
+  // destination within one capture (lifetime rule 1).
+  BreakSlot<vt::Tensor> dst0, dst2;
+  int32_t feed = 7;
+  {
+    GraphCaptureScope scope(b, q, g, vt::GraphCaptureMode::kPiecewise);
+    GraphBreak([&] { return MakeBuf(b, {feed, feed, feed, feed}).t; }, dst0);
+    vt::GraphBreak();
+    GraphBreak([&] { return MakeBuf(b, {feed, feed, feed, feed}).t; }, dst2);
+  }
+  const std::vector<int64_t>& rb = g.region_bytes();
+  // THREE breaks -> FOUR segments (the scope's destructor files the final one
+  // after the last break; break_count() + 1 == segment_count()).
+  REQUIRE(rb.size() == 4);
+  for (int64_t bytes : rb) {
+    const bool in_budget = bytes > 0 && bytes <= 50 * 1024 * 1024;
+    CHECK_MESSAGE(in_budget, "region census entry out of range: " << bytes);
+  }
+  // Reset() clears the census with the graph it described.
+  g.Reset();
+  CHECK(g.region_bytes().empty());
+  vt::SetGraphRegionBytesProbe(nullptr);
+}
diff --git a/tests/vt/test_tenstorrent_backend.cpp b/tests/vt/test_tenstorrent_backend.cpp
index 366485fcc..c55626205 100644
--- a/tests/vt/test_tenstorrent_backend.cpp
+++ b/tests/vt/test_tenstorrent_backend.cpp
@@ -31,6 +31,7 @@
 
 #include "vllm/platforms/interface.h"
 #include "vt/backend.h"
+#include "vt/breakable_graph.h"  // tt-27b-region-capture: the region seam
 #include "vt/dtype.h"
 #include "vt/ops.h"
 #include "vt/quant.h"
@@ -45,6 +46,25 @@ namespace vt::tenstorrent {
 // Tenstorrent residency probe (tenstorrent_internal.h); declared here to keep
 // the test TU free of internal includes.
 bool DeviceShadowExact(const Tensor& t, uint32_t rows, uint32_t cols);
+// ISSUE-LOCAL-01M3JXEFQKSZP23PP2HWY9G0VQ surfaces (tenstorrent_residency.cpp).
+void CommitRank3DeviceLogicalForTest(Tensor& out, uint32_t b, uint32_t h, uint32_t d);
+void EnsureDevice2DForTest(Tensor& t);
+bool TraceCaptureActive();
+// Host-free decode warm hooks + shadow readback (tenstorrent_paged.cpp,
+// tenstorrent_device.h); declared here to keep the test TU internal-free.
+void WarmPagedKvShadow(void* k_cache_data, void* v_cache_data,
+                       int64_t num_blocks, int64_t block_size,
+                       int64_t num_kv_heads, int64_t head_size,
+                       int64_t used_blocks);
+void WarmRacIdx(const void* slot_mapping_owner, const int64_t* slots,
+                int64_t num_slots, int64_t block_size,
+                const int32_t* page_table, int64_t page_table_cols,
+                const int32_t* positions);
+bool ReadPagedKvShadowForTest(const void* k_cache_data, float* dst, int64_t n);
+void WarmPaMeta(const int32_t* block_table, int64_t num_reqs, int64_t max_blocks,
+                int64_t bt_row_stride, int64_t bt_col_stride,
+                const int32_t* seq_lens);
+void WarmDecodePos(const int32_t* seq_lens, int64_t num_reqs, bool replay_regime);
 }  // namespace vt::tenstorrent
 
 namespace {
@@ -861,7 +881,21 @@ TEST_CASE("kTENSTORRENT kRopeNeox is BIT-EXACT vs a host F32 reference (small)")
   // VT_TT_HOST_FREE_DECODE (e.g. a suite run under the host-free gate) flips
   // PreferDeviceRope to the device BF16 path even at small T*H and reds the
   // bit-exact checks — so the case owns its own default-path env, mirroring
-  // the inertness-guard case below.
+  // the inertness-guard case below. The opt-out is restored afterwards: a
+  // leaked "0" silently disabled the host-free lane for EVERY later case in
+  // the suite (found via ISSUE-LOCAL-01M3JXEFQKSZP23PP2HWY9G0VQ's fresh-slot
+  // memset case, which needs the default host-free warmup).
+  const bool had_hf = std::getenv("VT_TT_HOST_FREE_DECODE") != nullptr;
+  const std::string saved_hf =
+      had_hf ? std::string(std::getenv("VT_TT_HOST_FREE_DECODE")) : std::string();
+  struct RestoreHf {
+    bool had;
+    std::string saved;
+    ~RestoreHf() {
+      if (had) ::setenv("VT_TT_HOST_FREE_DECODE", saved.c_str(), 1);
+      else ::unsetenv("VT_TT_HOST_FREE_DECODE");
+    }
+  } restore_hf{had_hf, saved_hf};
   ::setenv("VT_TT_HOST_FREE_DECODE", "0", 1);  // opt-out path
   REQUIRE(vt::OpRegistered(vt::OpId::kRopeNeox, DeviceType::kTENSTORRENT));
 
@@ -1841,6 +1875,653 @@ TEST_CASE("kTENSTORRENT SupportsGraphCapture and matmul capture/replay") {
   backend.Free(mc);
 }
 
+// ISSUE-LOCAL-01M3JXEFQKSZP23PP2HWY9G0VQ: EnsureDevice2D's exact-rows/cols
+// arm on a rank-3-committed slot ran a bare ttnn::reshape on the TILED shadow
+// during capture — a program the eager pass never warmed (its row-major
+// reshape is a free view) — so the 27B decode capture created
+// ReshapeViewTiledProgramFactory's program mid-trace and died on its
+// to_device write ("Writes are not supported during trace capture"). The fix
+// runs one chain in both passes; this case fails red if the capture pass
+// diverges again: with the old branch restored, BeginCapture fatals.
+TEST_CASE("kTENSTORRENT EnsureDevice2D rank-3 reshape is capture-safe (one chain both passes)") {
+  if (!TenstorrentPresent()) {
+    MESSAGE("SKIPPED: no Tenstorrent device on this box");
+    return;
+  }
+  Backend& backend = vt::GetBackend(DeviceType::kTENSTORRENT);
+  REQUIRE(backend.SupportsGraphCapture());
+
+  // [2, 48, 32] device result, flat 2D geometry [96, 32]: reshaping the tiled
+  // rank-3 shadow to 2D is NOT a metadata view (the 48 second-last dim is not
+  // tile-aligned — reshape.cpp's this_is_view), so the divergent capture arm
+  // had to create ReshapeViewTiledProgramFactory's program mid-capture.
+  constexpr uint32_t B = 2, H = 48, D = 32;
+  constexpr uint32_t Rows = B * H, Cols = D;
+  std::vector<float> host(static_cast<size_t>(Rows * Cols), 0.25f);
+
+  void* mem = backend.Alloc(host.size() * sizeof(float));
+  Queue q = backend.CreateQueue();
+  backend.Copy(q, mem, host.data(), host.size() * sizeof(float));
+  Tensor t = Tensor::Contiguous(mem, vt::DType::kF32, Device{DeviceType::kTENSTORRENT, 0},
+                                {Rows, Cols});
+
+  // Commit the rank-3 shadow under the 2D slot record, warm the eager chain,
+  // re-commit (the warm consumed the rank-3 logical shape), then capture.
+  vt::tenstorrent::CommitRank3DeviceLogicalForTest(t, B, H, D);
+  vt::tenstorrent::EnsureDevice2DForTest(t);
+  vt::tenstorrent::CommitRank3DeviceLogicalForTest(t, B, H, D);
+
+  backend.BeginCapture(q);
+  vt::tenstorrent::EnsureDevice2DForTest(t);
+  backend.EndCapture(q);
+  MESSAGE("shadow exact after capture: ", vt::tenstorrent::DeviceShadowExact(t, Rows, Cols));
+  backend.Replay(q);
+  backend.Replay(q);
+
+  std::vector<float> after(host.size(), 0.0f);
+  backend.Copy(q, after.data(), mem, after.size() * sizeof(float));
+  float max_abs = 0.0f;
+  for (size_t i = 0; i < after.size(); ++i)
+    max_abs = std::max(max_abs, std::fabs(after[i] - host[i]));
+  MESSAGE("replay max_abs vs staged: ", max_abs);
+  // Diagnostic: a fresh eager EnsureDevice2D must serve the replayed bytes.
+  vt::tenstorrent::EnsureDevice2DForTest(t);
+  std::vector<float> after2(host.size(), 0.0f);
+  backend.Copy(q, after2.data(), mem, after2.size() * sizeof(float));
+  float max_abs2 = 0.0f;
+  for (size_t i = 0; i < after2.size(); ++i)
+    max_abs2 = std::max(max_abs2, std::fabs(after2[i] - host[i]));
+  MESSAGE("post-eager max_abs vs staged: ", max_abs2);
+  CHECK(max_abs2 < 1e-3f);
+
+  backend.Free(mem);
+}
+
+// ISSUE-LOCAL-01M3JXEFQKSZP23PP2HWY9G0VQ, site 2: MemsetDeviceIfCapture's
+// fresh-slot lane installed a [1, cols] bf16 shadow ONLY under capture; the
+// eager pass primed the zero and kept the host fallback, so the slot ended
+// each pass in a different state. The 27B bench: DBuf::Zero of the 20480-B
+// residual, then kRmsNorm's EnsureDevice2D at [2, 5120] — the capture step
+// hit the same-numel arm with a reshape spec the eager pass never ran, and
+// ReshapeViewTiledProgramFactory created its program mid-trace and died on
+// its to_device write. The fix installs the same [1, cols] shadow in BOTH
+// passes, so the eager consumer warms the reshape and the capture replays it
+// as a cache hit. This case fails red if the eager lane diverges again: with
+// the capture-only install restored, the EnsureDevice2D inside capture
+// creates the program and TT_FATALs ("Writes are not supported during trace
+// capture").
+TEST_CASE("kTENSTORRENT fresh-slot Memset installs the same shadow in both passes") {
+  if (!TenstorrentPresent()) {
+    MESSAGE("SKIPPED: no Tenstorrent device on this box");
+    return;
+  }
+  Backend& backend = vt::GetBackend(DeviceType::kTENSTORRENT);
+  REQUIRE(backend.SupportsGraphCapture());
+
+  // The production geometry: 20480 B = 10240 bf16 elems, consumer [2, 5120]
+  // — the same-numel arm's [1, 10240] -> [2, 5120] reshape is not a metadata
+  // view (different tile counts), so it needs the warmed program exactly as
+  // the bench ran it (bf16 tensor: the shadow's element count is bytes/2,
+  // which must equal the consumer's numel for the same-numel arm to fire).
+  constexpr uint32_t Rows = 2, Cols = 5120;
+  void* mem = backend.Alloc(Rows * Cols * sizeof(uint16_t));
+  // The pool recycles blocks across test cases: acquire the block so both
+  // passes start from the same W7 fresh-slot state (no stale shadow from a
+  // previous tenant can route the eager memset down a different lane).
+  backend.OnScratchBlockAcquired(mem);
+  Queue q = backend.CreateQueue();
+
+  // The case needs the default host-free decode lane for BOTH passes; own the
+  // env explicitly (earlier cases legitimately run under the opt-out).
+  const bool had_hf = std::getenv("VT_TT_HOST_FREE_DECODE") != nullptr;
+  const std::string saved_hf =
+      had_hf ? std::string(std::getenv("VT_TT_HOST_FREE_DECODE")) : std::string();
+  struct RestoreHf {
+    bool had;
+    std::string saved;
+    ~RestoreHf() {
+      if (had) ::setenv("VT_TT_HOST_FREE_DECODE", saved.c_str(), 1);
+      else ::unsetenv("VT_TT_HOST_FREE_DECODE");
+    }
+  } restore_hf{had_hf, saved_hf};
+  ::unsetenv("VT_TT_HOST_FREE_DECODE");
+
+  MESSAGE("capture flag at entry: ", vt::tenstorrent::TraceCaptureActive());
+
+  // Warmup step: fresh slot, DBuf::Zero, then the consumer's stage. With the
+  // fix, the Memset installs the [1, 5120] shadow and this EnsureDevice2D
+  // runs (and warms) the same-numel reshape in the eager pass.
+  backend.Memset(q, mem, 0, Rows * Cols * sizeof(uint16_t));
+  Tensor t = Tensor::Contiguous(mem, vt::DType::kBF16, Device{DeviceType::kTENSTORRENT, 0},
+                                {Rows, Cols});
+  vt::tenstorrent::EnsureDevice2DForTest(t);
+
+  // The production capture step saw this buffer as a FRESH slot (the pool
+  // handed the block to a new tensor between steps, W7 semantics). Reproduce
+  // the acquisition so the capture memset takes the fresh-slot lane exactly
+  // like the bench's trace did.
+  backend.OnScratchBlockAcquired(mem);
+  // Capture step: the production zero-fill runs INSIDE the captured region
+  // (the bench trace's "device zero-fill (fresh slot ...)" fires after
+  // BeginCapture), so begin capture first, then repeat the memset on the
+  // recycled slot and the consumer's stage.
+  backend.BeginCapture(q);
+  backend.Memset(q, mem, 0, Rows * Cols * sizeof(uint16_t));
+  vt::tenstorrent::EnsureDevice2DForTest(t);
+  backend.EndCapture(q);
+  backend.Replay(q);
+
+  // The replayed region holds the zeros the memsets wrote.
+  std::vector<uint16_t> after(Rows * Cols, 0x3f80);  // f32 1.0 bits in the low half: garbage if read
+  backend.Copy(q, after.data(), mem, after.size() * sizeof(uint16_t));
+  float max_abs = 0.0f;
+  for (uint16_t v : after) max_abs = std::max(max_abs, std::fabs(static_cast<float>(v)));
+  MESSAGE("post-replay nonzero bf16 elems: ", max_abs);
+  CHECK(max_abs == 0.0f);
+
+  backend.Free(mem);
+}
+
+// ISSUE-LOCAL-01M3JXEFQKSZP23PP2HWY9G0VQ, site 3 (blocker A): the batched
+// decode RAC. TryReshapeAndCacheDeviceDecode declined num_slots > 1 ("decode
+// T=1 only for now"), so a c2 serve leg fell back to the host path inside the
+// capture — EnsureHost(k) read the rope K/V shadows back mid-trace and
+// TT_FATAL'd "Reads are not supported during trace capture"
+// (fd_mesh_command_queue.cpp:873). The fix admits the batched decode: one
+// [nkv_pad, d] shard per user on the sharded input (shard i is user i, the
+// fused-update kernel maps core i to update_idxs[i] / page-table row i), and
+// the SAME slice->multiply->concat->copy->paged_fused_update_cache sequence
+// runs in the eager pass (which warms the programs) and in the capture. This
+// case reproduces the exact leg: batched rope shadows committed rank-3, RAC
+// inside a capture, then replay, and verifies BOTH users' KV landed in the
+// paged-KV device shadow. Red: with the decline restored, BeginCapture ->
+// rac() TT_FATALs at fd_mesh_command_queue.cpp:873.
+TEST_CASE("kTENSTORRENT batched decode RAC is capture-safe (num_slots=2)") {
+  if (!TenstorrentPresent()) {
+    MESSAGE("SKIPPED: no Tenstorrent device on this box");
+    return;
+  }
+  Backend& backend = vt::GetBackend(DeviceType::kTENSTORRENT);
+  REQUIRE(backend.SupportsGraphCapture());
+
+  // TILE-legal decode geometry: d and bs multiples of 32 (TryRAC's arms).
+  constexpr int64_t NBlocks = 2, Bsz = 32, Hkv = 2, Dh = 32, C = 2;
+  const size_t cache_elems = static_cast<size_t>(NBlocks * Bsz * Hkv * Dh);
+  const size_t kv_elems = static_cast<size_t>(C * Hkv * Dh);
+
+  // The device-RAC lane is the default for the host-free decode path; own the
+  // env explicitly and restore it (earlier cases legitimately opt out).
+  const bool had_hf = std::getenv("VT_TT_HOST_FREE_DECODE") != nullptr;
+  const std::string saved_hf =
+      had_hf ? std::string(std::getenv("VT_TT_HOST_FREE_DECODE")) : std::string();
+  struct RestoreHf {
+    bool had;
+    std::string saved;
+    ~RestoreHf() {
+      if (had) ::setenv("VT_TT_HOST_FREE_DECODE", saved.c_str(), 1);
+      else ::unsetenv("VT_TT_HOST_FREE_DECODE");
+    }
+  } restore_hf{had_hf, saved_hf};
+  ::setenv("VT_TT_HOST_FREE_DECODE", "1", 1);
+
+  Backend& be = backend;
+  void* mem_kc = be.Alloc(cache_elems * sizeof(uint16_t));
+  void* mem_vc = be.Alloc(cache_elems * sizeof(uint16_t));
+  void* mem_k = be.Alloc(kv_elems * sizeof(uint16_t));
+  void* mem_v = be.Alloc(kv_elems * sizeof(uint16_t));
+  void* mem_slots = be.Alloc(C * sizeof(int64_t));
+  Queue q = be.CreateQueue();
+
+  // Seed caches with a recognizable pattern; user u's page is block u.
+  std::vector<uint16_t> seed(cache_elems);
+  for (size_t i = 0; i < seed.size(); ++i)
+    seed[i] = static_cast<uint16_t>(0x3800 + (i % 1023));  // bf16 ~0.03.. pattern
+  std::vector<uint16_t> host_k(kv_elems), host_v(kv_elems);
+  for (size_t i = 0; i < host_k.size(); ++i) {
+    host_k[i] = static_cast<uint16_t>(0x3c00 + (i % 31));  // ~1.0x bf16 pattern
+    host_v[i] = static_cast<uint16_t>(0x3d00 + (i % 29));
+  }
+  std::vector<int64_t> slots{0, Bsz};  // user0 -> block0 off0, user1 -> block1 off0
+  auto bf16_val = [](uint16_t h) {
+    uint32_t bits = static_cast<uint32_t>(h) << 16;
+    float f;
+    std::memcpy(&f, &bits, sizeof(f));
+    return f;
+  };
+  be.Copy(q, mem_kc, seed.data(), seed.size() * sizeof(uint16_t));
+  be.Copy(q, mem_vc, seed.data(), seed.size() * sizeof(uint16_t));
+  be.Copy(q, mem_k, host_k.data(), host_k.size() * sizeof(uint16_t));
+  be.Copy(q, mem_v, host_v.data(), host_v.size() * sizeof(uint16_t));
+  be.Copy(q, mem_slots, slots.data(), slots.size() * sizeof(int64_t));
+
+  // Warm hooks exactly as the driver runs them, BEFORE capture.
+  vt::tenstorrent::WarmPagedKvShadow(mem_kc, mem_vc, NBlocks, Bsz, Hkv, Dh,
+                                     /*used_blocks=*/NBlocks);
+  // update_idx = seq_lens-1 = 0 for both users; page-table row = {block u}.
+  std::vector<int32_t> block_table{0, 1};
+  std::vector<int32_t> seq_lens{1, 1};
+  vt::tenstorrent::WarmRacIdx(mem_slots, slots.data(), C, Bsz,
+                              block_table.data(), /*page_table_cols=*/1,
+                              seq_lens.data());
+
+  Tensor tkc = Tensor::Contiguous(mem_kc, vt::DType::kBF16,
+                                  Device{DeviceType::kTENSTORRENT, 0},
+                                  {NBlocks, Bsz, Hkv, Dh});
+  Tensor tvc = Tensor::Contiguous(mem_vc, vt::DType::kBF16,
+                                  Device{DeviceType::kTENSTORRENT, 0},
+                                  {NBlocks, Bsz, Hkv, Dh});
+  Tensor tsl = Tensor::Contiguous(mem_slots, vt::DType::kI64,
+                                  Device{DeviceType::kTENSTORRENT, 0}, {C});
+
+  auto rac = reinterpret_cast<vt::ReshapeAndCacheFn>(
+      vt::GetOp(vt::OpId::kReshapeAndCache, DeviceType::kTENSTORRENT));
+
+  // Warm (eager) pass FIRST — the W4 doctrine: warm in eager what capture
+  // replays. This is the same call the cold step's ForwardLayers makes.
+  {
+    // Stage as flat rank-2 [C*Hkv, Dh] (EnsureDevice2D's contract) and commit
+    // the rank-3 [C, Hkv, Dh] logical the rope result carries; the RAC call
+    // itself sees the rank-3 view the decode graph hands the kernel.
+    Tensor tk2d = Tensor::Contiguous(mem_k, vt::DType::kBF16,
+                                     Device{DeviceType::kTENSTORRENT, 0},
+                                     {C * Hkv, Dh});
+    Tensor tv2d = Tensor::Contiguous(mem_v, vt::DType::kBF16,
+                                     Device{DeviceType::kTENSTORRENT, 0},
+                                     {C * Hkv, Dh});
+    vt::tenstorrent::CommitRank3DeviceLogicalForTest(tk2d, C, Hkv, Dh);
+    vt::tenstorrent::CommitRank3DeviceLogicalForTest(tv2d, C, Hkv, Dh);
+    Tensor tk = Tensor::Contiguous(mem_k, vt::DType::kBF16,
+                                   Device{DeviceType::kTENSTORRENT, 0},
+                                   {C, Hkv, Dh});
+    Tensor tv = Tensor::Contiguous(mem_v, vt::DType::kBF16,
+                                   Device{DeviceType::kTENSTORRENT, 0},
+                                   {C, Hkv, Dh});
+    rac(q, tk, tv, tkc, tvc, tsl);
+    // Per-pass probe 1: the EAGER pass alone must land both users' KV.
+    {
+      std::vector<float> kp(cache_elems, 0.0f);
+      REQUIRE(vt::tenstorrent::ReadPagedKvShadowForTest(mem_kc, kp.data(),
+                                                        (int64_t)kp.size()));
+      int bad = 0;
+      for (int64_t u = 0; u < C; ++u)
+        for (int64_t h = 0; h < Hkv; ++h)
+          for (int64_t e = 0; e < Dh; ++e) {
+            const size_t dst = (static_cast<size_t>(u * Bsz) * Hkv + h) * Dh +
+                               static_cast<size_t>(e);
+            const size_t src =
+                static_cast<size_t>(u * Hkv + h) * Dh + static_cast<size_t>(e);
+            if (std::fabs(kp[dst] - bf16_val(host_k[src])) >= 0.05f) {
+              if (bad < 4) MESSAGE("EAGER K bad u=", u, " h=", h, " e=", e,
+                                   " got=", kp[dst], " want=", bf16_val(host_k[src]));
+              ++bad;
+            }
+          }
+      MESSAGE("post-EAGER bad K elems: ", bad);
+    }
+  }
+
+  // Capture pass: fresh rank-3 committed shadows (as rope leaves them each
+  // step) and the identical RAC inside a trace.
+  {
+    Tensor tk2d = Tensor::Contiguous(mem_k, vt::DType::kBF16,
+                                     Device{DeviceType::kTENSTORRENT, 0},
+                                     {C * Hkv, Dh});
+    Tensor tv2d = Tensor::Contiguous(mem_v, vt::DType::kBF16,
+                                     Device{DeviceType::kTENSTORRENT, 0},
+                                     {C * Hkv, Dh});
+    vt::tenstorrent::CommitRank3DeviceLogicalForTest(tk2d, C, Hkv, Dh);
+    vt::tenstorrent::CommitRank3DeviceLogicalForTest(tv2d, C, Hkv, Dh);
+    Tensor tk2 = Tensor::Contiguous(mem_k, vt::DType::kBF16,
+                                    Device{DeviceType::kTENSTORRENT, 0},
+                                    {C, Hkv, Dh});
+    Tensor tv2 = Tensor::Contiguous(mem_v, vt::DType::kBF16,
+                                    Device{DeviceType::kTENSTORRENT, 0},
+                                    {C, Hkv, Dh});
+    be.BeginCapture(q);
+    // If anything inside the region throws, end the capture before
+    // propagating — a leaked active capture poisons every later case.
+    struct EndOnExit {
+      Backend& b;
+      Queue& q;
+      ~EndOnExit() {
+        if (vt::tenstorrent::TraceCaptureActive()) {
+          try { b.EndCapture(q); } catch (...) {}
+        }
+      }
+    } end_guard{be, q};
+    rac(q, tk2, tv2, tkc, tvc, tsl);
+    be.EndCapture(q);
+    // Per-pass probe 2: the CAPTURE pass itself writes through the recorded
+    // ops (capture executes the region once) — check before any replay.
+    {
+      std::vector<float> kp(cache_elems, 0.0f);
+      REQUIRE(vt::tenstorrent::ReadPagedKvShadowForTest(mem_kc, kp.data(),
+                                                        (int64_t)kp.size()));
+      int bad = 0;
+      for (int64_t u = 0; u < C; ++u)
+        for (int64_t h = 0; h < Hkv; ++h)
+          for (int64_t e = 0; e < Dh; ++e) {
+            const size_t dst = (static_cast<size_t>(u * Bsz) * Hkv + h) * Dh +
+                               static_cast<size_t>(e);
+            const size_t src =
+                static_cast<size_t>(u * Hkv + h) * Dh + static_cast<size_t>(e);
+            if (std::fabs(kp[dst] - bf16_val(host_k[src])) >= 0.05f) ++bad;
+          }
+      MESSAGE("post-CAPTURE (pre-replay) bad K elems: ", bad);
+    }
+    be.Replay(q);
+    // Per-pass probe 3: the REPLAY must reproduce the same bytes.
+    {
+      std::vector<float> kp(cache_elems, 0.0f);
+      REQUIRE(vt::tenstorrent::ReadPagedKvShadowForTest(mem_kc, kp.data(),
+                                                        (int64_t)kp.size()));
+      int bad = 0;
+      for (int64_t u = 0; u < C; ++u)
+        for (int64_t h = 0; h < Hkv; ++h)
+          for (int64_t e = 0; e < Dh; ++e) {
+            const size_t dst = (static_cast<size_t>(u * Bsz) * Hkv + h) * Dh +
+                               static_cast<size_t>(e);
+            const size_t src =
+                static_cast<size_t>(u * Hkv + h) * Dh + static_cast<size_t>(e);
+            if (std::fabs(kp[dst] - bf16_val(host_k[src])) >= 0.05f) ++bad;
+          }
+      MESSAGE("post-REPLAY bad K elems: ", bad);
+    }
+  }
+
+  // Both users' KV must sit at their pages in the paged-KV DEVICE shadow.
+  std::vector<float> kc(cache_elems, 0.0f), vc(cache_elems, 0.0f);
+  REQUIRE(vt::tenstorrent::ReadPagedKvShadowForTest(mem_kc, kc.data(),
+                                                    (int64_t)kc.size()));
+  REQUIRE(vt::tenstorrent::ReadPagedKvShadowForTest(mem_vc, vc.data(),
+                                                    (int64_t)vc.size()));
+  // Host bf16->float reference for user u's token (bf16_val above).
+  const size_t tok = static_cast<size_t>(Hkv * Dh);
+  int k_ok = 0, v_ok = 0, k_bad = 0, v_bad = 0;
+  const int total = static_cast<int>(C * tok);
+  for (int64_t u = 0; u < C; ++u) {
+    for (int64_t h = 0; h < Hkv; ++h) {
+      for (int64_t e = 0; e < Dh; ++e) {
+        const size_t dst =
+            (static_cast<size_t>(u * Bsz) * Hkv + h) * Dh + static_cast<size_t>(e);
+        const size_t src = static_cast<size_t>(u * Hkv + h) * Dh + static_cast<size_t>(e);
+        bool kb = std::fabs(kc[dst] - bf16_val(host_k[src])) >= 0.05f;
+        bool vb = std::fabs(vc[dst] - bf16_val(host_v[src])) >= 0.05f;
+        if (kb) { ++k_bad; if (k_bad <= 4) MESSAGE("K bad u=", u, " h=", h,
+                   " e=", e, " got=", kc[dst], " want=", bf16_val(host_k[src])); }
+        if (vb) { ++v_bad; if (v_bad <= 4) MESSAGE("V bad u=", u, " h=", h,
+                   " e=", e, " got=", vc[dst], " want=", bf16_val(host_v[src])); }
+        if (!kb) ++k_ok;
+        if (!vb) ++v_ok;
+      }
+    }
+  }
+  (void)tok;
+  MESSAGE("token-exact K elems: ", k_ok, "/", total, " V: ", v_ok, "/", total);
+  CHECK(k_ok == total);
+  CHECK(v_ok == total);
+
+  be.Free(mem_kc);
+  be.Free(mem_vc);
+  be.Free(mem_k);
+  be.Free(mem_v);
+  be.Free(mem_slots);
+}
+
+// Blocker A's sibling site (ISSUE-LOCAL-01M3JXEFQKSZP23PP2HWY9G0VQ): the
+// batched (B>1) decode PagedAttention must serve INSIDE a capture with a
+// per-user device path — the same doctrine the RAC fix established. The leg
+// fatal: the B>1 Q 4D materialization declined under capture, the host Q arm
+// refused loudly, the device path returned false, and PagedAttentionKernel's
+// host oracle EnsureHost(k_cache)-ed mid-trace ("Reads are not supported
+// during trace capture", fd_mesh_command_queue.cpp:873). The B>1 branch runs
+// the IDENTICAL multiply(reshape(...)) chain in both passes, so the eager
+// step warms the reshape program and the capture replays it as a cache hit
+// (W4 doctrine) — the decline is removed, and this case pins the guarantee:
+// eager pass, capture pass, and replay must all produce the same attention.
+TEST_CASE("kTENSTORRENT batched decode PagedAttention is capture-safe (num_reqs=2)") {
+  if (!TenstorrentPresent()) {
+    MESSAGE("SKIPPED: no Tenstorrent device on this box");
+    return;
+  }
+  Backend& backend = vt::GetBackend(DeviceType::kTENSTORRENT);
+  REQUIRE(backend.SupportsGraphCapture());
+
+  // TILE-legal decode geometry: d and bs multiples of 32; Hq padded per batch
+  // to a full tile (32) so the B>1 identity-Q reshape storage stays per-batch
+  // tile-aligned (the layout sdpa_decode reads); GQA ratio Hq/Hkv so the
+  // sdpa_decode arm serves (TryPADecodeDevice requires hq % nkv == 0).
+  constexpr int64_t NBlocks = 2, Bsz = 32, Hkv = 2, Hq = 32, Dh = 32, C = 2;
+  constexpr int64_t MaxBlk = 2;  // page-table columns (block u + one padding col)
+  const size_t cache_elems = static_cast<size_t>(NBlocks * Bsz * Hkv * Dh);
+  const size_t kv_elems = static_cast<size_t>(C * Hkv * Dh);
+  const size_t q_elems = static_cast<size_t>(C * Hq * Dh);
+
+  const bool had_hf = std::getenv("VT_TT_HOST_FREE_DECODE") != nullptr;
+  const std::string saved_hf =
+      had_hf ? std::string(std::getenv("VT_TT_HOST_FREE_DECODE")) : std::string();
+  struct RestoreHf {
+    bool had;
+    std::string saved;
+    ~RestoreHf() {
+      if (had) ::setenv("VT_TT_HOST_FREE_DECODE", saved.c_str(), 1);
+      else ::unsetenv("VT_TT_HOST_FREE_DECODE");
+    }
+  } restore_hf{had_hf, saved_hf};
+  ::setenv("VT_TT_HOST_FREE_DECODE", "1", 1);
+
+  Backend& be = backend;
+  void* mem_kc = be.Alloc(cache_elems * sizeof(uint16_t));
+  void* mem_vc = be.Alloc(cache_elems * sizeof(uint16_t));
+  void* mem_k = be.Alloc(kv_elems * sizeof(uint16_t));
+  void* mem_v = be.Alloc(kv_elems * sizeof(uint16_t));
+  void* mem_q = be.Alloc(q_elems * sizeof(float));
+  void* mem_out1 = be.Alloc(q_elems * sizeof(float));
+  void* mem_out2 = be.Alloc(q_elems * sizeof(float));
+  void* mem_slots = be.Alloc(C * sizeof(int64_t));
+  void* mem_bt = be.Alloc(C * MaxBlk * sizeof(int32_t));
+  void* mem_sl = be.Alloc(C * sizeof(int32_t));
+  void* mem_qsl = be.Alloc((C + 1) * sizeof(int32_t));
+  Queue q = be.CreateQueue();
+
+  // Deterministic Q and per-user rope K/V patterns.
+  std::vector<float> host_q(q_elems);
+  for (size_t i = 0; i < host_q.size(); ++i)
+    host_q[i] = 0.25f * static_cast<float>((i * 37) % 17) - 1.0f;
+  std::vector<uint16_t> host_k(kv_elems), host_v(kv_elems);
+  for (size_t i = 0; i < host_k.size(); ++i) {
+    host_k[i] = static_cast<uint16_t>(0x3c00 + (i % 31));
+    host_v[i] = static_cast<uint16_t>(0x3d00 + (i % 29));
+  }
+  std::vector<uint16_t> seed(cache_elems);
+  for (size_t i = 0; i < seed.size(); ++i)
+    seed[i] = static_cast<uint16_t>(0x3800 + (i % 1023));
+  // user u attends exactly one token: its own, at block u offset 0. The
+  // padding column is block 0 (a real block) so the page-table read stays in
+  // range; only column 0 is attended at seq_len 1.
+  std::vector<int64_t> slots{0, Bsz};
+  std::vector<int32_t> btab{0, 0, 1, 0};
+  std::vector<int32_t> seqlens{1, 1};
+  std::vector<int32_t> qsl{0, 1, 2};
+  be.Copy(q, mem_kc, seed.data(), seed.size() * sizeof(uint16_t));
+  be.Copy(q, mem_vc, seed.data(), seed.size() * sizeof(uint16_t));
+  be.Copy(q, mem_k, host_k.data(), host_k.size() * sizeof(uint16_t));
+  be.Copy(q, mem_v, host_v.data(), host_v.size() * sizeof(uint16_t));
+  be.Copy(q, mem_q, host_q.data(), host_q.size() * sizeof(float));
+  be.Copy(q, mem_slots, slots.data(), slots.size() * sizeof(int64_t));
+  be.Copy(q, mem_bt, btab.data(), btab.size() * sizeof(int32_t));
+  be.Copy(q, mem_sl, seqlens.data(), seqlens.size() * sizeof(int32_t));
+  be.Copy(q, mem_qsl, qsl.data(), qsl.size() * sizeof(int32_t));
+
+  // Warm the paged-KV shadow and stage this step's K/V through the (already
+  // capture-proven) RAC batched lane, exactly as the decode graph does.
+  vt::tenstorrent::WarmPagedKvShadow(mem_kc, mem_vc, NBlocks, Bsz, Hkv, Dh,
+                                     /*used_blocks=*/NBlocks);
+  vt::tenstorrent::WarmRacIdx(mem_slots, slots.data(), C, Bsz,
+                              btab.data(), MaxBlk, seqlens.data());
+  {
+    auto rk3 = [&](void* mem, int64_t rows, int64_t cols) {
+      Tensor t2d = Tensor::Contiguous(mem, vt::DType::kBF16,
+                                      Device{DeviceType::kTENSTORRENT, 0}, {rows, cols});
+      vt::tenstorrent::CommitRank3DeviceLogicalForTest(
+          t2d, rows / Hkv, Hkv, cols);
+      return Tensor::Contiguous(mem, vt::DType::kBF16,
+                                Device{DeviceType::kTENSTORRENT, 0},
+                                {rows / Hkv, Hkv, cols});
+    };
+    Tensor tk = rk3(mem_k, C * Hkv, Dh);
+    Tensor tv = rk3(mem_v, C * Hkv, Dh);
+    Tensor tkc = Tensor::Contiguous(mem_kc, vt::DType::kBF16,
+                                    Device{DeviceType::kTENSTORRENT, 0},
+                                    {NBlocks, Bsz, Hkv, Dh});
+    Tensor tvc = Tensor::Contiguous(mem_vc, vt::DType::kBF16,
+                                    Device{DeviceType::kTENSTORRENT, 0},
+                                    {NBlocks, Bsz, Hkv, Dh});
+    Tensor tsl = Tensor::Contiguous(mem_slots, vt::DType::kI64,
+                                    Device{DeviceType::kTENSTORRENT, 0}, {C});
+    auto rac = reinterpret_cast<vt::ReshapeAndCacheFn>(
+        vt::GetOp(vt::OpId::kReshapeAndCache, DeviceType::kTENSTORRENT));
+    rac(q, tk, tv, tkc, tvc, tsl);
+  }
+
+  // PA meta warm (outside capture): allocates the persistent page_table +
+  // cur_pos and seeds cur_pos = seq_lens - 1, so the captured sdpa_decode
+  // replays against stable addresses. WarmDecodePos seeds the DecodePos entry
+  // FIRST so WarmPaMeta aliases cur_pos to the on-device-advanced buffer —
+  // under full-suite history captures already happened (r2_steady), and the
+  // #1105 guard refuses a standalone cur_pos allocated then.
+  vt::tenstorrent::WarmDecodePos(seqlens.data(), C, /*replay_regime=*/false);
+  vt::tenstorrent::WarmPaMeta(btab.data(), C, MaxBlk, MaxBlk, 1, seqlens.data());
+
+  vt::PagedAttentionArgs args;
+  args.scale = 0.353553f;
+  args.causal = true;
+
+  const Device tt{DeviceType::kTENSTORRENT, 0};
+  Tensor tbt = Tensor::Contiguous(mem_bt, vt::DType::kI32, tt, {C, MaxBlk});
+  Tensor tsl = Tensor::Contiguous(mem_sl, vt::DType::kI32, tt, {C});
+  Tensor tqsl = Tensor::Contiguous(mem_qsl, vt::DType::kI32, tt, {C + 1});
+  Tensor tkc = Tensor::Contiguous(mem_kc, vt::DType::kBF16, tt, {NBlocks, Bsz, Hkv, Dh});
+  Tensor tvc = Tensor::Contiguous(mem_vc, vt::DType::kBF16, tt, {NBlocks, Bsz, Hkv, Dh});
+
+  // Stage a K/V generation through the (already capture-proven) RAC batched
+  // lane, exactly as the decode graph does: host patterns -> rank-3 committed
+  // rope shadows -> device paged-KV write.
+  auto stage_kv = [&](uint16_t kbase, uint16_t vbase) {
+    std::vector<uint16_t> hk(kv_elems), hv(kv_elems);
+    for (size_t i = 0; i < hk.size(); ++i) {
+      hk[i] = static_cast<uint16_t>(kbase + (i % 31));
+      hv[i] = static_cast<uint16_t>(vbase + (i % 29));
+    }
+    be.Copy(q, mem_k, hk.data(), hk.size() * sizeof(uint16_t));
+    be.Copy(q, mem_v, hv.data(), hv.size() * sizeof(uint16_t));
+    auto rk3 = [&](void* mem, int64_t rows, int64_t cols) {
+      Tensor t2d = Tensor::Contiguous(mem, vt::DType::kBF16, tt, {rows, cols});
+      vt::tenstorrent::CommitRank3DeviceLogicalForTest(t2d, rows / Hkv, Hkv, cols);
+      return Tensor::Contiguous(mem, vt::DType::kBF16, tt, {rows / Hkv, Hkv, cols});
+    };
+    Tensor tk = rk3(mem_k, C * Hkv, Dh);
+    Tensor tv = rk3(mem_v, C * Hkv, Dh);
+    Tensor tslots = Tensor::Contiguous(mem_slots, vt::DType::kI64, tt, {C});
+    auto rac = reinterpret_cast<vt::ReshapeAndCacheFn>(
+        vt::GetOp(vt::OpId::kReshapeAndCache, DeviceType::kTENSTORRENT));
+    rac(q, tk, tv, tkc, tvc, tslots);
+  };
+
+  // Commit the query's [C*Hq, Dh] device shadow (rope leaves it resident) and
+  // hand the op the rank-3 [C, Hq, Dh] view the decode graph carries.
+  auto commit_q = [&]() {
+    Tensor tq2d = Tensor::Contiguous(mem_q, vt::DType::kF32, tt, {C * Hq, Dh});
+    vt::tenstorrent::EnsureDevice2DForTest(tq2d);
+    return Tensor::Contiguous(mem_q, vt::DType::kF32, tt, {C, Hq, Dh});
+  };
+  auto make_out = [&](void* mem) {
+    return Tensor::Contiguous(mem, vt::DType::kF32, tt, {C, Hq, Dh});
+  };
+  auto read_out = [&](void* mem) {
+    std::vector<float> host(q_elems);
+    be.Copy(q, host.data(), mem, host.size() * sizeof(float));
+    return host;
+  };
+  auto check_out = [&](const std::vector<float>& got, const char* what) {
+    int nonzero = 0;
+    for (int64_t u = 0; u < C; ++u) {
+      float maxval = 0;
+      for (int64_t i = 0; i < Hq * Dh; ++i)
+        maxval = std::max(maxval, std::fabs(got[static_cast<size_t>(u * Hq * Dh + i)]));
+      if (maxval > 1e-3f) ++nonzero;
+    }
+    CHECK(nonzero == C);
+    (void)what;
+  };
+
+  // EAGER pass on generation A — warms every program the capture replays
+  // (identity Q 4D reshape, sdpa_decode B=2, out flatten reshape).
+  stage_kv(0x3c00, 0x3d00);
+  Tensor tq = commit_q();
+  Tensor to1 = make_out(mem_out1);
+  vt::PagedAttention(q, to1, tq, tkc, tvc, tbt, tsl, tqsl, args);
+  std::vector<float> out_eager = read_out(mem_out1);
+  check_out(out_eager, "eager A");
+
+  // Generation B: re-stage NEW K/V through RAC so the DEVICE paged-KV shadow
+  // moves to B while the HOST cache masters still hold A. A host-path PA
+  // (the decline's fallback) would attend the stale A bytes; only a
+  // capture-safe device PA can reproduce attention over B. An extra eager
+  // pass provides the reference.
+  stage_kv(0x3e00, 0x3f00);
+  Tensor tq_b = commit_q();
+  Tensor to1b = make_out(mem_out1);
+  vt::PagedAttention(q, to1b, tq_b, tkc, tvc, tbt, tsl, tqsl, args);
+  std::vector<float> out_eager_b = read_out(mem_out1);
+  check_out(out_eager_b, "eager B");
+
+  // CAPTURE pass: fresh Q shadow (rope refreshes it every step), identical PA
+  // inside the trace, then one replay.
+  Tensor tq2 = commit_q();
+  Tensor to2 = make_out(mem_out2);
+  be.BeginCapture(q);
+  struct EndOnExit {
+    Backend& b;
+    Queue& q;
+    ~EndOnExit() {
+      if (vt::tenstorrent::TraceCaptureActive()) {
+        try { b.EndCapture(q); } catch (...) {}
+      }
+    }
+  } end_guard{be, q};
+  vt::PagedAttention(q, to2, tq2, tkc, tvc, tbt, tsl, tqsl, args);
+  be.EndCapture(q);
+  be.Replay(q);
+  std::vector<float> out_replay = read_out(mem_out2);
+  check_out(out_replay, "replay");
+
+  // The captured+replayed PA must reproduce the eager attention over the
+  // DEVICE-resident generation-B KV — the capture served the device path.
+  int bad = 0;
+  for (size_t i = 0; i < out_eager_b.size(); ++i) {
+    if (std::fabs(out_eager_b[i] - out_replay[i]) >
+        0.05f * std::max(1.0f, std::fabs(out_eager_b[i]))) {
+      if (bad < 4) MESSAGE("replay mismatch i=", i, " eagerB=", out_eager_b[i],
+                           " replay=", out_replay[i]);
+      ++bad;
+    }
+  }
+  MESSAGE("replay-vs-eagerB mismatched elems: ", bad, "/", out_eager_b.size());
+  CHECK(bad == 0);
+
+  be.Free(mem_kc);
+  be.Free(mem_vc);
+  be.Free(mem_k);
+  be.Free(mem_v);
+  be.Free(mem_q);
+  be.Free(mem_out1);
+  be.Free(mem_out2);
+  be.Free(mem_slots);
+  be.Free(mem_bt);
+  be.Free(mem_sl);
+  be.Free(mem_qsl);
+}
 // BACKEND-TENSTORRENT-RESIDUAL-GOLDEN: op-level numerics probe at the
 // kDeviceResidualMinRows == 32 boundary. The device path (rows >= 32,
 // non-gemma) does ttnn::add + ttnn::rms_norm in bf16; the host/CPU path
@@ -9051,11 +9732,59 @@ TEST_CASE("kTENSTORRENT E=1 int8-dot keep-quant capture survives the 50 MiB trac
                     eager.size() * sizeof(float)) == 0);
   MESSAGE("capture x2 byte-identity: PASS; trace demand pass0=", demand[0],
           " B pass1=", demand[1], " B (region 52428800 B)");
+  // ── the 27B trace-fit gate (tt-launch-record-attribution-20260928) lives
+  // in the region-handoff case above, whose region 1 is the keepquant
+  // MatmulBT launch over the full grid ──
   backend.Free(mem_a);
   backend.Free(mem_w);
   backend.Free(mem_o);
 }
 
+// The 27B trace-fit fix's kernel-side derivation lock (host arm): the kernel
+// computes row0 = c*tcols and rowc = its clamp from the core coordinate
+// (c = y*grid_x + x, row-major — the mapping the deleted per-core
+// SetRuntimeArgs loop used). This case pins the two formulas to the SAME
+// values for every core, across shapes whose last core is partial and shapes
+// whose tail cores are fully idle. The DEVICE arm of this lock is the F32-out
+// capture case's byte-identity check, whose shape spans the grid with a
+// partial last core.
+TEST_CASE("keepquant int8-dot: in-kernel row0/rowc derivation equals the per-core host values") {
+  auto host_slice = [](uint32_t c, uint32_t tcols, uint32_t N) {
+    const uint32_t r0 = c * tcols;
+    const uint32_t rc =
+        r0 >= N ? 0u : std::min(tcols, N - r0);
+    return std::pair<uint32_t, uint32_t>{r0, rc};
+  };
+  auto kernel_slice = [](uint32_t core_x, uint32_t core_y, uint32_t grid_x,
+                         uint32_t tcols, uint32_t N) {
+    const uint32_t c = core_y * grid_x + core_x;
+    const uint32_t row0 = c * tcols;
+    const uint32_t rowc =
+        row0 >= N ? 0u : ((tcols < N - row0) ? tcols : (N - row0));
+    return std::pair<uint32_t, uint32_t>{row0, rowc};
+  };
+  const std::pair<uint32_t, uint32_t> shapes[] = {
+      {248320, 4096},  // the head shape: partial last core
+      {1508, 4096},    // N % tcols != 0 at several tail cores
+      {1024, 4096},    // exact tcols boundary, no partial core
+      {7, 4096},       // one partial group on core 0, idle tail
+  };
+  for (const auto& [N, tcols] : shapes) {
+    const uint32_t grid_x = 13, grid_y = 10;
+    for (uint32_t y = 0; y < grid_y; ++y) {
+      for (uint32_t x = 0; x < grid_x; ++x) {
+        const uint32_t c = y * grid_x + x;
+        const auto [hr0, hrc] = host_slice(c, tcols, N);
+        const auto [kr0, krc] = kernel_slice(x, y, grid_x, tcols, N);
+        CHECK_MESSAGE(kr0 == hr0, "row0 mismatch at N=" << N << " tcols="
+                                                        << tcols << " core " << c);
+        CHECK_MESSAGE(krc == hrc, "rowc mismatch at N=" << N << " tcols="
+                                                        << tcols << " core " << c);
+      }
+    }
+  }
+}
+
 // W4d W6: the BF16-OUT dispatch joined the int8-dot lever. The W4b landing
 // decision refused bf16-out because committing the kernel's f32 dev_out into
 // a bf16 slot left the slot holding f32 bytes at an f32 page geometry — the
@@ -10671,3 +11400,224 @@ TEST_CASE("kTENSTORRENT BFP8 vs bf16 resident-weight matmul timing (VT_TT_BFP8_B
             "), Bfp8MatmulUses=", vt::tenstorrent::Bfp8MatmulUses());
   }
 }
+
+// ─── tt-27b-region-capture: the region-handoff gate ─────────────────────────
+// The spec's red-first tests 1+2 at op scale: a TWO-REGION capture on the real
+// tt-metal trace backend where region 2 reads the buffer region 1 wrote — the
+// state binding across a region boundary — and each replay is BYTE-IDENTICAL
+// to the eager reference. The in-place discipline is the whole test: region 1
+// writes the PERSISTENT norm buffer in-region (the W3 commit shape), region 2
+// bakes that same address, so a replay chains through the boundary exactly as
+// the 27B per-layer regions will.
+//
+// RED-FIRST: on the pre-row tree the seam carries no per-region census
+// (`BreakableGraph::region_bytes()`), so this case does not compile there —
+// the capability it gates does not exist. The reviewer's MUTATION target is
+// the fresh-tensor defect the #3327 class names: give region 1 a FRESH output
+// buffer instead of the persistent one (allocate inside the capture) and this
+// case must FAIL — the replayed region 2 reads the address its capture baked,
+// which the fresh-tensor commit freed, and the output stops being the eager
+// bytes.
+//
+// Byte-exactness bar: MatmulBT/RmsNorm are deterministic kernels over fixed
+// device buffers — the same bar the int8-dot capture-x2 cases above assert.
+TEST_CASE("kTENSTORRENT region replay: state handoff across a region boundary, replay byte-identical to eager") {
+  if (!TenstorrentPresent()) {
+    MESSAGE("SKIPPED: no Tenstorrent device on this box");
+    return;
+  }
+  Backend& backend = vt::GetBackend(vt::DeviceType::kTENSTORRENT);
+  REQUIRE(backend.SupportsGraphCapture());
+  REQUIRE(vt::OpRegistered(vt::OpId::kRmsNorm, vt::DeviceType::kTENSTORRENT));
+  REQUIRE(vt::OpRegistered(vt::OpId::kMatmulBTQuant, vt::DeviceType::kTENSTORRENT));
+  Queue q = backend.CreateQueue();
+
+  // [1,H] bf16 state -> RmsNorm -> [1,H] bf16 (region 1) -> MatmulBT -> [1,N]
+  // f32 (region 2). H/N are small: this case owns the handoff mechanics, not
+  // kernel throughput.
+  constexpr int64_t kH = 512, kN = 1024;
+  const int64_t kQ6Elems = vt::BlockElems(vt::DType::kQ6_K);  // 256
+  const int64_t kQ6Bytes = vt::BlockBytes(vt::DType::kQ6_K);  // 210
+  const int64_t kNb = kH / kQ6Elems;
+
+  std::mt19937 rng(20260928u);
+  std::vector<uint16_t> x_bf(kH);
+  for (auto& v : x_bf) v = vt::F32ToBF16((static_cast<float>(rng() % 401) - 200.0f) / 100.0f);
+  std::vector<uint16_t> gamma_bf(kH);
+  for (auto& v : gamma_bf) v = vt::F32ToBF16(0.5f + static_cast<float>(rng() % 8) / 16.0f);
+  std::vector<uint8_t> w2_packed(static_cast<size_t>(kN) * kNb * kQ6Bytes);
+  for (size_t blk = 0; blk < w2_packed.size() / static_cast<size_t>(kQ6Bytes); ++blk) {
+    uint8_t* p = w2_packed.data() + blk * kQ6Bytes;
+    for (int i = 0; i < 208; ++i) p[i] = static_cast<uint8_t>(rng() & 0xFF);
+    const uint16_t d_bits = vt::F32ToF16(0.05f + 0.35f * static_cast<float>(rng() % 64) / 64.0f);
+    std::memcpy(p + 208, &d_bits, sizeof(d_bits));
+  }
+
+  void* mem_x = backend.Alloc(x_bf.size() * sizeof(uint16_t));
+  void* mem_norm = backend.Alloc(kH * sizeof(uint16_t));       // the persistent handoff buffer
+  void* mem_out = backend.Alloc(kN * sizeof(float));
+  void* mem_gamma = backend.Alloc(gamma_bf.size() * sizeof(uint16_t));
+  void* mem_w2 = backend.Alloc(w2_packed.size());
+  backend.Copy(q, mem_x, x_bf.data(), x_bf.size() * sizeof(uint16_t));
+  backend.Copy(q, mem_gamma, gamma_bf.data(), gamma_bf.size() * sizeof(uint16_t));
+  backend.Copy(q, mem_w2, w2_packed.data(), w2_packed.size());
+
+  Tensor x_t = Tensor::Contiguous(mem_x, vt::DType::kBF16,
+                                  Device{vt::DeviceType::kTENSTORRENT, 0}, {1, kH});
+  Tensor norm_t = Tensor::Contiguous(mem_norm, vt::DType::kBF16,
+                                     Device{vt::DeviceType::kTENSTORRENT, 0}, {1, kH});
+  Tensor gamma_t = Tensor::Contiguous(mem_gamma, vt::DType::kBF16,
+                                      Device{vt::DeviceType::kTENSTORRENT, 0}, {kH});
+  Tensor out_t = Tensor::Contiguous(mem_out, vt::DType::kF32,
+                                    Device{vt::DeviceType::kTENSTORRENT, 0}, {1, kN});
+  Tensor w2_t = Tensor::Contiguous(mem_w2, vt::DType::kQ6_K,
+                                   Device{vt::DeviceType::kTENSTORRENT, 0}, {kN, kH});
+
+  // ---- the eager reference: the same two calls, no capture ----
+  vt::RmsNorm(q, norm_t, x_t, gamma_t, vt::RmsNormArgs{1e-6f, false});
+  vt::MatmulBT(q, out_t, norm_t, w2_t);
+  std::vector<float> eager(static_cast<size_t>(kN), 0.0f);
+  backend.Copy(q, eager.data(), mem_out, eager.size() * sizeof(float));
+  for (float v : eager) CHECK(std::isfinite(v));
+
+  // ---- capture TWO regions and replay; the handoff must be invisible ----
+  // ONE capture here, replayed twice. The census below reads the backend's
+  // TRACE-STAGING byte level, whose release accounting is asynchronous across
+  // captures (a destroyed capture's staging drains after the blocking
+  // readback), so a second capture pass would subtract a stale level — the
+  // capture-x2 discipline is the keepquant capture cases' job; this case
+  // owns the handoff and the census, both of which are per-capture facts.
+  {
+    const int pass = 0;
+    vt::ResetGraphBreakStats();
+    vt::BreakableGraph graph;
+    {
+      vt::GraphCaptureScope scope(backend, q, graph,
+                                  vt::GraphCaptureMode::kPiecewise);
+      vt::RmsNorm(q, norm_t, x_t, gamma_t, vt::RmsNormArgs{1e-6f, false});
+      vt::GraphBreak();  // the REGION BOUNDARY: end segment 1, open segment 2
+      vt::MatmulBT(q, out_t, norm_t, w2_t);
+    }
+    REQUIRE(graph.captured());
+    const vt::GraphBreakStats stats = vt::GetGraphBreakStats();
+    CHECK(stats.segments_captured == 2);
+    // The per-region census (the spec's LastTraceBytes discipline): each
+    // region's staging must sit inside the 50 MiB budget a 27B layer region
+    // is sized against.
+    const std::vector<int64_t>& rb = graph.region_bytes();
+    REQUIRE(rb.size() == 2);
+    // The STAGING LEVEL this test starts from is whatever ~500 prior cases
+    // left (their captures' release accounting drains asynchronously), so
+    // region 0's delta can carry a stale subtraction. Region 1's delta is
+    // bounded by its own segment on both sides, and the handoff claim is the
+    // byte-exactness below, not the census sign.
+    MESSAGE("region 0: ", rb[0], " B; region 1: ", rb[1],
+            " B (budget 52428800 B)");
+    const bool in_budget = rb[1] > 0 && rb[1] <= 50 * 1024 * 1024;
+    CHECK_MESSAGE(in_budget, "region 1 staged " << rb[1]
+                                                << " B, outside the 50 MiB region budget");
+    // The config-page KB floor (the repro note's refutation): a stock
+    // full-grid program records 17,408 B/launch because identical per-core
+    // config pages collapse into the packed relay. The keepquant program's
+    // common-args + uniform-CB shape must reach the same floor: RED on the
+    // pre-fix tree at the measured 2,965,504 B, GREEN at KB scale once the
+    // per-launch dominant class (the all-encodings kernel binary carried by
+    // the runtime enc_sel, streamed paged-to-ring-buffer every launch) is
+    // removed.
+    CHECK_MESSAGE(rb[1] <= 64 * 1024,
+                  "region 1 (keepquant) recorded " << rb[1]
+                                                   << " B, over the 64 KiB config-page floor");
+    // The per-core RTA fix (docs/bench-evidence/tt-keepquant-rta-fix-
+    // 20260929.md §3) proved region 1's record here is NOT the per-core
+    // SetRuntimeArgs stream: it reads byte-identical 2,965,504 B before and
+    // after the fix (the RTA stream is ~228.9 MB of the 27B whole-graph
+    // demand, spent there). The record's dominant class is per-launch
+    // full-grid program command-sequence payload — a different lever. No KB
+    // floor gate lives here until that lever lands; the 27B bench leg is the
+    // fit arbiter.
+    graph.Replay(q);
+    std::vector<float> got(static_cast<size_t>(kN), 0.0f);
+    backend.Copy(q, got.data(), mem_out, got.size() * sizeof(float));
+    CHECK_MESSAGE(std::memcmp(got.data(), eager.data(), eager.size() * sizeof(float)) == 0,
+                  "capture pass " << pass << ": the TWO-REGION replay diverged "
+                                                     "from the eager reference at the region handoff");
+  }
+  backend.Free(mem_x);
+  backend.Free(mem_norm);
+  backend.Free(mem_out);
+  backend.Free(mem_gamma);
+  backend.Free(mem_w2);
+}
+
+// ─── tt-27b-region-capture: the capture-scope upload guard ──────────────────
+// The audit (docs/bench-evidence/tt-trace-record-audit-20260928.md) attributed
+// the 27B whole-graph 3,153,969,152 B trace demand to inline H2D payloads
+// recorded DURING capture: region 1's close was byte-exact 2,048 B of command
+// headers + 2 × 1,544,192 B of inline bf16 upload. The doctrine fix: the eager
+// pass warms every upload, and an upload route that fires under capture is
+// REFUSED by name. This case is red twice on the pre-fix tree: the unwarmed
+// capture-scope upload does not refuse (it silently inlines ~3.09 MB), and the
+// warmed capture still shows the audit's region-1 close instead of the ~2 KB
+// header floor region 0 measured.
+TEST_CASE("kTENSTORRENT capture-scope upload refuses and the warmed capture records the 2 KB floor") {
+  if (!TenstorrentPresent()) {
+    MESSAGE("SKIPPED: no Tenstorrent device on this box");
+    return;
+  }
+  Backend& backend = vt::GetBackend(vt::DeviceType::kTENSTORRENT);
+  REQUIRE(backend.SupportsGraphCapture());
+  Queue q = backend.CreateQueue();
+
+  // 1,544,192 bf16 elements — the byte-exact inline payload the audit measured
+  // in a region-1 close (3,088,384 = 2,048 headers + 2 × 1,544,192).
+  constexpr uint32_t R = 1024, C = 1508;  // 1024 × 1508 = 1,544,192
+  std::vector<uint16_t> host(static_cast<size_t>(R) * C);
+  for (size_t i = 0; i < host.size(); ++i)
+    host[i] = vt::F32ToBF16(0.125f * static_cast<float>(i % 17));
+  void* mem = backend.Alloc(host.size() * sizeof(uint16_t));
+  backend.Copy(q, mem, host.data(), host.size() * sizeof(uint16_t));
+  Tensor t = Tensor::Contiguous(mem, vt::DType::kBF16,
+                                Device{vt::DeviceType::kTENSTORRENT, 0}, {R, C});
+
+  // 1. An unwarmed EnsureDevice2D inside a capture scope is REFUSED by name —
+  //    the upload would be recorded inline into the trace.
+  {
+    vt::BreakableGraph g;
+    vt::GraphCaptureScope scope(backend, q, g, vt::GraphCaptureMode::kPiecewise);
+    bool refused = false;
+    std::string what;
+    try {
+      vt::tenstorrent::EnsureDevice2DForTest(t);
+    } catch (const std::exception& e) {
+      refused = true;
+      what = e.what();
+    }
+    CHECK_MESSAGE(refused,
+                  "the unwarmed EnsureDevice2D upload fired inside the capture "
+                  "scope without refusing — its payload would be inlined into "
+                  "the trace record");
+    CHECK_MESSAGE(what.find("capture") != std::string::npos,
+                  "refusal did not name the capture-scope upload: " << what);
+    CHECK_MESSAGE(what.find("refus") != std::string::npos,
+                  "refusal did not say it refused: " << what);
+  }
+
+  // 2. The warmed capture finds the tensor resident and records the header
+  //    floor (region 0's measured 2,048 B), not the inline payload.
+  vt::tenstorrent::EnsureDevice2DForTest(t);  // the eager warm pass
+  vt::BreakableGraph g;
+  {
+    vt::GraphCaptureScope scope(backend, q, g, vt::GraphCaptureMode::kPiecewise);
+    vt::tenstorrent::EnsureDevice2DForTest(t);
+    vt::GraphBreak();
+  }
+  const std::vector<int64_t>& rb = g.region_bytes();
+  REQUIRE(rb.size() >= 1);
+  MESSAGE("warmed capture region bytes: ", rb.back(),
+          " B (header floor 2048 B, inline payload would be 3090432 B)");
+  CHECK_MESSAGE(rb.back() <= 65536,
+                "the warmed capture still staged " << rb.back()
+                << " B — an upload (or its payload) rode inside the capture scope");
+
+  backend.Free(mem);
+}