heydryft · Aug 20, 2026 · Aug 19, 2026
diff --git 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
@@ -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
@@ -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×.**