diff --git a/memory/mission/00_RESUME_HERE.md b/memory/mission/00_RESUME_HERE.md
index 0d44e0beb..7bba714ff 100644
--- a/memory/mission/00_RESUME_HERE.md
+++ b/memory/mission/00_RESUME_HERE.md
@@ -50,6 +50,15 @@ reordered by it, with Rung 0 (~3.75 sessions, **zero GPU hours**) gating
 everything above. **Both are on branch `docs/census-session8` until it merges.**
 Measured numbers → `FACTS.md` §2026-08-19. New rules → `KERNEL_RULES.md` top block.
 
+🎯 **`memory/mission/FRONTIER_BITS_FOR_DECODE.md`** — **THE BIT-RATE FRONTIER.
+Read it before quoting the 3.46× decode win, before proposing any codebook work
+at 2 bpw, and before choosing a geometry.** The cheap decode is real but is
+**NOT at 2 bpw**; the shipping geometry is **K9/V4/L12 at 2.25 bpw**. Carries the
+measured reason 2 bpw is closed (trellis freedom, not codebook coverage), the
+exact byte-extraction bound, the census's structural blind spot, and a
+`CONFIRMED FALSE` on the "GPU computes the codebook" comment. **CPU only, $0,
+every table re-run from scratch.**
+
 ## MEASURED this session — zero GPU hours, four results
 
 - **Trellis inst/weight, COMPILED** (`nvcc -cubin`, CUDA 12.4.131, inner loop by
@@ -61,6 +70,18 @@ Measured numbers → `FACTS.md` §2026-08-19. New rules → `KERNEL_RULES.md` to
   exactly **32,768 B**, under the 48 KB static limit, **no `cudaFuncSetAttribute`
   needed**; occupancy **62.5%** (not the predicted 31%) and
   **register-limited, not shared-limited**.
+  🔴 **AND THEN A FIFTH RESULT SUPERSEDED THE GEOMETRY, SAME DAY, SAME $0:**
+  **K8/V4/L12 is quality-CLOSED** — **−0.00698** `w_cos` against a **±0.0008**
+  band, and **nine codebook designs top out at −0.00307**, still 3.8× outside.
+  The 3.46×-shaped decode is available **for a quarter of a bit more**:
+  **K9/V4/L12, 2.25 bpw, +0.00402 — 5.0× the band on the GOOD side of the shipped
+  control** — at the **same 32,768 B table** and the **same decode shape**, since
+  at fixed L=12/V=4 **the codebook does not read K**. **The price is capacity: KV
+  58.8 → 49.6 GB, and that effect is UNMEASURED.** ⚠️ **Do not quote 4.375 for
+  K9 — the kernel has never been compiled at any K but 8.**
+  ✅ **The one-card wedge SURVIVES the extra quarter-bit:** 2.25 bpw ⇒ **83.4 GB
+  of 141**, still **1.8× under MXFP4's 151 GB**, which does not fit at all.
+  Full record + anchors: **`memory/mission/FRONTIER_BITS_FOR_DECODE.md`**.
 - **V4 KV over-retention:** 361.4 MB/seq → **6.35–11.96 MB/seq at 8k = 30–57×**;
   **B=32 @ 8k: 11.57 GB → 0.38 GB.** Compaction copies 1.12 rows per row stored.
 - **Launch reductions shipped:** 9,131 → **8,650/token** (−430 dead mHC casts,
@@ -581,6 +602,17 @@ K4/V2/L16 shipped **15.125** inst/wt (sm_90) → **K8/V4/L12 + row-scale hoist 4
 **3.46× fewer instructions per weight**, LUT exactly 32,768 B (fits static smem, no
 `cudaFuncSetAttribute`), occupancy 62.5% and **register-limited**. **Still 3.1–3.9× short of
 the 1.13–1.41 budget** — that is scope, not a sentence (D21). Full table: `FACTS.md` §2026-08-19.
+🔴 **THAT GEOMETRY IS QUALITY-CLOSED AT 2 bpw AND MUST NOT BE BAKED — the winning
+rung is K9/V4/L12 at 2.25 bpw.** K8/V4/L12 costs **−0.00698** `w_cos` (band
+±0.0008) and **nine codebook designs stop at −0.00307**; the measured reason is
+**trellis freedom, not codebook coverage** (trellis OFF, LBG halves the loss vs
+random; trellis ON, the same LBG loses to random). K9 buys **+0.00402** at the
+**same 32,768 B table and same decode shape** for **0.25 bpw**, costing **KV 58.8
+→ 49.6 GB (UNMEASURED as a batch/context effect)**. ⚠️ **K9's inst/weight is
+UNMEASURED — the kernel has never been compiled at any K but 8. Do not quote
+4.375 for K9.** Runner-up if capacity binds: **K4/V2/L13, 2.00 bpw, −0.00206,
+11.250 inst/wt (1.34×)** — misses the band, costs no bit and no KV.
+Full record: **`memory/mission/FRONTIER_BITS_FOR_DECODE.md`**.
 🔴 **PRICE IT HONESTLY: the re-bake gets ~8× MORE expensive, not cheaper.** The production
 baker is **beam, not exhaustive** (213 vs 8,257 s/layer on A100) and is **issue-bound**, so cost
 `(n/V)×W×2^K` ⇒ **~213 → ~1,700 s/layer ≈ 20 h ≈ $30**. Affordable — but it **inverts** the
diff --git a/memory/mission/COMPETITIVE_TEARDOWN.md b/memory/mission/COMPETITIVE_TEARDOWN.md
index d36bb13e7..52bac7b35 100644
--- a/memory/mission/COMPETITIVE_TEARDOWN.md
+++ b/memory/mission/COMPETITIVE_TEARDOWN.md
@@ -145,6 +145,9 @@ matmul.**
 > `sum2`'s 1,477 is also not the ceiling of the format work: the compiled ladder puts K8/V4/L12 at
 > **4.375 inst/wt**, and the honest read of the whole program is **~15–19× over today's 659 with
 > re-bake plus kernel work** — scope, not a promise, and priced at ~$30 of re-bake.
+> ⚠️ **That 4.375 is K=8, which is quality-closed at 2 bpw; the shippable rung is K=9 at 2.25 bpw
+> and its instruction count is UNMEASURED. Treat the 15–19× as a shape, not a number, until K9 is
+> compiled** (see below and `memory/mission/FRONTIER_BITS_FOR_DECODE.md`).
 **But the trellis tax is a LOW-CONCURRENCY tax that AMORTISES with B:** at B=256, `sum2` lands within
 **1.07×** of llm-d's 1,573-per-decode-GPU — **which takes 32 H200 + InfiniBand. Arc's claim is one
 card.** Path: measured 19.6 @B=32 → ×3.25 ragged (**MEASURED 18.83→61.22**) → ×1.82 MTP per-seq KV →
@@ -169,6 +172,13 @@ all three geometries **2 bpw**, so the scoring error cannot recur):
 
 **3.46× fewer instructions than shipped — and still 3.1–3.9× short of the budget.** Both LUTs fit
 the 48 KB static limit (no `cudaFuncSetAttribute`); occupancy **62.5%**, **register**-limited.
+🔴 **BUT NOT AT 2 bpw — the K8/V4/L12 rows are instruction counts for a geometry that is
+quality-CLOSED** (**−0.00698** `w_cos` vs a ±0.0008 band; nine codebook designs stop at −0.00307).
+**The 3.46×-shaped decode costs a quarter of a bit more: K9/V4/L12 at 2.25 bpw, +0.00402, same
+32,768 B table, same decode shape.** That moves the model to **83.4 GB of 141** — **the one-card
+wedge survives** (MXFP4 is 151 GB and does not fit at all) — but cuts KV headroom **58.8 → 49.6 GB**,
+**UNMEASURED as a batch/context effect**. ⚠️ **K9's inst/weight is UNMEASURED; do not quote 4.375
+for K9.** Full record: **`memory/mission/FRONTIER_BITS_FOR_DECODE.md`**.
 🔴 **And the re-bake gets ~8× MORE expensive, not cheaper** — the production baker is **beam**
 (213 s/layer A100), not exhaustive (8,257), and is issue-bound; `(n/V)×W×2^K` ⇒ **~1,700 s/layer
 ≈ 20 h ≈ $30**. **`L=16` remains the disease** (512 KB never fits shared, which is why RUN-161 went
diff --git a/memory/mission/FACTS.md b/memory/mission/FACTS.md
index 97643f6dc..202fe82e4 100644
--- a/memory/mission/FACTS.md
+++ b/memory/mission/FACTS.md
@@ -65,6 +65,11 @@ unroller ceiling**, so linearity is unverified **for those two rows only**. The
 headline does not depend on the pending re-run.
 
 - **HEADLINE: 15.125 → 4.375 = 3.46× fewer instructions per weight, COMPILED.**
+  🔴 **BUT NOT AT 2 bpw — SUPERSEDED SAME DAY, see §1b.** The geometry those
+  instruction counts were taken at, **K8/V4/L12, is quality-CLOSED**: −0.00698
+  `w_cos` against a ±0.0008 band, and **nine codebook designs cannot rescue it**.
+  The cheap decode is real but costs **0.25 bpw more** (K9/V4/L12). **Do not quote
+  4.375 for K9 — the kernel has never been compiled at any K but 8.**
 - **Still 3.1–3.9× short of the 1.13–1.41 budget on sm_90** (2.6–3.2× on sm_80),
   **and that is the inner loop alone** — it excludes serving scaffolding (expert
   gather, shared staging, tail guards). ⚠️ **Geometry-to-geometry ratios are
@@ -80,6 +85,57 @@ headline does not depend on the pending re-run.
   K=8 warm-up is only `L/K = 2` byte-aligned symbols, cheaper than reconstructing
   the window. **Its value is confined to K=4.**
 
+### 1b. THE QUALITY SIDE OF THAT SAME GEOMETRY — **the 3.46× is at 2.25 bpw, not 2.00**
+
+**Full record: `memory/mission/FRONTIER_BITS_FOR_DECODE.md`.** CPU only, zero GPU
+hours, every table re-run from scratch the day it was written. Control for all of
+it: **K4/V2/L16, `w_cos`, gaussian σ=0.02, exhaustive Viterbi, rot=128, max/3,
+n=48, k=2048, 5 draws.** Ship band **±0.0008**.
+
+| geometry | bpw | **Δ w_cos** | table B | model GB | **KV GB** | band |
+|---|---|---|---|---|---|---|
+| K8/V4/L12 | 2.00 | **−0.00698** | 32,768 | 74.2 | 58.8 | **fails, 8.7× out** |
+| **K9/V4/L12** | **2.25** | **+0.00402** | 32,768 | **83.4** | **49.6** | **CLEARS, 5.0× the band on the good side** |
+| K10/V4/L12 | 2.50 | +0.01166 | 32,768 | 92.7 | 40.3 | CLEARS |
+| K12/V4/L12 (L==K ⇒ memoryless) | 3.00 | +0.02057 | 32,768 | 111.2 | 21.8 | CLEARS |
+
+At fixed **L=12/V=4 the codebook does not read K** (asserted in the probe, not
+assumed), so the table stays **32,768 B** and the decode shape stays one
+`LDS.64` per 4 weights across every row. **Only the bit rate moves.**
+⚠️ **The price is CAPACITY — KV 58.8 → 49.6 GB, −15.6%, and its batch/context
+effect is UNMEASURED.** Anchor: `probe_bit_rate_frontier`, 6.63 s, branch
+`arcquant/designed-codebook` @ `269e95988`.
+
+- 🔴 **2 bpw is CLOSED, and the mechanism is measured.** Nine codebook designs at
+  K8/V4/L12 top out at **−0.00307** (converged trellis-Lloyd, 40 sweeps): random
+  −0.00714 · amplitude re-fit −0.00597 · **LBG + per-block Haar −0.00457** · LBG
+  set-partitioned −0.00877 · LBG clustered −0.00895 · **D4 lattice cosets
+  −0.01008** · LBG memoryless −0.01443. Means over 3 fixtures × 5 draws.
+- 🔑 **THE MECHANISM: trellis OFF ⇒ LBG HALVES the loss vs random (−0.01443 vs
+  −0.02905); trellis ON ⇒ the same LBG LOSES to random (−0.00877 vs −0.00714).**
+  ⇒ **the binding constraint is TRELLIS FREEDOM, not codebook coverage.** More
+  codebook design at 2 bpw is not worth an hour.
+- **Structural result: K8/V4 at depth L behaves exactly like K4/V2 at L−2**
+  (−0.00115/−0.00119 and −0.00325/−0.00320). ⇒ **`L/K` lookback is NOT the
+  predictor.** Quote this before re-deriving it.
+- **Runner-up, still live: K4/V2/L13 — 2.00 bpw, −0.00206, 32,768 B, compiled
+  11.250 inst/wt (1.34×).** Misses the band but costs **no extra bit and no KV**.
+- ⚠️ **UNMEASURED, do not soften:** K9's inst/weight (the kernel has never been
+  compiled at any K but the K=8 control) · the pad-vs-clamp delta · the
+  batch/context effect of 49.6 vs 58.8 GB · bf16 vs f32 codebook values.
+
+### 1c. 🔴 CONFIRMED FALSE — "the GPU decode paths compute the codebook"
+
+`materialize` (`mistralrs-quant/src/qtip/mod.rs:596-598`) claims the GPU decode
+paths compute the table instead of storing it. **True only for `Mcg`.**
+`QtipCodebook::DEFAULT` is `Gaussian` (`mistralrs-quant/src/qtip/mod.rs:570`), and
+on that path the kernel **gathers from the stored 512 KiB table**
+(`mistralrs-quant/kernels/qtip/qtip_gemv.cu:213-215`) — which does not fit 48 KB
+shared at any occupancy, so **every symbol lookup is a dependent, scattered global
+load**, the 388 GB/s ≈ 8%-of-HBM decode limiter already recorded below.
+🔑 **That is the disease the 32,768 B table exists to kill — a reason the L=12
+family matters that is independent of any instruction count.**
+
 ### 2. V4 KV OVER-RETENTION — **30–57×**, source-derived
 
 **361.4 MB/seq → 6.35–11.96 MB/seq at 8k context = 30–57×.**
diff --git a/memory/mission/FRONTIER_BITS_FOR_DECODE.md b/memory/mission/FRONTIER_BITS_FOR_DECODE.md
new file mode 100644
index 000000000..5db06825c
--- /dev/null
+++ b/memory/mission/FRONTIER_BITS_FOR_DECODE.md
@@ -0,0 +1,365 @@
+# 🎯 THE BIT-RATE FRONTIER — the cheap decode is real, at 2.25 bpw not 2.00
+
+**Parent system: ArcQuant / QTIP.** Measured 2026-08-19, **CPU only, zero GPU
+hours, $0**. Every table below was re-run from scratch on the machine that wrote
+this file; nothing here is transcribed from a prior session.
+
+> **THE ONE SENTENCE.** The 3.46× cheaper decode geometry **is** available — but
+> **not at 2 bpw**, where it is quality-closed by −0.00698 w_cos against a
+> ±0.0008 band. It is available for **a quarter of a bit more**: **K=9/V=4/L=12,
+> 2.25 bpw**, which lands **+0.00402 ABOVE the shipped control** at the **same
+> 32,768 B table** and the **same decode shape**. The price is not quality. The
+> price is **capacity**: KV headroom 58.8 → 49.6 GB.
+
+---
+
+## 0. The control, stated once, and true of every number in this file
+
+Unless a row says otherwise, every `d(w_cos)` below is measured against **one**
+control and **one** configuration:
+
+| knob | value |
+|---|---|
+| control geometry | **K4/V2/L16** — the shipped rung, 2.008 bpw |
+| metric | **`w_cos`** — weight-space cosine, the tight one |
+| fixture | **gaussian** (`gen_gaussian`), σ = 0.02 |
+| search | **exhaustive Viterbi** (not greedy, not beam) |
+| rotation | Hadamard block **128** |
+| scale policy | **max\|row\|/3** |
+| shape | **n = 48, k = 2048** |
+| draws | **5 weight draws** (seeds 1, 101, 202, 303, 404) |
+| codebook | pseudorandom Box-Muller Gaussian, `gaussian_lut_geo` |
+
+**Ship band: ±0.0008 `d(w_cos)`.** It is not arbitrary. Both edges are already in
+the record (`FACTS.md`, quantization section):
+
+* **−0.0008** is what **shipped**: the `sum2` codebook variant, *"−0.0008 cos
+  worst family, +0.0002 mean — 5× smaller than the W=128 delta Jish rejected"*.
+  That is the largest loss the program has ever called quality-neutral.
+* **−0.0013 … −0.0021** is the far side: the measured beam-vs-exhaustive loss on
+  `fp4_dequant` (exhaustive wins 8 of 9 fixture cells). It is a **recorded
+  regression**, shipped knowingly and **only because beam is faster**. Nothing
+  has ever bought that magnitude with quality alone.
+
+So a candidate at −0.002 is not "close". It is at the top of a range the program
+already knows it can see.
+
+**Where the numbers live.** `mistralrs-quant/src/qtip/bake_quality_tests.rs` on
+branch `arcquant/designed-codebook` @ `269e95988` (which contains
+`arcquant/trellis-geometry-sweep` @ `38a03cc14` as an ancestor). Re-run any row:
+
+```
+cargo test -p mistralrs-quant --release <probe_name> -- --ignored --nocapture
+```
+
+Total cost to reproduce **every table in this document**: about **2.5 minutes**
+of laptop CPU on a warm build.
+
+---
+
+## 1. THE ANSWER — the bit-rate frontier at fixed L=12/V=4
+
+Anchor: `probe_bit_rate_frontier`, **6.63 s**, output reproduced verbatim.
+
+| geometry | bpw | **Δ w_cos** | [min, max] | table B | ≤32 KB | model GB | **KV GB** | band |
+|---|---|---|---|---|---|---|---|---|
+| **K8**/V4/L12 | 2.00 | **−0.00698** | [−0.00703, −0.00692] | 32,768 | yes | 74.2 | 58.8 | **fails** |
+| **K9**/V4/L12 ⬅ **THE ANSWER** | **2.25** | **+0.00402** | [+0.00391, +0.00420] | 32,768 | yes | **83.4** | **49.6** | **CLEARS** |
+| **K10**/V4/L12 | 2.50 | +0.01166 | [+0.01150, +0.01175] | 32,768 | yes | 92.7 | 40.3 | CLEARS |
+| **K11**/V4/L12 | 2.75 | +0.01699 | [+0.01679, +0.01717] | 32,768 | yes | 101.9 | 31.1 | CLEARS |
+| **K12**/V4/L12 (L==K) | 3.00 | +0.02057 | [+0.02032, +0.02077] | 32,768 | yes | 111.2 | 21.8 | CLEARS |
+
+**K9 is not "inside the band". It is 5.0× the band's own width on the good side
+of the shipped control** — `+0.00402 / 0.0008 = 5.03`. The 2 bpw row is
+**8.7× outside** it in the other direction.
+
+### Why the codebook does not move across that table
+
+At fixed **L=12, V=4** the table is `2^L × V` bf16 = **32,768 B exactly** and
+**does not read K at all**. The probe asserts this rather than assuming it: it
+builds the table at K=8…12 and requires bit-equality, failing with *"codebook
+moved with K; frontier is confounded"* otherwise. So every row above uses the
+**same 4096-entry 4-D table** and the **same one-`LDS.64`-per-4-weights decode**.
+**The only thing varying is how many bits index it** — `bpw = K/V`
+(`mistralrs-quant/src/qtip/mod.rs:376-380`).
+
+That is what makes this a *frontier* and not a sweep: quality is priced **in
+bits, with decode cost held constant**.
+
+### The price is capacity, and it is UNMEASURED as a serving effect
+
+H200 accounting, from the in-tree reserve (`FACTS.md`: *"141 − 74.2 − ~8 reserve
+⇒ ~59 GB usable"*): usable KV = `141 − model − 8`.
+
+**58.8 → 49.6 GB is a −15.6% cut in KV headroom.** Its effect on batch size and
+context length **has not been measured**. Do not trade it away in a sentence —
+see §6.
+
+### The three-bpw diagnostic: buying table size back does NOT recover the loss
+
+Anchor: `probe_bits_for_decode_ladder`, **14.76 s**. Rows over the 32 KB bar are
+diagnostics, not candidates:
+
+| geometry | role | bpw | Δ w_cos | table B | ≤32 KB |
+|---|---|---|---|---|---|
+| K8/V4/L12 | anchor, 2.0 bpw | 2.00 | −0.00698 | 32,768 | yes |
+| K4/V2/L13 | anchor, 2.0 bpw | 2.00 | **−0.00202** | 32,768 | yes |
+| K10/V4/L12 | 2.5 bpw V4 | 2.50 | +0.01166 | 32,768 | yes |
+| K5/V2/L13 | 2.5 bpw V2 | 2.50 | +0.01717 | 32,768 | yes |
+| K12/V4/L12 | 3.0 bpw V4, **memoryless** | 3.00 | +0.02057 | 32,768 | yes |
+| K6/V2/L13 | 3.0 bpw V2 | 3.00 | +0.02681 | 32,768 | yes |
+| K6/V2/L12 | 3.0 bpw V2, half table | 3.00 | +0.02625 | 16,384 | yes |
+| K3/V1/L14 | 3.0 bpw V1 | 3.00 | +0.02752 | 32,768 | yes |
+| K12/V4/L13 | over bar, 64 KB | 3.00 | +0.02287 | 65,536 | **no** |
+| K12/V4/L14 | over bar, 128 KB | 3.00 | +0.02441 | 131,072 | **no** |
+| K6/V2/L14 | over bar, 64 KB | 3.00 | +0.02723 | 65,536 | **no** |
+
+**At equal bit rate — each family at its largest table that still clears the
+32 KB bar — V=1 > V=2 > V=4 in quality, and the ordering never inverts:**
+K3/V1 (+0.02752) > K6/V2 (+0.02681) > K12/V4 (+0.02057) at 3 bpw; K5/V2
+(+0.01717) > K10/V4 (+0.01166) at 2.5 bpw; K4/V2/L13 (−0.00202) > K8/V4/L12
+(−0.00698) at 2 bpw. **V=4 is bought WITH quality, and the bit rate is what buys
+it back.** Note also that **going over the 32 KB bar does not recover the loss**:
+K12/V4 at 64 KB (+0.02287) and 128 KB (+0.02441) both stay below the V=2 row that
+fits in 32 KB (+0.02681).
+
+`K12/V4/L12` has `L == K`, so `state = ((state << K) | sym) & mask` shifts every
+history bit out and the code is **memoryless — a plain 4096-point 4-D VQ with no
+trellis**. That is asserted, not assumed, by `l_equals_k_is_a_memoryless_code`.
+
+---
+
+## 2. WHY 2 bpw IS CLOSED — record the mechanism, it is why we stopped searching
+
+The geometry sweep drew every codebook at *random*, which bounds the loss of a
+**random** codebook at K8/V4/L12, not the geometry's own limit. So the table was
+**designed** instead — nine ways, including the two strongest constructions in
+the literature. Anchor: `probe_designed_codebook_at_k8v4l12`, **49.0 s**; column
+is the **mean over all 3 fixtures × 5 draws** (gaussian, student-t(4),
+fp4-dequant), which is why these differ slightly from §1's gaussian-only rows.
+
+| codebook at K8/V4/L12 | mean Δ w_cos | vs random |
+|---|---|---|
+| gaussian (random) — **the control design** | **−0.00714** | — |
+| gaussian + amplitude re-fit | −0.00597 | better |
+| **LBG + per-block Haar rotation** — best design | **−0.00457** | better |
+| trellis-Lloyd ← random, 8 sweeps | −0.00450 | better |
+| **trellis-Lloyd ← random, CONVERGED (40 sweeps)** | **−0.00307** | **best of all** |
+| trellis-Lloyd ← LBG, 8 sweeps | −0.00727 | **worse** |
+| LBG set-partitioned (Ungerboeck) | −0.00877 | **worse** |
+| LBG clustered (negative control) | −0.00895 | **worse** |
+| **D4 lattice cosets** — best known 4-D quantizer | **−0.01008** | **worse** |
+| LBG memoryless VQ | −0.01443 | **worse** |
+| random memoryless VQ | −0.02905 | **worse** |
+
+**The best design in the world at this geometry is −0.00307. The band is
+±0.0008. That is still 3.8× outside.** K8/V4/L12 is **quality-CLOSED**, and it is
+closed by the *geometry*, not by the table.
+
+The convergence is not truncation: `probe_trellis_lloyd_convergence`, **60.9 s**,
+walks 0 → 40 sweeps at −0.00597 → −0.00365 → −0.00331 → −0.00318 → −0.00312 →
+**−0.00307**, and a further amplitude re-fit gives **−0.00308** (worse). It has
+converged.
+
+### 🔑 THE MECHANISM — measured, not inferred
+
+Two rows in that table are a controlled experiment, and they are the reason the
+search stopped:
+
+> **Trellis OFF** (all blocks identical, so the memory is switched off on
+> purpose): **LBG −0.01443 vs random −0.02905 — LBG halves the loss.** The LBG
+> training is sound and codebook design works.
+>
+> **Trellis ON**: the *same* LBG centroids, laid into the trellis by the textbook
+> set partition, give **−0.00877 — WORSE than random's −0.00714.**
+
+**So the binding constraint at K8/V4/L12 is TRELLIS FREEDOM, not codebook
+coverage.** A better-covering table cannot help, because coverage is not what is
+short. This is why more codebook design at 2 bpw is not worth an hour: the
+instrument says the lever is not attached to the load.
+
+The one design that *does* beat random — LBG + per-block **Haar** rotation
+(−0.00457) — beats it precisely by restoring per-block *reach* rather than
+per-block *quality*, which is the same finding from the other side.
+
+---
+
+## 3. THE GEOMETRY LADDER, and its structural result
+
+Anchor: `probe_geometry_delta_noise`, **42.65 s**, gaussian fixture, 5 draws.
+
+| geometry | trellis depth `L/K` | Δ w_cos |
+|---|---|---|
+| K4/V2/L14 | 3.50 | −0.00115 |
+| K8/V4/L16 | 2.00 | −0.00119 |
+| K4/V2/L12 | 3.00 | −0.00325 |
+| K8/V4/L14 | 1.75 | −0.00320 |
+| K8/V4/L12 | 1.50 | −0.00698 |
+
+Read the pairs, not the rows:
+
+> **K8/V4 at depth L behaves exactly like K4/V2 at depth L−2.**
+> L=14 ↔ L=16: −0.00115 vs −0.00119. L=12 ↔ L=14: −0.00325 vs −0.00320.
+
+**Therefore `L/K` lookback is NOT the predictor**, even though it is the natural
+one — those pairs have depths 3.50 vs 2.00 and 3.00 vs 1.75 and still land on top
+of each other. **`L` alone (offset by V) predicts; the depth ratio does not.**
+That result survives to every future geometry question and should be quoted
+before anyone re-derives it.
+
+The ladder holds on all three fixtures — the same pairing appears in student-t(4)
+(−0.00114/−0.00123, −0.00336/−0.00338) and fp4-dequant (−0.00121/−0.00127,
+−0.00338/−0.00341) — so it is a property of the code, not of the weights.
+
+### The runner-up, and it is a live option
+
+**K4/V2/L13 — 2.00 bpw, Δ w_cos −0.00206, 32,768 B table, compiled at 11.250
+inst/wt on sm_90 (1.34× cheaper than the shipped 15.125).** Anchor for the
+instruction count: `FACTS.md` §2026-08-19 §1, the compiled `nvcc -cubin` table.
+
+It **does not clear the ship band** — −0.00206 lands at the top of the
+−0.0013…−0.0021 range §0 anchors as a recorded regression — but it is the only
+2.00-bpw candidate within reach of it, it costs **no extra bit**, and it keeps KV
+at 58.8 GB. If capacity turns out to bind harder than §1 assumes, this is the
+fallback: **1.34× instead of 3.46×, and no capacity loss at all.**
+
+---
+
+## 4. KERNEL SIDE — exact integers, not estimates
+
+### 4.1 The byte-extraction bound is `ceil((8 − gcd(K,8) + K)/8)`, and it is EXACT
+
+The naive bound is `ceil((7+K)/8)`, which maximises over **all** bit offsets 0..7.
+But symbol `t` starts at bit `K·t`, so the **reachable** offsets of `K·t mod 8`
+are only the multiples of `gcd(K,8)`. The true bound is therefore
+
+```
+bytes(K) = ceil((8 − gcd(K,8) + K) / 8)      K=8 → 1     K=9 → 2     K=10 → 2
+```
+
+**Two bytes is exact for both new rungs, not merely sufficient. This retires the
+3-byte route for every K in play.** Anchor: `census_sym_ext` in
+`arc-tools/sass_census/qtip_sass_census.cu` on branch `perf/k9-alignment-census`
+@ `c3994cc14`, where the naive 3-byte route (`EXT_B3`) is still compiled — solely
+so a defensively-written kernel can see what the unnecessary generality costs.
+
+### 4.2 The warm-up tax is set by SLICE LENGTH, and one warp per row collapses it
+
+Every lane replays `W` warm-up symbols before decoding its own slice of `S`, so
+extraction cost per weight is `(S + W) / (S·V)` against an ideal `1/V` — a tax of
+**`1 + W/S`**, driven entirely by slice length, and slice length is
+`num_symbols / lanes_per_row`. Spreading a row across a whole 128-thread block
+makes slices short and the tax large:
+
+| `in_features` | 128 lanes/row (old) | **32 lanes/row (warp-per-row)** |
+|---|---|---|
+| 4096 | S=8, **1.25×** | S=32, **1.06×** |
+| 1024 | S=2, **1.99×** | S=8, **1.24×** |
+| 512 | S=1, **2.98×** | S=4, **1.48×** |
+
+Same parallelism, four times the slice length. It **also deletes the cross-warp
+butterfly, the `warp_sums` shared array, and BOTH `__syncthreads` from the row
+loop**, since a warp reduces with shuffles alone. Resulting rate: **0.2651
+extractions per weight at `in_features=4096`, 0.3105 at 1024** (theoretical floor
+0.25). The tax is **K-independent** — it multiplies the extraction count at every
+K, so it scales the alignment penalty too.
+
+Anchor: `KERNEL_LANES_PER_ROW` and `the_warmup_tax_is_bounded_by_slice_length` in
+`mistralrs-quant/src/qtip/trellis_v4l12.rs`, branch `feat/qtip-k8v4l12` @
+`b6023cd2e` (PR #168).
+
+### 4.3 ⚠️ THE SASS CENSUS IS STRUCTURALLY BLIND TO §4.2 — do not misread a flat row
+
+The alignment census measures `inst/weight` as a **differential over unroll
+depth**. That cancels prologue and epilogue **exactly** — which is what makes it
+robust, and is **precisely why it cannot see per-row or per-block overhead**.
+
+**A barrier removal, a deleted shared array and a deleted cross-warp reduction
+can never appear in a census number. The instrument can neither credit nor charge
+them.**
+
+> **A flat census result does NOT mean the reshape did nothing.** It means the
+> census measured steady-state inner-loop cost per weight, in which the reshape
+> appears only indirectly, through the extraction rate it changes.
+
+And when the census does report a per-route cost, multiply it by **this rung's
+own extraction rate at the shape in question** — 0.2651 at 4096 — not by an
+illustrative endpoint of 0.25 or 0.375. Anchor:
+`Rung::structural_per_row_overhead` and `route_cost_per_weight_x10000`, same file
+and branch as §4.2, pinned there deliberately so the structural half has an
+accounting that is never presented as an instruction count.
+
+### 4.4 🔴 CONFIRMED FALSE — "the GPU decode paths compute the codebook"
+
+`materialize` (`mistralrs-quant/src/qtip/mod.rs:596-598`) says the table is
+*"Always materialized for the CPU search and for the artifact; **the GPU decode
+paths compute instead**."*
+
+**That is true only for `Mcg`.** `DEFAULT` (`mistralrs-quant/src/qtip/mod.rs:570`)
+is `QtipCodebook::Gaussian`, and on that path `qtip_cb_value_ldg`
+(`mistralrs-quant/kernels/qtip/qtip_gemv.cu:213-215`) **gathers from the stored
+table** — *"LUT path: 512 KiB FP32, L2-resident, `__ldg` read-only."* The kernel
+header states the size outright: *"The LUT (2^16 × V floats = 512 KiB) lives in
+L2 cache"* (`mistralrs-quant/kernels/qtip/qtip_gemv.cu:11-12`).
+
+**512 KiB does not fit 48 KB shared memory at any occupancy. So on the SHIPPED
+DEFAULT path, every symbol lookup is a dependent, data-scattered global load** —
+measured at **388 GB/s ≈ 8% of H200 HBM** (`FACTS.md`), and it is the decode
+limiter.
+
+🔑 **That is the disease the 32,768 B table exists to kill, and it is the real
+reason the L=12 family matters — independent of any instruction count.** Anyone
+reading `mod.rs`'s comment and concluding "the codebook is free on GPU" has
+mis-costed the entire rung.
+
+---
+
+## 5. WHAT THIS CHANGES IN THE EXISTING RECORD
+
+| previously recorded | status now |
+|---|---|
+| "V=4 is compiled: K8/V4/L12 + row-scale hoist = **4.375** inst/wt, **3.46×** fewer than shipped" — stated with **no quality caveat** | **The instruction counts stand. The implied ship candidate does not.** K8/V4/L12 is **quality-CLOSED at −0.00698**, 8.7× outside the band, and nine codebook designs cannot rescue it (§2). The 3.46× decode is reachable at **K9/V4/L12, 2.25 bpw** — same table, same shape, **+0.00402** |
+| "the GPU decode paths compute the codebook" (`mistralrs-quant/src/qtip/mod.rs:596-598`) | **CONFIRMED FALSE for the default path** (§4.4) |
+| a K-field extract needs `ceil((7+K)/8)` = 3 bytes | **superseded** — `ceil((8 − gcd(K,8) + K)/8)`; K=9 and K=10 are **exactly 2** (§4.1) |
+| a flat SASS census row ⇒ "the reshape bought nothing" | **invalid inference** — the instrument is structurally blind to it (§4.3) |
+
+---
+
+## 6. 🚫 STILL UNMEASURED — recorded as such, not softened
+
+1. **K9's `inst/weight`.** The kernel has **never been compiled at any K other
+   than the K=8 control**. nvcc lanes are queued. **Do NOT quote 4.375 for K9** —
+   K=9 is not byte-aligned and pays the §4.1 two-byte extract that K=8 does not.
+2. **The pad-vs-clamp delta.** An unpadded row stride forces `min(idx,
+   row_bytes−1)` on every byte index; padding the stride by `MAX_BYTES−1` buys
+   that back. The delta is a real, removable cost and is **the open format
+   decision on the serving side**. Unpriced.
+3. **The KV-capacity effect of 49.6 vs 58.8 GB.** −15.6% headroom. Its effect on
+   achievable batch and context is **not measured**, and capacity is the wedge —
+   this is the one number that could still close K9.
+4. **bf16 vs f32 codebook values.** The 32,768 B figure assumes bf16. The quality
+   probes ran in f32. The gap is unmeasured.
+5. **Any of this on a card.** Everything above is CPU arithmetic and compiler
+   output. No GPU hour was spent, and none of it is a wall-clock claim.
+
+---
+
+## 7. WHERE THE CODE IS
+
+| what | where |
+|---|---|
+| geometry parameterisation + K/V/L sweep | branch `arcquant/trellis-geometry-sweep` @ `38a03cc14` |
+| designed codebooks + the bit-rate frontier | branch `arcquant/designed-codebook` @ `269e95988` (contains the above) |
+| K8/V4/L12 decode family, K as a parametric seam | **PR #168** — `feat/qtip-k8v4l12` |
+| UQFF geometry discriminator (K is an explicit wire field) | **PR #170** — `feat/qtip-k8v4l12-format`, **stacked on #168** |
+| geometry-parametric beam encoder, ready to bake | **PR #169** — `wave-encoder-k8l12` |
+| byte-alignment / extraction-route SASS census | branch `perf/k9-alignment-census` @ `c3994cc14` |
+
+Both `arcquant/*` branches carry the evidence in
+`mistralrs-quant/src/qtip/bake_quality_tests.rs` as runnable probes, with
+anti-silent-success guards on every table: `designed_codebook_reaches_the_search`
+(no two designs may reconstruct alike), `geo_pipeline_is_byte_identical_at_shipped_geometry`
+(the parameterisation is a refactor at the shipped geometry, byte for byte),
+`generic_bit_packing_reproduces_the_production_layouts`, and a per-row guard that
+fails if a geometry's reconstruction matches the control's. **Five guards fired
+on unfixed code while this was being built** — the tables are not self-reported.
diff --git a/memory/mission/KERNEL_RULES.md b/memory/mission/KERNEL_RULES.md
index a0c7039c4..91723a8a3 100644
--- a/memory/mission/KERNEL_RULES.md
+++ b/memory/mission/KERNEL_RULES.md
@@ -40,6 +40,14 @@ claim that the full ceiling was one re-bake away.
 are not allowed a shared budget line. The compiled ladder is safe *because* all
 three of its geometries are 2 bpw.
 
+⚠️ **AND A LADDER THAT HOLDS bpw FIXED HAS NOT PICKED A WINNER — IT HAS ONLY
+PRICED ONE AXIS.** Holding all three geometries at 2 bpw made the instruction
+comparison valid and simultaneously hid the fact that **the cheapest of the three
+is quality-closed at that bit rate** (K8/V4/L12: −0.00698 `w_cos` vs a ±0.0008
+band). The winning rung turned out to be **off the ladder entirely**, at 2.25 bpw.
+⇒ **Price quality on the same sweep, or state in the same sentence that you have
+not.** `memory/mission/FRONTIER_BITS_FOR_DECODE.md`.
+
 ## 3. A DEPRECATED CLI THAT PRINTS HELP AND EXITS **0** IS SILENT SUCCESS IN A NEW COSTUME
 
 `huggingface-cli` → `hf`. The old binary still exists, still prints, still exits