diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md
index fc211c1..37f4215 100644
--- a/CONTRIBUTING.md
+++ b/CONTRIBUTING.md
@@ -40,7 +40,7 @@ those commits.
 **Foreign HIP** (leapdragon, a17t, anyone who is not this lab) keeps
 **their** Author **and** Committer. Same rule as extras
 [PR #1](https://github.com/opengfx1030/vllm-rdna/pull/1) (leapdragon
-`rdna_ar`, now dest extras @ `a4060647` squash / tip `2a5e893`; pick unique Aron Hsiao commits,
+`rdna_ar`, now dest extras @ `a4060647` squash / tip `700753d9`; pick unique Aron Hsiao commits,
 not the GitHub squash). Do not rewrite those as Blivion or Cursor.
 
 | Rule | Meaning |
diff --git a/docs/ARCHITECTURE.md b/docs/ARCHITECTURE.md
index 07b13c4..c82468d 100644
--- a/docs/ARCHITECTURE.md
+++ b/docs/ARCHITECTURE.md
@@ -200,7 +200,7 @@ One V1 id per tile (`HIPPIHX_V1_OP_*`). Serve wraps as
 `torch.ops.hippihx.<op>` — never a second Triton path in this library.
 Python mirror: `hippihx.v1` (`V1OpId`, `ABI_REVISION`). Caps include
 optional activation `dtype` (revision **2**). Qualnames `attention.*`
-(revision **3**; ids unchanged). Dest extras tip `2a5e89368272` still
+(revision **3**; ids unchanged). Dest extras tip `700753d9add5` still
 has no `torch.ops.hippihx.*` rewire — see [`BACKPORT.md`](BACKPORT.md).
 Do not edit `opengfx1030/vllm-rdna` from this tree.
 
