diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md
index 9b3a733..151593e 100644
--- a/CONTRIBUTING.md
+++ b/CONTRIBUTING.md
@@ -48,7 +48,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`; dest squash `a4060647` / tip `ac9dfd8f`; pick unique Aron
+`rdna_ar`; dest squash `a4060647` / tip `bc5fbee5`; pick unique Aron
 Hsiao commits, not the GitHub squash). Do not rewrite those as Blivion
 or Cursor.
 
diff --git a/docs/ARCHITECTURE.md b/docs/ARCHITECTURE.md
index e81c279..52ebed7 100644
--- a/docs/ARCHITECTURE.md
+++ b/docs/ARCHITECTURE.md
@@ -234,7 +234,7 @@ optional activation `dtype` (revision **2**). Qualnames `attention.*`
 plan, params, tensor descriptors, the stream and `caps.fabric`. Per-op
 params, tensor slots and scratch rules live in the catalog, and
 `python -m hippihx._lib.codegen` writes the header enums and C tables.
-See [`CONSUME.md`](CONSUME.md). Dest extras tip `ac9dfd8fba6d` still
+See [`CONSUME.md`](CONSUME.md). Dest extras tip `bc5fbee5d0a1` still
 has no `torch.ops.hippihx.*` rewire. Dest extras **#27** is opt-in
 `VLLM_HIPPIHX` ctypes consume (default **off**; plans not ready) — see
 [`BACKPORT.md`](BACKPORT.md).
diff --git a/docs/BACKPORT.md b/docs/BACKPORT.md
index 346208e..f7910c0 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` @ `ac9dfd8fba6d` (2026-10-01 11:35 UTC). Dest **default
+`rdna_extras` @ `bc5fbee5d0a1` (2026-10-01 13:08 UTC). Dest **default
 branch is `rdna_extras`**, not `main` (`c00091e02670` upstream vLLM — no
 dest HIP). **No** `torch.ops.hippihx.*`. Dest extras **#27** is opt-in
 `VLLM_HIPPIHX` ctypes consume (default **off**; plans not ready). Observe
@@ -69,6 +69,7 @@ until a body migrates and extras binds it.
 | `a5059cf0480e` | extras **#32** dest-presence (EXL3 mul1 / K=1..8 + MoE staging bound) | Stay extras. ATen HIP. Unique Blivion dest landing. GitHub PR closed-unmerged. Do not dump. Do not copy tok/s. |
 | `30b0bd4e3604` | dest EXL3 v0.30 loader + fused project + MoE method + PLE n-gram | Stay extras. Python + dest HIP `exl3_project_rdna2`. Unique Blivion. extras **#32** GitHub closed-unmerged. Do not dump. Do not copy tok/s. |
 | `ac9dfd8fba6d` | dest TunableOp unify + EXL3 27B rows + gemma_rms fake | Stay extras. Serve scripts / rocBLAS tables / inductor fake. Unique Blivion. Do not dump. Do not copy tok/s. |
+| `bc5fbee5d0a1` | dest serve_rdna recipes + EXL3 MTP draft + TunableOp profiles | Stay extras. Python/serve. Unique Blivion. Dest now serves EXL3 quantized MTP; no zoo MTP tile. Do not dump. Do not copy tok/s. |
 
 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
@@ -84,7 +85,7 @@ unmerged, **#11** dest-reverted wvSplitK (**no zoo tile**), **#12**
 superseded by dest **#15**. Merged
 **#13/#14/#15/#17/#19/#20/#22/#24/#26/#27/#28/#29** dest *presence*;
 dest **#30** @ `3fb9d43a` / dest **#32** @ `a5059cf0` dest-presence.
-Dest tip `ac9dfd8f` is TunableOp/serve (not a new extras PR).
+Dest tip `bc5fbee5` is serve recipes + EXL3 MTP (not a new extras PR).
 Observe, do not pick. `rdna_extras_wip_20260910` is not dest.
 
 ## Action
