From 93dfb7aa29bf6f33d0eb461398a8202bea4fd921 Mon Sep 17 00:00:00 2001 From: Nirupam Bhowmick <48842933+heydryft@users.noreply.github.com> Date: Wed, 19 Aug 2026 15:34:39 +0100 Subject: [PATCH] =?UTF-8?q?docs(ArcQuant):=20the=20cheap=20decode=20is=20a?= =?UTF-8?q?t=202.25=20bpw,=20not=202.00=20=E2=80=94=20the=20bit-rate=20fro?= =?UTF-8?q?ntier?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The compiled ladder (FACTS §2026-08-19 §1) put K8/V4/L12 at 4.375 inst/wt, 3.46x under the shipped K4/V2/L16, and recorded it with no quality caveat. That geometry is quality-CLOSED: -0.00698 w_cos against a +-0.0008 ship band, 8.7x outside, and nine codebook designs — including D4 lattice cosets and a converged trellis-Lloyd — top out at -0.00307, still 3.8x outside. The same decode shape IS reachable, for a quarter of a bit more. At fixed L=12/V=4 the codebook does not read K (asserted, not assumed), so the table stays 32,768 B and the decode stays one LDS.64 per 4 weights while only the bit rate moves. K9/V4/L12 at 2.25 bpw lands +0.00402 — 5.0x the band's own width on the good side of the shipped control. The price is capacity: KV headroom 58.8 -> 49.6 GB, which is UNMEASURED as a batch/context effect. Records the mechanism, which is why the 2 bpw search stopped: with the trellis OFF, LBG halves the loss vs random (-0.01443 vs -0.02905); with the trellis ON, the same LBG loses to random (-0.00877 vs -0.00714). The binding constraint is trellis freedom, not codebook coverage. Also records, as exact integers: the byte-extraction bound is ceil((8 - gcd(K,8) + K)/8), not ceil((7+K)/8), so K=9 and K=10 are exactly two bytes and the 3-byte route is retired; the warm-up tax is set by slice length, so warp-per-row takes 1.25x/2.98x to 1.06x/1.48x; and the SASS census is structurally blind to that win, because its inst/weight is a differential over unroll depth that cancels prologue and epilogue exactly. CONFIRMED FALSE: qtip/mod.rs's claim that the GPU decode paths compute the codebook holds only for Mcg. DEFAULT is Gaussian, and qtip_gemv.cu gathers from the stored 512 KiB table on that path — which does not fit 48 KB shared, so every symbol lookup is a dependent scattered global load. That is the disease the 32 KB table exists to kill, independent of any instruction count. All CPU, zero GPU hours, $0; every table re-run from scratch, ~2.5 min total. Evidence: arcquant/trellis-geometry-sweep @ 38a03cc14 and arcquant/designed-codebook @ 269e95988, both now pushed. Prunes five stale claims across FACTS.md, 00_RESUME_HERE.md, COMPETITIVE_TEARDOWN.md and KERNEL_RULES.md so no superseded statement sits next to its correction. Co-Authored-By: Claude Opus 5 (1M context) Claude-Session: https://claude.ai/code/session_01SpVNMpb13HkUXqSqbN1o9H --- memory/mission/00_RESUME_HERE.md | 32 ++ memory/mission/COMPETITIVE_TEARDOWN.md | 10 + memory/mission/FACTS.md | 56 ++++ memory/mission/FRONTIER_BITS_FOR_DECODE.md | 365 +++++++++++++++++++++ memory/mission/KERNEL_RULES.md | 8 + 5 files changed, 471 insertions(+) create mode 100644 memory/mission/FRONTIER_BITS_FOR_DECODE.md 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 -- --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