diff --git a/docs/BACKPORT.md b/docs/BACKPORT.md
index bf70acc..740d440 100644
--- a/docs/BACKPORT.md
+++ b/docs/BACKPORT.md
@@ -1,7 +1,7 @@
 # extras → hippihx backport review
 
 Lock: [`opengfx1030/vllm-rdna`](https://github.com/opengfx1030/vllm-rdna)
-`rdna_extras` @ `2a5e89368272` (2026-09-24 21:52 UTC). Dest **default
+`rdna_extras` @ `700753d9add5` (2026-09-25 05:48 UTC). Dest **default
 branch is `rdna_extras`**, not `main` (`c00091e02670` upstream vLLM — no
 dest HIP). **No** `torch.ops.hippihx.*`. Observe only: no `.cu` dump, no
 tok/s, no extras maxdiff.
@@ -32,9 +32,9 @@ until a body migrates and extras binds it.
 | `3bddd3c9a5d1` | Flash-Next production launcher (`serve_gfx1030_flashnext.sh`) | extras scripts. PIECEWISE + seq cap 6. leapdragon `VLLM_RDNA_AR` unset (still **0**). Zoo still no Finegrained. |
 | `50120e13b468` | HC `_contig()` capture-safe per-shape cache | extras product Python. Gate still **off**. Remaining same-shape clobber. No V1 op. |
 | `5c4ab9891910` | dest-reverted gfx1030 wvSplitK decode port | extras dense GEMM. Kernel asserts on gfx1030 under capture. Keeps `gemv_f16_rdna2` for decode `M<=8`. **No zoo tile.** Do not reintroduce `gemv_f16`. |
-| `e45dd5cb2de8` | QSA `CircularBufferManager` empty-ring prefix hits | extras v1 core. Observe `attention/qsa_indexer`. Empty compression-group ring is valid. Do not copy tok/s. |
+| `e45dd5cb2de8` | QSA `CircularBufferManager` empty-ring prefix hits | extras v1 core. Observe `attention/qsa_indexer`. Empty compression-group ring is valid. |
 | `c59a23f625e0` | recovered extras probes / benches / HC Triton WIP | extras ops. HC Triton + graph-keepalive diagnostics **not** on the serve path. Do not dump. |
-| `31003ff` / `ae5c7edc` / `37345ee9` / `aaa4775a` / `0dd38115` / `657bdba8` / `9c602943` / `609c9c0d` | Flash-Next launcher serve knobs | extras scripts. Vision-on, then `609c9c0d` FULL_AND_PIECEWISE (ROCm executes as PIECEWISE). Stay extras. Do not copy tok/s. |
+| `31003ff` / `ae5c7edc` / `37345ee9` / `aaa4775a` / `0dd38115` / `657bdba8` / `9c602943` / `609c9c0d` | Flash-Next launcher serve knobs | extras scripts. Vision-on, then `609c9c0d` FULL_AND_PIECEWISE (ROCm executes as PIECEWISE). Stay extras. |
 | `741e5bc31ae5` | mamba spec-decode tables index by `req_idx` | extras worker (upstream vLLM #55506 port; Author Karl0007). Persistent per-request-slot tables. V1 `req_idx == batch_idx`. Stay extras. |
 | `3b59ee16e553` | extras PR **#13** squash (T44b `rdna_ar` VRAM flags + wedge) | Later `comm.pcie`. Still opt-in (`VLLM_RDNA_AR=0`). Dest extras `MAX_KB` default **64**. Zoo lock **512**. Do **not** pick squash (Cursor rewrite). |
 | `dbb1e7764aba` | extras PR **#14** merge (V620 Triton MoE JSON / ROCR amdsmi / PLE fp8) | Stay extras. HIP MoE ignores the JSON. Cursor rewrite — do not pick. |
@@ -42,18 +42,19 @@ until a body migrates and extras binds it.
 | `3d6df9ed617a` | extras `exl3_dot2_*` `__HIP__RDNA__` guard for docker multi-arch | Stay extras. Zoo still one `--offload-arch` per fatbin. gfx1150 / gfx12xx not dest. Do not dump. |
 | `ed94e3f3d299` | Qwen4Exp MTP proposer allowlist + skinny `w2_zp` | Stay extras. MTP still not dest. No HIP body. |
 | `b33f9b66eb2b` | `get_device_name` torch fallback when amdsmi has no handles | Stay extras. Platform only. |
-| `f3dd65fa7063` | extras PR **#15** merge (QSA live-context bound + folded **#12** PLE/MTP/graph-redirect) | Stay extras. Python/Triton/serve. No HIP body. Foreign George + Codex. Do not pick merge. Do not copy tok/s. |
-| `e131562936f9` | extras PR **#17** merge (resident W4A16 MoE / Mamba #55450 retire) | Stay extras. ATen HIP (`moe_resident_decode.cu`). `VLLM_RDNA_MOE_RESIDENT*` default **off**. Dest `VLLM_RDNA_AR` still **0**. Foreign George + Codex. Do not pick merge. Do not dump. Do not copy tok/s. |
-| `48c56efbe80b` | pin `AttentionConfig.backend = RDNA_ATTN` when `VLLM_USE_RDNA2_FA=1` | Stay extras. Python platform (`vllm/platforms/rocm.py`). API-server env; workers inherit config. Dest `VLLM_RDNA_AR` still **0**. Unique Blivion. No HIP. Do not copy tok/s. |
-| `68a635ed8d78` | extras **#20** capture + dest-integrated **#21** FULL+PIECEWISE | Stay extras. Python/serve. Keeps FULL decode graphs; piecewise for mixed/prefill. Foreign Waldecir + Cursor. Do not pick merges. Do not copy tok/s. Dest `VLLM_RDNA_AR` still **0**. Zoo does not own graph mode. |
-| `2a5e89368272` | extras PR **#19** merge (MTP unquantized-weight detect) | Stay extras. Python only (`qwen3_5_mtp.py`). MTP still not dest. Foreign a17t + opencode. Do not pick merge. Do not copy tok/s. Dest `VLLM_RDNA_AR` still **0**. |
+| `f3dd65fa7063` | extras PR **#15** merge (QSA live-context bound + folded **#12** PLE/MTP/graph-redirect) | Stay extras. Python/Triton/serve. Foreign George + Codex. Do not pick. |
+| `e131562936f9` | extras PR **#17** merge (resident W4A16 MoE / Mamba #55450 retire) | Stay extras. ATen HIP. `VLLM_RDNA_MOE_RESIDENT*` default **off**. Foreign George + Codex. Do not pick. |
+| `48c56efbe80b` | pin `AttentionConfig.backend = RDNA_ATTN` when `VLLM_USE_RDNA2_FA=1` | Stay extras. Python platform. Unique Blivion. |
+| `68a635ed8d78` | extras **#20** capture + dest-integrated **#21** FULL+PIECEWISE | Stay extras. Keeps FULL decode graphs; piecewise mixed/prefill. Foreign Waldecir + Cursor. Zoo does not own graph mode. |
+| `2a5e89368272` | extras PR **#19** merge (MTP unquantized-weight detect) | Stay extras. Python only. MTP still not dest. Foreign a17t + opencode. Do not pick. |
+| `700753d9add5` | MTP draft unbreak (amdsmi shutdown swallow + local-argmax) | Stay extras. Python platform + MTP model. MTP still not dest. Unique Blivion. |
 
 Other dest-fixed ISA (keep in tile locks, not a dump): GDN prefill `o`
 `i_t_local`; causal_conv FIR pre-shift then shift; ConfigH
 (`K_STEP=64`) reverted. Live dest bugs: [Dest extras defects](#dest-extras-defects).
 
-Open extras PRs **#2** (GLM Later), **#18**. Draft **#9/#10**
-sdot — skip. Closed **#3/#16** unmerged, **#21** dest-integrated
+Open extras PRs **#2** (GLM Later), **#18**, **#22**, **#24**. Draft **#9/#10**
+sdot, **#23** PCIe P2P KV — skip. Closed **#3/#16** unmerged, **#21** dest-integrated
 unmerged, **#11** dest-reverted wvSplitK (**no zoo tile**), **#12**
 superseded by dest **#15**. Merged **#13/#14/#15/#17/#19/#20** dest
 *presence*, observe, do not pick.
@@ -70,7 +71,7 @@ superseded by dest **#15**. Merged **#13/#14/#15/#17/#19/#20** dest
 | leapdragon `rdna_ar` + PIX helpers + dest T44b (`3b59ee16`) | **Later** `comm.pcie`. Still opt-in. Dest extras `MAX_KB` **64**; zoo **512**. Do not pick squash. |
 | PR **#2** GLM-5.3 KDA/DSA | **Later** `kda_scan` / `dsa_nope` / `qsa_indexer`. Do not name tiles `glm5_*`. |
 | GDN arenas, `rdna2_graph_keepalive.cuh`, breakable cudagraphs, Hybrid W4 gfx10, Flash-Next HC/QSA/PLE/M-RoPE HIP, skinny GEMM / `gemv_f16_rdna2`, PLE schema, W4 MoE oracle, `new_zeros`/`zeros_like`, V1 FULL→PIECEWISE, vLLM custom AR, `bench_report.py`, seq cap 6, Flash-Next launcher knobs, HC `_contig()` cache, QSA Triton bounds, mamba spec-decode, T44b wedge, V620 MoE JSON / amdsmi / PLE, ROCm platform init, extras EXL3 docker arch-guard, Qwen4Exp MTP, recovered extras probes, resident MoE, FA RDNA_ATTN pin, `qwen4_exp/**` | **Stay extras.** Product gates default off (S6 revert). No new V1 op until dest locks a class. |
-| PR **#3** a17t / explore **#9/#10** sdot / closed **#12** / closed **#16** / PR **#18** | **Skip.** Second W4 family, not dest, serve-only. **#16** unmerged. **#18** GPTQ `BLOCK_KN_SIZE` 256. |
+| PR **#3** a17t / explore **#9/#10** sdot / closed **#12** / closed **#16** / PR **#18** / PR **#22** / PR **#23** / PR **#24** | **Skip.** Second W4 family, not dest, serve-only. **#16** unmerged. **#18** GPTQ `BLOCK_KN_SIZE` 256. **#22** two-shot `rdna_ar` (Cursor; library still **0**; zoo **512**; do not dump). **#23** PCIe P2P KV (Cursor; HIP IPC SDMA; zoo `comm.pcie` is AR, not KV disagg). **#24** RAM KV offload (George/Codex; no HIP). |
 | PR **#5** / **#8** Flash-Next, extras **#11** wvSplitK, extras **#13** T44b, extras **#14**, extras **#15**, extras **#17**, extras **#19**, extras **#20**, extras **#21** | **Closed.** Dest-landed T44b / #14 / #15 / #17 / #19 / #20 / dest-integrated #21 are observe-only. Dest reverted wvSplitK (`5c4ab989`). Do not re-merge. Do not grow `gemv_f16`. |
 
 Bodies stay in extras because every dest HIP file includes `torch/all.h`.
@@ -78,7 +79,7 @@ A dump is not a migrate. CONTRIBUTING: do not edit extras from this repo.
 
 ## Dest file → tile
 
-Observed at `2a5e89368272`. Numbers are extras observations, not dest
+Observed at `700753d9add5`. Numbers are extras observations, not dest
 locks. Fill tile READMEs.
 
 | extras path | hippihx tile | Notes |
@@ -108,7 +109,7 @@ extras consume.
    journals FPP13 16k c=8 green via serve arenas — still extras.
 2. Tile README LDS / `__launch_bounds__` / wave / DOT unit are filled.
 3. hippihx ships one HIP entry extras can bind. **Started:** V1
-   `plan`/`run`. `run` is `NOT_READY`. extras @ `2a5e89368272` has no
+   `plan`/`run`. `run` is `NOT_READY`. extras @ `700753d9add5` has no
    `torch.ops.hippihx.*`.
 4. extras is rewired **in extras**. The extras copy is then deleted.
 
@@ -117,10 +118,10 @@ Until then: observe, lock numbers, keep stubs. Tracker:
 
 ## Dest extras defects
 
-Dest tip `2a5e89368272`. Live dest bugs / rolled-back paths — hippihx
+Dest tip `700753d9add5`. Live dest bugs / rolled-back paths — hippihx
 must not reproduce them.
 
-| Defect | Where | Status @ `2a5e893` | Zoo lock |
+| Defect | Where | Status @ `700753d9` | Zoo lock |
 |---|---|---|---|
 | `llvm.amdgcn.fdot2.bf16.bf16` ISel abort | Hybrid W4 gfx10 + bf16. Dest HIP has **no** `fdot2.bf16`. | **Serve-mitigated** (`59237b3`). | Never `fdot2.bf16`. DOT + GDN HIP = **fp16 act**. |
 | Scale-baked W4 zero-point | `qdq_4_rdna2.cuh` `prep_zero_scale_fp16`: `0xE400 \| zero` then `scale * (-1024 - zero)` in `half`. | **Still live** | Integer `q - zero`, then `* scale`. |
@@ -183,4 +184,4 @@ This review is documentation, not a kernel migrate.
 | PR #3 a17t unique W4 (`d53572644`, later Simon Siebert) | **Not taken** | Closed unmerged. If dest ever locks that family, pick **their** commits |
 | Explore PRs **#9/#10** sdot | **Not taken** | Not dest |
 | extras dest-presence serve (**#13** T44b / **#14** / **#15** / **#17** / **#19** / **#20** / dest-integrated **#21** / mamba `741e5bc3` / closed **#12**) | observe | Python/Triton/serve + dest-landed ATen HIP. No HIP migrate. Do not pick Cursor/George/Codex/Karl/Waldecir/a17t rewrites. Unique Aron Hsiao `rdna_ar` still pickable. |
-| extras not dest (**#16** closed / **#18**) | **Not taken** | GPTQ `BLOCK_KN_SIZE` 256. |
+| extras not dest (**#16** closed / **#18** / **#22** / **#23** / **#24**) | **Not taken** | GPTQ `BLOCK_KN_SIZE` 256. Two-shot `rdna_ar`. PCIe P2P KV. RAM KV offload. |
diff --git a/docs/CONSUME.md b/docs/CONSUME.md
index 1bd983b..5d3f567 100644
--- a/docs/CONSUME.md
+++ b/docs/CONSUME.md
@@ -28,7 +28,7 @@ fa_fdot2.run(binding)
 Register **one** `torch.ops.hippihx.<op>` that calls `hippihx_v1_run`.
 No Triton→HIP double-fire. Bind keys on **arch + wave**, not GFX name
 alone. extras V1 consume is the HIP fatbin. Dest extras tip
-`2a5e89368272` still has no `torch.ops.hippihx` bind. FlyDSL extras
+`700753d9add5` still has no `torch.ops.hippihx` bind. FlyDSL extras
 consume waits on `FLYDSL_V1_CONSUME` (graph-safe JIT) — see
 [`FLYDSL.md`](FLYDSL.md).
 
diff --git a/docs/EXTRAS.md b/docs/EXTRAS.md
index dc02350..a12093c 100644
--- a/docs/EXTRAS.md
+++ b/docs/EXTRAS.md
@@ -7,12 +7,13 @@ one V1 op. Review: [`BACKPORT.md`](BACKPORT.md).
 ## Unvalidated extras inventory
 
 Snapshot of [`opengfx1030/vllm-rdna`](https://github.com/opengfx1030/vllm-rdna)
-`rdna_extras` @ `2a5e89368272` (2026-09-24 21:52 UTC). Dest default
+`rdna_extras` @ `700753d9add5` (2026-09-25 05:48 UTC). Dest default
 branch is **`rdna_extras`**. `main` is upstream vLLM `c00091e02670`.
 Merged extras [PR #1](https://github.com/opengfx1030/vllm-rdna/pull/1)
-squash is `a4060647cfbb`. Open **#2**, **#18**. Closed **#3/#16**
-unmerged; **#21** dest-integrated unmerged; **#12** superseded by
-**#15**. Merged **#13**, **#14**, **#15**, **#17**, **#19**, **#20**.
+squash is `a4060647cfbb`. Open **#2**, **#18**, **#22**, **#24**. Draft **#23**.
+Closed **#3/#16** unmerged; **#21** dest-integrated unmerged; **#12**
+superseded by **#15**. Merged **#13**, **#14**, **#15**, **#17**, **#19**,
+**#20**.
 **No** `torch.ops.hippihx.*`.
 
 **Unvalidated.** Not dest. Not silicon-signed. No tok/s. hippihx still
@@ -50,6 +51,9 @@ Status: **extras** = live on dest tip · **Later** = side branch ·
 | leapdragon push AR | `rdna_allreduce.{cu,cuh}` (merged PR **#1**) | `comm/pcie` | Dest extras, default **off**. Occupancy pin closed. **unvalidated** |
 | a17t extra AWQ GEMM / GEMV | `awq_gemm_rdna2.cu` etc. (closed PR **#3**) | — | **skip** — closed unmerged, second W4 family |
 | GPTQ exllama `BLOCK_KN_SIZE` 256 | `q_gemm.cu` (PR **#18**) | — | **skip** — not dest DOT W4 |
+| two-shot `rdna_ar` | `rdna_allreduce.{cu,cuh}` (PR **#22**) | — | **skip** — not dest. Library `VLLM_RDNA_AR` still `"0"`. Zoo `AR_MAX_KB` **512**. Do not dump |
+| PCIe P2P KV disagg | `PcieP2pConnector` (draft PR **#23**) | — | **skip** — not dest. HIP IPC SDMA. Zoo `comm.pcie` is AR, not KV. Do not dump |
+| RAM KV offload | offloading connector (PR **#24**) | — | **skip** — not dest. George/Codex. Serve only. No HIP. Do not pick |
 | Explore W4A8 sdot4 / W4A4 sdot8 | extras PRs **#9/#10** | — | **skip** — not dest |
 | Resident W4A16 MoE skinny decode | `moe_resident_decode.cu` (dest **#17** @ `e1315629`) | `moe/routed` (watch) | Stay extras. Opt-in `VLLM_RDNA_MOE_RESIDENT*`. ATen HIP. Do not dump. Closed **#16** unmerged. **unvalidated** |
 
@@ -128,23 +132,15 @@ stay extras (no D2H under capture in the zoo).
 
 ### Not taken / leave in extras
 
-Capture plumbing (`torch.zeros` / `new_zeros` / `zeros_like`, persist
-keepalive, GDN arenas, `eager_break_during_capture`), product serve
-(`qwen4_exp/**`, TunableOp, PLE offload), skinny GEMM, ConfigH, V1
-FULL→PIECEWISE, vLLM custom AR, `bench_report.py`, Qwen4Exp HIP gates /
-wrappers / HC compute / `_contig()` cache, seq cap 6, Flash-Next
-launcher knobs, dest-reverted wvSplitK, QSA Triton bounds, recovered
-extras ops, mamba spec-decode `req_idx`, dest T44b `rdna_ar` (still
-opt-in; do not pick Cursor squash), dest PR **#14** (HIP MoE ignores
-the JSON; do not pick), dest **#17** (resident MoE ATen HIP; do not
-dump; do not pick), dest PR **#20** / dest-integrated PR **#21** (FULL+PIECEWISE
-serve; zoo does not own graph mode; do not pick; do not copy tok/s),
-dest PR **#19** (MTP unquantized-weight detect; MTP still not dest;
-do not pick a17t/opencode), a17t PR **#3** (closed unmerged), PR **#18**
-(GPTQ `BLOCK_KN_SIZE` 256),
-closed **#5/#11/#13/#14/#15/#17/#19/#20/#21**, explore **#9/#10**, closed PR **#12**
-(superseded by **#15**), closed PR **#16**, ROCm platform init, extras
-EXL3 docker arch-guard (one `--offload-arch` per fatbin; gfx1150 /
-gfx12xx not dest), Qwen4Exp MTP (not dest), amdsmi `get_device_name`
-fallback, dest FA `RDNA_ATTN` pin (`48c56ef`). Produce (`-cb 3inst`,
-AWQ pack) stays outside hippihx.
+Stay extras: capture plumbing, product serve, skinny GEMM, ConfigH,
+graph mode, vLLM custom AR, `bench_report.py`, dest T44b `rdna_ar`
+(still opt-in), dest PR **#13** / PR **#14** / PR **#15** / PR **#17** /
+PR **#19** / PR **#20** (observe, do not pick), dest-integrated PR **#21**,
+dest FA `RDNA_ATTN` pin, resident MoE ATen HIP, Qwen4Exp MTP (not dest),
+ROCm platform init, extras EXL3 docker arch-guard. Skip: a17t PR **#3**,
+PR **#18** GPTQ `BLOCK_KN_SIZE` 256, PR **#22** two-shot `rdna_ar`
+(Cursor; library `VLLM_RDNA_AR` still `"0"`; zoo `AR_MAX_KB` **512**;
+do not dump), draft PR **#23** PCIe P2P KV (Cursor; HIP IPC SDMA; do not
+dump), PR **#24** RAM KV offload (George/Codex; no HIP; do not pick),
+explore **#9/#10**, closed PR **#12**, closed PR **#16**,
+closed **#5/#11**. Produce stays outside hippihx.
diff --git a/tests/test_no_b12x_cuda.py b/tests/test_no_b12x_cuda.py
index 07f6371..6c44068 100644
--- a/tests/test_no_b12x_cuda.py
+++ b/tests/test_no_b12x_cuda.py
@@ -112,7 +112,7 @@ def test_readme_unvalidated_inventory() -> None:
     assert "e131562" in extras
     assert "48c56ef" in extras
     assert "68a635" in extras
-    assert "2a5e893" in extras
+    assert "700753d9" in extras
     assert "e45dd5c" in extras
     assert "d1b200" in extras
     assert "1046782" in extras
@@ -146,6 +146,7 @@ def test_readme_unvalidated_inventory() -> None:
     assert "48c56ef" in bp
     assert "68a635" in bp
     assert "2a5e893" in bp
+    assert "700753d9" in bp
     assert "dbb1e77" in bp
     assert "e45dd5c" in bp
     assert "d1b200" in bp
@@ -181,13 +182,17 @@ def test_readme_unvalidated_inventory() -> None:
     assert "PR **#19**" in extras or "PR **#19**" in bp
     assert "PR **#20**" in extras or "PR **#20**" in bp
     assert "PR **#21**" in extras or "PR **#21**" in bp
+    assert "PR **#22**" in extras or "PR **#22**" in bp
+    assert "PR **#23**" in extras or "PR **#23**" in bp
+    assert "PR **#24**" in extras or "PR **#24**" in bp
     assert "VLLM_RDNA_MOE_RESIDENT" in extras
     assert "moe_resident" in extras
     contrib = (ROOT / "CONTRIBUTING.md").read_text(encoding="utf-8")
-    assert "2a5e893" in contrib
+    assert "700753d9" in contrib
     assert "a4060647" in contrib
     # CONTRIBUTING is tip-only. Historical dest SHAs live in BACKPORT.
     for token in (
+        "2a5e893",
         "68a635",
         "48c56ef",
         "e131562",
diff --git a/tiles/attention/fa_fdot2/README.md b/tiles/attention/fa_fdot2/README.md
index 26577bd..3d87b0d 100644
--- a/tiles/attention/fa_fdot2/README.md
+++ b/tiles/attention/fa_fdot2/README.md
@@ -33,11 +33,11 @@ body yet — when you do, cherry-pick BlivionIaG, do not re-author
 | `__launch_bounds__` | extras @ `a4060647`: decode D=128 `__launch_bounds__(128)`; decode D=256 + GQA-aware D=256 `__launch_bounds__(256)`; prefill `THREADS_PREFILL=128` | **not dest-locked** — occupancy pin closed |
 | GQA D=256 decode | extras @ `6c5ff94`: idle waves still shuffle; skip online-softmax when `m_new == -inf` (NaN guard); zero `O_partial` on empty seq. `GQA_MAX_G=8`, `GQA_BC=32`, `GQA_DSK=256+8` | observed; not dest-locked |
 | fp16 flash KV writer | extras: `reshape_and_cache_flash_rdna2`, `__launch_bounds__(128, 4)`, used for non-native KV write, fp16 only | extras HIP; do not dump ATen wrapper. Persist workspaces stay extras. |
-| GQA-subgroup prefill | extras @ `ecfec4e412ad`: `fa_rdna2_prefill_paged_varlen_gqa`, `HEADS_PER_CTA=2` / `BR=8`, dual-acc fdot2. Env `VLLM_FA_RDNA2_GQA_MODE` default `subgroup`. True-GQA (`HEADS_PER_CTA=6`) dropped. Dest @ `d1b200b1`: GQA prefill O accumulator is register-resident (was smem). Occupancy pin closed. Do not dump `.cu`. Do not copy tok/s. | observed; occupancy pin closed. Do not copy tok/s. |
+| GQA-subgroup prefill | extras @ `ecfec4e412ad`: `fa_rdna2_prefill_paged_varlen_gqa`, `HEADS_PER_CTA=2` / `BR=8`, dual-acc fdot2. Env `VLLM_FA_RDNA2_GQA_MODE` default `subgroup`. True-GQA (`HEADS_PER_CTA=6`) dropped. Dest @ `d1b200b1`: GQA prefill O accumulator is register-resident (was smem). Occupancy pin closed. Do not dump `.cu`. | observed; occupancy pin closed. |
 
 Dest extras @ `48c56ef` pins `AttentionConfig.backend = RDNA_ATTN` from
 the API-server `VLLM_USE_RDNA2_FA` env so workers actually select
-FA-RDNA2. Stay extras. No HIP. Do not copy tok/s.
+FA-RDNA2. Stay extras. No HIP.
 
 Scratch is sized by `plan`. C consume id: `HIPPIHX_V1_OP_ATTN_FA_FDOT2`
 (`include/hippihx/v1.h`). Serve wraps as one `torch.ops.hippihx.*` — no