@@ -97,7 +98,7 @@ Observe, do not pick. `rdna_extras_wip_20260910` is not dest.
 | GDN decode fp16 SSM state | **Later** `attention/gdn_scan`. Recurrence stays 16 fp32 VGPR. |
 | leapdragon `rdna_ar` + PIX helpers + dest T44b (`3b59ee16`) | **Later** `comm.pcie`. Still opt-in. Dest extras `MAX_KB` **64** (two-shot dest-off @ `101a16c8`). 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_*`. |
-| Serve, product HIP, dest-landed extras PRs, MTP, graph mode, `new_zeros`/`zeros_like`, `qwen4_exp/**` | **Stay extras.** Product gates default off (S6 revert). Dest **#27** opt-in ctypes consume; plans not ready. Dest **#28** / dest **#29** split decode / dest row-gate (`e0112c55`) / dest W4A8 (`3a0786ea`) / dest **#30** compile-dispatch (`3fb9d43a`) / dest MoE epilogue (`9cad8792`) / dest **#32** EXL3 mul1 (`a5059cf0`) / dest EXL3 loader (`30b0bd4e`) / dest TunableOp (`ac9dfd8f`) stay extras; do not dump. No body migrate. See SHA index. |
+| Serve, product HIP, dest-landed extras PRs, MTP, graph mode, `new_zeros`/`zeros_like`, `qwen4_exp/**` | **Stay extras.** Product gates default off (S6 revert). Dest **#27** opt-in ctypes consume; plans not ready. Dest **#28** / dest **#29** split decode / dest row-gate (`e0112c55`) / dest W4A8 (`3a0786ea`) / dest **#30** compile-dispatch (`3fb9d43a`) / dest MoE epilogue (`9cad8792`) / dest **#32** EXL3 mul1 (`a5059cf0`) / dest EXL3 loader (`30b0bd4e`) / dest TunableOp (`ac9dfd8f`) / dest EXL3 MTP + `serve_rdna` (`bc5fbee5`) stay extras; do not dump. No body migrate. See SHA index. |
 | PR **#3** a17t / closed **#9** / explore **#10** sdot8 / closed **#12** / closed **#16** / PR **#18** / PR **#23** / PR **#25** / PR **#31** / PR **#33** / PR **#34** | **Skip.** Second W4 family, not dest, serve-only. **#9** closed-unmerged (superseded by dest W4A8). **#16** unmerged. **#18** GPTQ `BLOCK_KN_SIZE` 256. **#23** PCIe P2P KV (HIP IPC SDMA; zoo `comm.pcie` is AR, not KV disagg). **#25** rdna_ar retained-output (George/Codex). **#31** W4 exact-dequant explore. **#33** FA decode scores. **#34** FA D=128 prefill. |
 | PR **#5** / **#8** Flash-Next, extras **#11** wvSplitK, extras **#13** T44b, extras **#14**, extras **#15**, extras **#17**, extras **#19**, extras **#20**, extras **#21**, extras **#22**, extras **#24**, extras **#26**, extras **#27**, extras **#28**, extras **#29** | **Closed.** Dest-landed observe-only. Dest reverted wvSplitK (`5c4ab989`). Do not re-merge. Do not grow `gemv_f16`. |
 
@@ -106,7 +107,7 @@ A dump is not a migrate. CONTRIBUTING: do not edit extras from this repo.
 
 ## Dest file → tile
 
-Observed at `ac9dfd8fba6d`. Numbers are extras observations, not dest
+Observed at `bc5fbee5d0a1`. Numbers are extras observations, not dest
 locks. Fill tile READMEs.
 
 | extras path | hippihx tile | Notes |
@@ -120,12 +121,13 @@ locks. Fill tile READMEs.
 | `csrc/rocm/moe_q_gemm_rdna2.cu` + `moe_accum_rdna2.cuh` | `moe/routed` | Reuses W4 helpers. Dest @ `e1315629` dequant/eight-row stay extras. Dest @ `9cad8792` packed-fp16 CAS default; `VLLM_RDNA2_MOE_FP32_ACCUM` opt-in. Do not dump. |
 | `csrc/rocm/moe_resident_decode.cu` | `moe/routed` (watch) | Dest @ `e1315629`. ATen HIP. `VLLM_RDNA_MOE_RESIDENT*` default **off**. Stay extras. Do not dump. |
 | `csrc/rocm/exl3_dot2_{dense,moe,dequant,hadamard}.*` | `gemm/exl3_3inst` | `LDS_PAD=8`. Dest @ `3d6df9ed` widened extras compile guard for docker multi-arch. Dest **#32** @ `a5059cf0` mul1 / K=1..8 + MoE n-tile staging bound. Dest @ `7be99d84` fused `exl3_project_rdna2`. Zoo still one arch per fatbin. Produce stays `-cb 3inst` outside. Do not dump. |
-| `vllm/.../quantization/exl3.py` | — | Dest @ `aaa374ad` / `30b0bd4e`. v0.30 loader (codebook markers, suh grouping, TP slicing, `fold_weight` lm_head only). Stay extras. Python. Do not dump. |
+| `vllm/.../quantization/exl3.py` | — | Dest @ `aaa374ad` / `30b0bd4e` / `b13effcb`. v0.30 loader + dest EXL3 MTP `mtp.fc`. Stay extras. Python. Do not dump. |
+| `scripts/serve_rdna.sh` + `scripts/recipes/*.env` | — | Dest @ `bc5fbee5`. One serve entry + recipes. Stay extras. Serve scripts. Do not copy tok/s. |
 | `vllm/.../layers/layernorm.py` | — | Dest @ `74315f4c`. `gemma_rms_norm_fake` contiguous empty. Stay extras. Inductor fake. Do not dump. |
 | `tunableop/rocblas-f30bb442e9b5/` + `scripts/serve_gfx1030_exl3_27b.sh` | — | Dest @ `ac9dfd8f`. Canonical TunableOp rows + EXL3 27B launcher. Stay extras. Serve/ops. Do not copy tok/s. |
 | `vllm/.../quantization/exl3_moe.py` | — | Dest @ `a50520a3`. `Exl3MoEMethod` for RoutedExperts. Stay extras. Python. Do not dump. |
 | `vllm/.../quantization/exl3_ple.py` | — | Dest @ `c7bc9381`. `Exl3NgramTable` gather. Inert until PLE offload rewrite. Stay extras. Python. Do not dump. |
-| `csrc/rocm/causal_conv1d_rdna2.cu` | `sequence/causal_conv` | Scalar FMA, wave32. FIR on pre-shift then shift. |
+| `csrc/rocm/causal_conv1d_rdna2.cu` | `sequence/causal_conv` | Scalar FMA, wave32. FIR on pre-shift then shift. Dest @ `b13effcb` gates HIP when `state_len != width-1`. Stay extras. Do not dump. |
 | `csrc/rocm/indexer_paged_mqa_rdna2.cu` | `attention/qsa_indexer` | DeepSeek-class indexer. |
 | `csrc/rocm/qsa_rdna2.cu` | `attention/qsa_indexer` (watch) | `VLLM_RDNA_QSA_HIP` default **off**. |
 | `csrc/rocm/sparse_mla_rdna2.cu` | `attention/dsa_nope` | Sparse MLA class. |
@@ -145,7 +147,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 @ `ac9dfd8fba6d` has opt-in
+   `plan`/`run`. `run` is `NOT_READY`. extras @ `bc5fbee5d0a1` has opt-in
    `VLLM_HIPPIHX` ctypes consume (dest **#27**). Still no
    `torch.ops.hippihx.*`. Plans not ready.
 4. extras is rewired **in extras**. The extras copy is then deleted.
@@ -155,10 +157,10 @@ Until then: observe, lock numbers, keep stubs. Tracker:
 
 ## Dest extras defects
 
-Dest tip `ac9dfd8fba6d`. Live dest bugs / rolled-back paths — hippihx
+Dest tip `bc5fbee5d0a1`. Live dest bugs / rolled-back paths — hippihx
 must not reproduce them.
 
-| Defect | Where | Status @ `ac9dfd8f` | Zoo lock |
+| Defect | Where | Status @ `bc5fbee5` | 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`. |
@@ -223,5 +225,5 @@ This review is documentation, not a kernel migrate.
 | Explore PRs **#9/#10** sdot | **Not taken** | **#9** closed-unmerged (superseded by dest W4A8 @ `3a0786ea`). **#10** still not dest. |
 | dest W4A8 sdot4 (`w4a8_sdot4_rdna2.*`, `moe_w4a8_rdna2.cu`) | dest @ `3a0786ea` BlivionIaG | Stay extras. Opt-in. Not a second W4 zoo family. Do not dump. |
 | dest MoE epilogue (`moe_accum_rdna2.cuh`) | dest @ `9cad8792` BlivionIaG | Stay extras. Packed-fp16 CAS default. `VLLM_RDNA2_MOE_FP32_ACCUM` opt-in. Do not dump. |
-| extras dest-presence serve (dest-landed extras PRs, dest **#29** split decode, dest row-gate `e0112c55`, dest W4A8 `3a0786ea`, dest **#30** `3fb9d43a`, dest MoE epilogue `9cad8792`, dest **#32** `a5059cf0`, dest EXL3 loader `30b0bd4e`, dest TunableOp `ac9dfd8f`, mamba `741e5bc3`) | observe | Python/Triton/serve + dest-landed ATen HIP. No HIP migrate. Do not pick Cursor/George/Codex/Karl/Waldecir/a17t/Claude rewrites. Unique Aron Hsiao `rdna_ar` still pickable. |
+| extras dest-presence serve (dest-landed extras PRs, dest **#29** split decode, dest row-gate `e0112c55`, dest W4A8 `3a0786ea`, dest **#30** `3fb9d43a`, dest MoE epilogue `9cad8792`, dest **#32** `a5059cf0`, dest EXL3 loader `30b0bd4e`, dest TunableOp `ac9dfd8f`, dest EXL3 MTP + `serve_rdna` `bc5fbee5`, mamba `741e5bc3`) | observe | Python/Triton/serve + dest-landed ATen HIP. No HIP migrate. Do not pick Cursor/George/Codex/Karl/Waldecir/a17t/Claude rewrites. Unique Aron Hsiao `rdna_ar` still pickable. |
 | extras not dest (**#16** closed / **#18** / **#23** / **#25** / **#31** / **#33** / **#34**) | **Not taken** | GPTQ `BLOCK_KN_SIZE` 256. PCIe P2P KV. rdna_ar retained-output. W4 exact-dequant explore. FA decode scores. FA D=128 prefill. |
diff --git a/docs/CONSUME.md b/docs/CONSUME.md
index bac6e9a..558c18e 100644
--- a/docs/CONSUME.md
+++ b/docs/CONSUME.md
@@ -79,7 +79,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
-`ac9dfd8fba6d` still has no `torch.ops.hippihx` bind. Dest extras **#27**
+`bc5fbee5d0a1` still has no `torch.ops.hippihx` bind. Dest extras **#27**
 (`a3f7e5da`) is opt-in `VLLM_HIPPIHX` ctypes consume via
 `hippihx.v1_ctypes` + `hippihx_v1_load` (default **off**; every plan
 still not ready). FlyDSL extras consume waits on `FLYDSL_V1_CONSUME`
diff --git a/docs/EXTRAS.md b/docs/EXTRAS.md
index 9cede89..b0e0128 100644
--- a/docs/EXTRAS.md
+++ b/docs/EXTRAS.md
@@ -7,7 +7,7 @@ one V1 op. Review: [`BACKPORT.md`](BACKPORT.md).
 ## Unvalidated extras inventory
 
 Snapshot of [`opengfx1030/vllm-rdna`](https://github.com/opengfx1030/vllm-rdna)
-`rdna_extras` @ `ac9dfd8fba6d` (2026-10-01 11:35 UTC). Dest default
+`rdna_extras` @ `bc5fbee5d0a1` (2026-10-01 13:08 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**. Draft **#10**, **#23**,
@@ -17,9 +17,9 @@ squash is `a4060647cfbb`. Open **#2**, **#18**. Draft **#10**, **#23**,
 dest-integrated unmerged; **#12** superseded by **#15**. Merged **#13**,
 **#14**, **#15**, **#17**, **#19**, **#20**, **#22**, **#24**, **#26**,
 **#27**, **#28**, **#29**. Dest **#30** @ `3fb9d43a` / dest **#32** @
-`a5059cf0` dest-presence. Dest tip `ac9dfd8f` is TunableOp/serve (not a
-new extras PR). **No** `torch.ops.hippihx.*`. Dest **#27** is opt-in
-`VLLM_HIPPIHX` ctypes consume (default **off**; plans not ready).
+`a5059cf0` dest-presence. Dest tip `bc5fbee5` is serve recipes + EXL3
+MTP (not a new extras PR). **No** `torch.ops.hippihx.*`. Dest **#27** is
+opt-in `VLLM_HIPPIHX` ctypes consume (default **off**; plans not ready).
 
 **Unvalidated.** Not dest. Not silicon-signed. No tok/s. hippihx still
 ships stubs.
@@ -49,7 +49,7 @@ Status: **extras** = live on dest tip · **Later** = side branch ·
 | EXL3 dense / MoE / dequant / Hadamard / trellis decode | `exl3_dot2_*.cu` | `gemm/exl3_3inst` | Consume `-cb 3inst`. Produce outside. Dest **#32** @ `a5059cf0` mul1 / K=1..8 + MoE n-tile bound. Dest @ `7be99d84` fused `exl3_project_rdna2`. UNC-26. **unvalidated** |
 | GDN packed decode | `gdn_decode_rdna2.cu` | `attention/gdn_scan` | Dest @ `02adbfd4`: SSM **fp16 or fp32**. Dest @ `388a61b6`: no one-shot `zero_()` wipe. **fp16 act only.** **unvalidated** |
 | GDN prefill chain | `gdn_prefill_*_rdna2.cu` | `attention/gdn_scan` | **Opt-in** (`VLLM_GDN_HIP_PREFILL=1`; default Triton/FLA @ `cd1231fd`). `o` varlen `i_t_local` dest-fixed. Dispatch still misses dtype. **unvalidated** |
-| causal_conv1d update + fwd | `causal_conv1d_rdna2.cu` | `sequence/causal_conv` | Scalar FMA. FIR on pre-shift then shift. BF16: fp32 mul, **not** `fdot2.bf16`. **unvalidated** |
+| causal_conv1d update + fwd | `causal_conv1d_rdna2.cu` | `sequence/causal_conv` | Scalar FMA. FIR on pre-shift then shift. Dest @ `b13effcb` gates HIP when `state_len != width-1`. BF16: fp32 mul, **not** `fdot2.bf16`. **unvalidated** |
 | Paged MQA indexer | `indexer_paged_mqa_rdna2.cu` | `attention/qsa_indexer` | gfx1030 BF16 6h×256: **4 warps**. **unvalidated** |
 | Flash-Next QSA store/compress/MQA | `qsa_rdna2.cu` | `attention/qsa_indexer` (watch) | `VLLM_RDNA_QSA_HIP` default **off**. Eager Triton serves (`8cf0dedb`). Dest @ `e45dd5cb`: empty QSA ring prefix hits. Dest @ `f3dd65fa`: live-context prefill scoring bound (Triton). **unvalidated** |
 | Sparse MLA decode / prefill | `sparse_mla_rdna2.cu` | `attention/dsa_nope` | **unvalidated** |
@@ -70,7 +70,8 @@ Status: **extras** = live on dest tip · **Later** = side branch ·
 | W4 exact-dequant explore | draft PR **#31** | — | **skip** — explore, not dest. Claude. Opt-in `VLLM_RDNA2_W4A16_EXACT_DEQUANT`. Zoo already locks integer `q-zero` then scale. Do not dump |
 | EXL3 mul1 decode / K=1..8 trellis | dest PR **#32** @ `a5059cf0` | `gemm/exl3_3inst` | Stay extras. Unsigned mul1 + K=1..8 + dest-fixed MoE staging bound. GitHub PR closed-unmerged. Do not dump. |
 | EXL3 v0.30 loader / fused project / MoE / PLE | dest @ `30b0bd4e` (`exl3.py`, `exl3_moe.py`, `exl3_ple.py`, `exl3_project_rdna2`) | `gemm/exl3_3inst` | Stay extras. Codebook markers, suh grouping, TP slicing, `fold_weight` lm_head only, `Exl3MoEMethod`, inert `Exl3NgramTable`. Unique Blivion. Do not dump. |
-| TunableOp unify / EXL3 27B rows / gemma_rms fake | dest tip `ac9dfd8f` (`tunableop/rocblas-f30bb442e9b5`, `serve_gfx1030_exl3_27b.sh`, `gemma_rms_norm_fake`) | — | Stay extras. Serve/ops + inductor fake. Unique Blivion. Do not dump. Do not copy tok/s. |
+| TunableOp unify / EXL3 27B rows / gemma_rms fake | dest @ `ac9dfd8f` (`tunableop/rocblas-f30bb442e9b5`, `gemma_rms_norm_fake`) | — | Stay extras. Serve/ops + inductor fake. Unique Blivion. Do not dump. Do not copy tok/s. |
+| EXL3 MTP draft + `serve_rdna` recipes | dest tip `bc5fbee5` (`exl3.py` `mtp.fc`, `serve_rdna.sh`, HIP conv `state_len` gate) | — | Stay extras. Dest serves EXL3 quantized MTP. No zoo MTP tile. Unique Blivion. Do not dump. Do not copy tok/s. |
 | FA decode scores / 4 barriers | draft PR **#33** | — | **skip** — not dest. Do not dump `fa_rdna2` |
 | FA D=128 prefill register-O GQA | draft PR **#34** | — | **skip** — not dest. Do not dump `fa_rdna2` |
 | Explore W4A8 sdot4 / W4A4 sdot8 | extras PRs **#9/#10** | — | **#9** closed-unmerged (superseded by dest W4A8 @ `3a0786ea`). **#10** **skip** — not dest |
@@ -169,7 +170,8 @@ Stay extras: serve, product HIP, dest-landed extras PRs, dest **#29**
 split decode / dest row-gate `e0112c55` / dest W4A8 `3a0786ea` / dest
 **#30** compile-dispatch `3fb9d43a` / dest MoE epilogue `9cad8792` /
 dest **#32** EXL3 mul1 `a5059cf0` / dest EXL3 loader `30b0bd4e` / dest
-TunableOp `ac9dfd8f` (observe, do not pick). Skip: a17t
+TunableOp `ac9dfd8f` / dest EXL3 MTP + `serve_rdna` `bc5fbee5`
+(observe, do not pick). Skip: a17t
 PR **#3**, PR **#18** GPTQ `BLOCK_KN_SIZE` 256, draft PR **#23** PCIe
 P2P KV, draft PR **#25** rdna_ar retained-output, draft PR **#31** W4
 exact-dequant explore, draft PR **#33** FA decode scores, draft PR
diff --git a/tests/test_no_b12x_cuda.py b/tests/test_no_b12x_cuda.py
index d99ba66..856860c 100644
--- a/tests/test_no_b12x_cuda.py
+++ b/tests/test_no_b12x_cuda.py
@@ -121,8 +121,10 @@ def test_readme_unvalidated_inventory() -> None:
     assert "a5059cf0" in extras
     assert "30b0bd4e" in extras
     assert "ac9dfd8f" in extras
+    assert "bc5fbee5" in extras
     assert "exl3_project_rdna2" in extras
     assert "gemma_rms_norm_fake" in extras
+    assert "serve_rdna" in extras
     assert "a3f7e5da" in extras
     assert "VLLM_HIPPIHX" in extras
     assert "VLLM_RDNA2_W4A8_SDOT4" in extras
@@ -178,8 +180,10 @@ def test_readme_unvalidated_inventory() -> None:
     assert "a5059cf0" in bp
     assert "30b0bd4e" in bp
     assert "ac9dfd8f" in bp
+    assert "bc5fbee5" in bp
     assert "exl3_project_rdna2" in bp
     assert "gemma_rms_norm_fake" in bp
+    assert "serve_rdna" in bp
     assert "d94e2209" in bp
     assert "dbb1e77" in bp
     assert "e45dd5c" in bp
@@ -232,10 +236,11 @@ def test_readme_unvalidated_inventory() -> None:
     assert "VLLM_RDNA_MOE_RESIDENT" in extras
     assert "moe_resident" in extras
     contrib = (ROOT / "CONTRIBUTING.md").read_text(encoding="utf-8")
-    assert "ac9dfd8f" in contrib
+    assert "bc5fbee5" in contrib
     assert "a4060647" in contrib
     # CONTRIBUTING is tip-only. Historical dest SHAs live in BACKPORT.
     for token in (
+        "ac9dfd8f",
         "30b0bd4e",
         "a5059cf0",
         "9cad8792",
diff --git a/tiles/attention/fa_fdot2/README.md b/tiles/attention/fa_fdot2/README.md
index 5ada77d..f921adb 100644
--- a/tiles/attention/fa_fdot2/README.md
+++ b/tiles/attention/fa_fdot2/README.md
@@ -39,7 +39,7 @@ Dest extras @ `48c56ef` pins `AttentionConfig.backend = RDNA_ATTN` from
 the API-server `VLLM_USE_RDNA2_FA` env so workers actually select
 FA-RDNA2. Dest extras **#28** @ `bfd5286d` and dest **#29** @ `d94e2209` /
 `e0112c55` (`cu_query_lens` split decode + row gate) stay extras.
-Dest tip `ac9dfd8f` does not change FA HIP. Do not dump. Do not pick
+Dest tip `bc5fbee5` does not change FA HIP. Do not dump. Do not pick
 Claude. No tok/s.
 
 ## V1 rev 4 contract (pinned)
diff --git a/tiles/gemm/exl3_3inst/README.md b/tiles/gemm/exl3_3inst/README.md
index 70e53a0..0ee5715 100644
--- a/tiles/gemm/exl3_3inst/README.md
+++ b/tiles/gemm/exl3_3inst/README.md
@@ -25,9 +25,9 @@ not. Do not dump the `.cu`.
 extras **#32** @ `a5059cf0` dest-lands unsigned mul1 + K=1..8 dispatch
 and dest-fixes MoE n-tile staging (GitHub closed-unmerged). Dest @
 `30b0bd4e` dest-lands the v0.30 loader, fused `exl3_project_rdna2`,
-`Exl3MoEMethod`, and inert `Exl3NgramTable` PLE gather. Dest tip
-`ac9dfd8f` is TunableOp/serve, not EXL3 HIP. Stay extras. Do not dump.
-UNC-26 is
+`Exl3MoEMethod`, and inert `Exl3NgramTable` PLE gather. Dest @
+`b13effcb` serves the EXL3 quantized MTP draft. Dest tip `bc5fbee5`
+is serve recipes. Stay extras. Do not dump. UNC-26 is
 still In Progress on extras — do not copy bodies until produce/consume
 and capture buffers stop moving. Port and later perf commits stay
 **BlivionIaG** `<kev29lt@gmail.com>` as Author **and** Committer. Do not
diff --git a/tiles/sequence/causal_conv/README.md b/tiles/sequence/causal_conv/README.md
index 87ac376..ec261bd 100644
--- a/tiles/sequence/causal_conv/README.md
+++ b/tiles/sequence/causal_conv/README.md
@@ -14,7 +14,7 @@ do not retarget one onto the other. This stub is the conv contract only.
 | Math | scalar FMA, fp16 or bf16 in/out, **fp32 acc / mul**. **Not** `fdot2`. **Not** `fdot2.bf16`. gfx1030 LLVM aborts on `llvm.amdgcn.fdot2.bf16.bf16` — BF16 must promote the multiply, not retarget DOT. |
 | Decode FIR order | extras @ `cafe95ef8` / tip `1ff735`: FIR on **pre-shift** state, then shift (matches fwd). Do not shift-first. |
 | Fwd out-stride / null-block | extras @ `82b6f183da2a`: write via `stride_o_token`; invalid slot (`null_block_id`) early-returns. |
-| State | `state_len = width-1`, typically **3 or 4** for GDN |
+| State | `state_len = width-1`, typically **3 or 4** for GDN. Dest @ `b13effcb` gates HIP when spec-decode sizes `state_len != width-1` (Triton fallback). Stay extras. |
 | Graph | scratch sized by `plan`, **zeroed** for page-commit; **no D2H under capture** |
 
 extras decode kernel is cudagraph-safe *as ISA* for **fp16**. Dest GDN