From 445a3dc4c7d561f63414e2c41b76292db36627ca Mon Sep 17 00:00:00 2001 From: TheTom Date: Wed, 25 Mar 2026 16:16:58 -0500 Subject: [PATCH] docs: move engineering docs here from llama.cpp fork MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit These belong in our repo, not in a fork of someone else's project. Added: - docs/quality-benchmarks.md — perplexity validation, bisection log, top-of-tree table - docs/turbo-speed-investigation.md — Metal gotchas, fp16 WHT, optimization history - docs/pre-rotate-queries-investigation.md — WHT/RoPE commutativity issue - docs/pre-rotate-queries-plan.md — original optimization plan - scripts/turbo-quality-gate.sh — pre-push perplexity check README links updated to point locally. Co-Authored-By: Claude Opus 4.6 (1M context) Co-Authored-By: tturney@psyguard.ai --- README.md | 9 +- docs/pre-rotate-queries-investigation.md | 77 +++++ docs/pre-rotate-queries-plan.md | 132 ++++++++ docs/quality-benchmarks.md | 312 ++++++++++++++++++ docs/turbo-speed-investigation.md | 400 +++++++++++++++++++++++ scripts/turbo-quality-gate.sh | 49 +++ 6 files changed, 975 insertions(+), 4 deletions(-) create mode 100644 docs/pre-rotate-queries-investigation.md create mode 100644 docs/pre-rotate-queries-plan.md create mode 100644 docs/quality-benchmarks.md create mode 100644 docs/turbo-speed-investigation.md create mode 100755 scripts/turbo-quality-gate.sh diff --git a/README.md b/README.md index e785bd1..003e586 100644 --- a/README.md +++ b/README.md @@ -235,10 +235,11 @@ benchmarks/ Detailed debugging logs, gotchas, and benchmarks from the llama.cpp port: -- [Quality Benchmarks](https://github.com/TheTom/llama-cpp-turboquant/blob/feature/turboquant-kv-cache/docs/quality-benchmarks.md) — perplexity validation, bisection log, top-of-tree quality+speed table -- [Speed Investigation](https://github.com/TheTom/llama-cpp-turboquant/blob/feature/turboquant-kv-cache/docs/turbo-speed-investigation.md) — Metal gotchas, fp16 WHT results, optimization history -- [Pre-Rotate-Queries Investigation](https://github.com/TheTom/llama-cpp-turboquant/blob/feature/turboquant-kv-cache/docs/pre-rotate-queries-investigation.md) — why WHT and RoPE don't commute (saves you weeks of debugging) -- [Quality Gate Script](https://github.com/TheTom/llama-cpp-turboquant/blob/feature/turboquant-kv-cache/scripts/turbo-quality-gate.sh) — pre-push perplexity check +- [Quality Benchmarks](docs/quality-benchmarks.md) — perplexity validation, bisection log, top-of-tree quality+speed table +- [Speed Investigation](docs/turbo-speed-investigation.md) — Metal gotchas, fp16 WHT results, optimization history +- [Pre-Rotate-Queries Investigation](docs/pre-rotate-queries-investigation.md) — why WHT and RoPE don't commute (saves you weeks of debugging) +- [Pre-Rotate-Queries Plan](docs/pre-rotate-queries-plan.md) — original optimization plan and approach +- [Quality Gate Script](scripts/turbo-quality-gate.sh) — pre-push perplexity check ## Contributing diff --git a/docs/pre-rotate-queries-investigation.md b/docs/pre-rotate-queries-investigation.md new file mode 100644 index 0000000..63c306f --- /dev/null +++ b/docs/pre-rotate-queries-investigation.md @@ -0,0 +1,77 @@ +# Pre-Rotate-Queries Investigation Log + +## Goal +Move WHT inverse rotation from per-block dequant (O(128) per 4-element access) to graph-level Q forward + V inverse rotation (O(1) per-access amortized). This would reclaim ~77 tok/s speed while maintaining PPL ~6.19. + +## Model Under Test +- Qwen3.5-35B-A3B-Q8_0 (MoE) +- n_embd_head_k = 256, n_embd_head_v = 256, n_head = 16, n_head_kv = 2 +- WHT rotation group = 128 (QK_TURBO3 = 128) +- Each head has 2 rotation groups (256/128 = 2) + +## ggml_mul_mat Semantics (Verified) +- ggml stores 2D tensors column-major: element(i,j) at offset i + j*ne[0] +- Storing a C row-major array M into ggml: ggml sees M^T +- ggml_mul_mat(A, x) computes A^T @ x +- NET EFFECT: storing row-major M and calling mul_mat gives M @ x +- VERIFIED with 2x2 rotation test: stored R row-major, got R@x output + +## Rotation Matrix Storage +- TURBO_ROTATION_R: R = diag(s2) * H/sqrt(128) * diag(s1) (forward rotation) +- TURBO_ROTATION_RT: R^T = R^{-1} (inverse rotation) +- For Q forward: store R (TURBO_ROTATION_R) -> mul_mat gives R @ q +- For V inverse: store R^T (TURBO_ROTATION_RT) -> mul_mat gives R^T @ cur = R^{-1} @ cur +- Python verified: R^T @ R = I, round-trip error = 1.2e-15 + +## Test Results (all on Qwen3.5-35B-A3B, wikitext-2, 8 chunks, turbo3 K+V, flash attn) + +### Baseline +| Config | PPL | Notes | +|--------|-----|-------| +| Dequant inverse ON, no graph rot | **6.194** | Known good baseline | +| No dequant inverse, no graph rot | 194 | Fully rotated, no compensation | + +### Graph Q Rotation Only (no V inverse) +| Storage | mul_mat gives | PPL | Notes | +|---------|---------------|-----|-------| +| TURBO_ROTATION_RT for Q (R^T@q) | R^T@q = R_inv@q | 157 | Wrong direction for K matching | +| TURBO_ROTATION_R for Q (R@q) | R@q | 157 | Correct direction but V still rotated | + +Both give ~157 because V output is still in rotated space (no V inverse). The attention scores differ but V corruption dominates. + +### Graph Q + V Rotation (both active) +| Q Storage | V Storage | Q gives | V gives | PPL | +|-----------|-----------|---------|---------|-----| +| RT (R^T@q) | R (R@cur) | R^T@q (wrong) | R@cur (wrong) | 26.6 | +| R (R@q) | RT (R^T@cur) | R@q (correct) | R^T@cur (correct) | **23.5** | + +Corrected storage is better (23.5 < 26.6), confirming direction matters. But 23.5 >> 6.19. + +### Isolation Tests (dequant inverse ON) +| Graph rotation | PPL | Notes | +|----------------|-----|-------| +| Q rot only (R@q) | 10.5 | Q rotation works - degrades quality vs un-rotated K | +| V inv only (R^T@cur) | 26.6 | V inverse works - degrades quality vs un-rotated V | + +### Key Finding +Both rotations mechanically work (modify output, verified with scale(2.0) test). But the full pre-rotate-queries approach (correct Q + correct V) gives PPL 23.5, NOT 6.19. + +## Unsolved: Why 23.5 Instead of 6.19? + +The math proves the approaches should be equivalent: +- Dequant inverse: x_dequant = R^{-1}(quant(R(x))) ~ x + R^{-1}(epsilon) +- Graph rotation: Q=R(q), K=quant(R(k)), V=quant(R(v)), out=R^{-1}(attn(Q,K,V)) +- Error magnitudes are identical (orthogonal rotation preserves norms) + +Hypotheses to investigate: +1. **Flash attention precision**: FA kernel may accumulate differently with rotated vs un-rotated values +2. **Block boundary effects**: 128-element rotation groups split across flash attention tiles differently +3. **Metal mul_mat precision**: GPU f32 mul_mat on (128,128)@(128,N) may have precision issues +4. **Non-contiguous tensor handling**: ggml_cont + reshape chain may not preserve data correctly on Metal +5. **Graph optimizer interference**: ggml graph optimizer may simplify/skip the rotation ops + +## Current Status +- Dequant inverse rotation RESTORED (PPL = 6.194, speed ~10.7 tok/s) +- Graph rotation code preserved as TODO comments for future investigation +- Virtual method infrastructure (get_turbo_rot_forward/inverse) remains in place +- Rotation tensor allocation and initialization remains in KV cache diff --git a/docs/pre-rotate-queries-plan.md b/docs/pre-rotate-queries-plan.md new file mode 100644 index 0000000..8442f4b --- /dev/null +++ b/docs/pre-rotate-queries-plan.md @@ -0,0 +1,132 @@ +# Pre-Rotate Queries Optimization Plan + +## The Trick + +`⟨q, R^T · centroids[idx]⟩ = ⟨R·q, centroids[idx]⟩` + +Instead of inverse-rotating every K vector during dequant (128 WHT ops × every KV position), +rotate Q once before attention (128 WHT ops × 1 query). + +## Changes Required + +### 1. Store rotation parameters in model/context +- [ ] Add WHT sign arrays to llama_context or llama_kv_cache +- [ ] Accessible during graph building via hparams or cparams + +### 2. Insert Q rotation in attention graph (llama-graph.cpp:1799) +- [ ] After `q = ggml_permute(ctx0, q, 0, 2, 1, 3);` +- [ ] Check: `if (k->type == GGML_TYPE_TURBO3_0 || k->type == GGML_TYPE_TURBO4_0)` +- [ ] Apply: `q = ggml_mul_mat(ctx0, rotation_matrix_tensor, q);` +- [ ] The rotation matrix is 128×128 per head_dim — shared across all heads/layers +- [ ] Q shape at this point: (head_dim, n_tokens, n_head, n_stream) +- [ ] Rotation operates on the head_dim dimension (dim 0) + +### 3. Simplify Metal dequant to skip rotation +- [ ] `dequantize_turbo3_0_t4`: just centroid lookup + norm scale (no WHT) +- [ ] `dequantize_turbo4_0_t4`: same +- [ ] `turbo3_dequantize_full_block`: remove turbo_rotate_inverse calls +- [ ] Keep QJL correction? Dejan.ai says NO for drop-in replacement + +### 4. Simplify Metal quantize similarly +- [ ] `quantize_turbo3_0`: still needs forward WHT (writing to KV cache) +- [ ] BUT: if we pre-rotate Q, the stored K values are in ROTATED space +- [ ] So quantize writes rotated centroids WITHOUT inverse rotation +- [ ] And dequant returns rotated centroids directly +- [ ] The Q·K dot product works because Q is also rotated + +### 5. Handle QJL decision +- [ ] Option A: Drop QJL entirely (MSE-only, all bits to Lloyd-Max) + - Simpler dequant (just centroid lookup) + - Dejan.ai found this gives BETTER quality than naive QJL add-back + - Saves 1 bit per element (or gives 1 extra bit to MSE) +- [ ] Option B: Keep QJL but handle in fused kernel + - More complex, but theoretically better IP preservation + - Requires keeping the two-part representation separate + +## Expected Performance Impact + +### Current (with WHT rotation in dequant): +- Dequant: unpack + centroid + inverse WHT + QJL WHT + scale = ~2000 ops/block +- MoE gen: 10.7 tok/s (8× slower than q8_0) + +### After pre-rotate-queries (no rotation in dequant): +- Dequant: unpack + centroid + scale = ~256 ops/block +- Expected: close to q4_0 speed (84.5 tok/s on MoE) +- The only overhead vs q4_0 is the 128-element block vs 32-element block + +### After dropping QJL (MSE-only): +- Dequant: unpack + centroid + scale = ~200 ops/block +- Block struct shrinks: no signs array needed +- Even simpler, even faster + +## Test Plan + +- [ ] A/B test 1: Current turbo3 vs MSE-only turbo3 (drop QJL) — quality + speed +- [ ] A/B test 2: Current turbo3 vs pre-rotated-Q turbo3 — speed +- [ ] A/B test 3: Pre-rotated-Q + MSE-only — the full optimization +- [ ] Quality: cosine sim, attention output comparison +- [ ] Speed: tok/s on MoE + Qwopus + +## Implementation Order + +1. **Quick win**: Remove WHT from dequant, test speed (broken quality but measures ceiling) +2. **MSE-only test**: Remove QJL from both quant and dequant, test quality +3. **Pre-rotate-Q**: Insert rotation in graph, full correct implementation + +## Codex Review Findings + +### Finding 1: R must be fixed and shared — VERIFIED ✅ +Our WHT sign arrays are `constant float` in the Metal shader — same for ALL blocks, +layers, and heads. The rotation is indeed global and fixed. The optimization is valid. + +### Finding 2: Use WHT for Q rotation, not dense matmul — AGREED +Insert a custom ggml op for WHT rotation, not ggml_mul_mat with a dense matrix. +Or implement as a Metal compute kernel that does in-place WHT on the Q tensor. + +### Finding 3: Normalization must be consistent — VERIFIED ✅ +Our WHT is normalized by 1/sqrt(128). The same normalization is applied in both +quantize (forward WHT) and dequant (inverse WHT). Pre-rotating Q with the same +normalized WHT preserves the dot product scale. + +### Finding 4: Gate on KV cache config type, not transient tensor type — GOOD POINT +The `k->type` at graph build time may be F16 (intermediate). Need to check the +actual KV cache type from `cparams.cache_type_k` or similar. + +### Finding 5: QJL removal is a new format — AGREED +Will test as a separate path, not modify existing turbo3 behavior. + +### Finding 6: Prompt/decode consistency — MUST VERIFY +The quantize path (SET_ROWS) and dequant path (flash_attn) must agree on whether +rotation is stored or applied at read time. + +### Finding 7: Backend parity — ACKNOWLEDGED +CPU backend still uses dense rotation in C code. Metal uses WHT. These produce +different results. Need to align or document the divergence. + +## Implementation Progress + +### Challenge: inserting rotation tensor in ggml graph +- ggml tensors don't have allocated data at graph build time +- Can't memcpy into tensor during graph construction +- Need to create the rotation matrix as a persistent tensor at init time +- Store in KV cache object or context, access via mctx in graph builder + +### Simplest path forward +1. Add a `ggml_tensor * turbo_rotation` member to `llama_kv_cache` +2. Allocate and fill during KV cache construction when type_k is turbo +3. In `build_attn_mha`, access via `mctx` and apply `ggml_mul_mat` to Q +4. Strip rotation from Metal dequant (already tested — gives 49 tok/s) + +### What we know for certain (from speed ceiling test) +- Removing rotation from dequant: 10.7 → 49.1 tok/s (confirmed on Metal) +- The math is sound (R is fixed, orthogonal, normalized — verified by codex) +- The insertion point is llama-graph.cpp line ~1799 +- k->type at that point IS the KV cache storage type (verified) + +### Files to modify +1. `src/llama-kv-cache.h` — add turbo_rotation tensor member +2. `src/llama-kv-cache.cpp` — allocate + fill rotation during construction +3. `src/llama-kv-cache-context.h` — expose rotation to graph context +4. `src/llama-graph.cpp` — apply rotation to Q before flash_attn +5. `ggml/src/ggml-metal/ggml-metal.metal` — strip rotation from dequant +6. `src/turbo-rotation-data.h` — pre-computed 128x128 rotation matrix diff --git a/docs/quality-benchmarks.md b/docs/quality-benchmarks.md new file mode 100644 index 0000000..f480acd --- /dev/null +++ b/docs/quality-benchmarks.md @@ -0,0 +1,312 @@ +# TurboQuant+ Quality Benchmarks + +## Why These Benchmarks Matter + +Speed without quality is useless. We claim 4.6× KV cache compression at 91-97% of q8_0 speed. +But we have ZERO quantitative quality data on the actual llama.cpp build. "Coherent text" and +Python cosine similarity (0.985) are not sufficient. + +The llama.cpp CONTRIBUTING.md requires for new quant types: +1. Perplexity vs FP16/BF16 +2. KL divergence data +3. Performance data via llama-bench + +Papers (KIVI, QJL, RotateKV) all use wikitext-2 perplexity as the primary metric. +Prince Canuma validated with NIAH at 8K/32K/64K context. + +## What We're Measuring + +### 1. Perplexity (wikitext-2) +- **What**: How well the model predicts next tokens over a standard text corpus +- **Why**: Gold standard for LLM quality. Lower = better. Sensitive to KV cache quality. +- **Target**: turbo3 within 1% of q8_0 perplexity. If >2%, quality problem. +- **Comparison**: f16, q8_0, q4_0, q4_1, q5_0, turbo3 + +### 2. KL Divergence vs f16 +- **What**: How different the output probability distribution is vs full precision +- **Why**: Measures distributional shift, not just top-token accuracy +- **Metrics**: mean KLD, delta-p RMS, same-top-p percentage +- **Required by**: llama.cpp CONTRIBUTING.md for upstream acceptance + +### 3. Passkey Retrieval (built-in NIAH) +- **What**: Can the model retrieve a specific passkey from a long haystack? +- **Why**: Tests attention pattern preservation over long context +- **Comparison**: f16 vs turbo3 at various context lengths and needle positions + +### 4. Generation Quality (qualitative) +- **What**: Side-by-side text generation comparison +- **Why**: Catches issues that aggregate metrics miss (repetition, coherence) + +## Test Configuration + +- **Model**: Qwen 3.5 35B-A3B MoE Q8_0 (primary), Qwopus v2 27B Q8_0 (secondary) +- **Hardware**: Apple M5 Max 128GB, Metal GPU +- **Dataset**: wikitext-2-raw (downloaded via scripts/get-wikitext-2.sh) +- **Context**: 512 tokens for perplexity (fast), 2048+ for NIAH +- **Chunks**: 8 for initial, 32 for final numbers + +## Commands + +### Download dataset +```bash +cd ~/local_llms/llama.cpp && bash scripts/get-wikitext-2.sh +``` + +### Perplexity suite +```bash +LLAMA=~/local_llms/llama.cpp/build-turbo/bin +MODEL=~/local_llms/models/Qwen3.5-35B-A3B-Q8_0.gguf +WIKI=~/local_llms/llama.cpp/wikitext-2-raw/wiki.test.raw + +for ct in f16 q8_0 q4_0 q4_1 q5_0 turbo3; do + echo "=== $ct ===" + $LLAMA/llama-perplexity -m $MODEL -f $WIKI -c 512 \ + -ctk $ct -ctv $ct -fa on --chunks 8 -ngl 99 \ + 2>&1 | tee results/ppl_${ct}.log +done +``` + +### KL Divergence +```bash +# Save f16 baseline logits +$LLAMA/llama-perplexity -m $MODEL -f $WIKI -c 512 \ + -ctk f16 -ctv f16 -fa on --chunks 8 -ngl 99 \ + --kl-divergence-base results/f16_logits.kld + +# Compute KL divergence for each cache type +for ct in q8_0 q4_0 turbo3; do + $LLAMA/llama-perplexity -m $MODEL -f $WIKI -c 512 \ + -ctk $ct -ctv $ct -fa on --chunks 8 -ngl 99 \ + --kl-divergence --kl-divergence-base results/f16_logits.kld \ + 2>&1 | tee results/kld_${ct}.log +done +``` + +## Results + +*To be filled after running benchmarks.* + +### Perplexity (wikitext-2, 512 context) + +| Cache Type | Bits/val | Perplexity | vs f16 | vs q8_0 | +|------------|----------|------------|--------|---------| +| f16 | 16 | — | baseline | — | +| q8_0 | 8 | — | — | baseline | +| q4_0 | 4 | — | — | — | +| q4_1 | 4.5 | — | — | — | +| q5_0 | 5.5 | — | — | — | +| **turbo3** | **3.5** | — | — | — | + +### KL Divergence vs f16 + +| Cache Type | Mean KLD | Delta-p RMS | Same Top-p % | +|------------|----------|-------------|-------------| +| q8_0 | — | — | — | +| q4_0 | — | — | — | +| **turbo3** | — | — | — | + +### Passkey Retrieval + +| Cache Type | 1K | 2K | 4K | 8K | +|------------|----|----|----|----| +| f16 | — | — | — | — | +| q8_0 | — | — | — | — | +| turbo3 | — | — | — | — | + +## INITIAL RESULTS — QUALITY FAILURE + +### Perplexity (wikitext-2, 512 context, 8 chunks) + +| Cache Type | Bits/val | Perplexity | vs f16 | +|------------|----------|------------|--------| +| f16 | 16 | 6.121 | baseline | +| q8_0 | 8 | 6.111 | -0.16% (better!) | +| q4_0 | 4 | 6.142 | +0.34% | +| **turbo3** | **3.5** | **165.6** | **+2607%** ❌ | + +**turbo3 perplexity is 27× worse than f16. This is catastrophic.** + +The model generates "coherent-looking" text but the actual predictions are garbage. +Speed benchmarks were meaningless — we were measuring how fast the model produces wrong answers. + +### Root cause investigation needed +The block size 32 change + MSE-only + pre-rotate-queries may have introduced a bug: +1. The norm handling: storing full 128-element group norm in each 32-element block + but dequant treats it as a per-block norm — scale mismatch? +2. The pre-rotate-queries: rotation matrix might not match the quantize rotation +3. The 3-bit index split (qs + signs) might have a packing error +4. The non-vec flash attention instantiation might use wrong nl parameter + +MUST FIX before claiming any quality results. + +## Quality Bisection Log + +Goal: find which change broke perplexity (165.6 vs 6.1 baseline). + +Changes applied (in order): +1. Pre-rotate-queries: rotation moved from dequant to Q in attention graph +2. MSE-only: dropped QJL, 3-bit centroids (2-bit qs + 1-bit signs) +3. Block size 32: storage blocks shrunk from 128 to 32 elements + +Bisection plan: +- [ ] Test A: revert block 32 → 128, keep MSE-only + pre-rotate +- [ ] Test B: revert MSE-only → 2bit+QJL, keep block 128 + pre-rotate +- [ ] Test C: revert pre-rotate → rotation in dequant, keep block 128 + 2bit+QJL +- [ ] If A/B/C all still bad: check the rotation matrix consistency + +## Bisection Results + +### Test A: Disable Q rotation → PPL 165.6 (STILL BAD) +Pre-rotate-queries is NOT the problem. Q rotation has no effect on perplexity. + +### Root Cause Found: V CACHE IN ROTATED SPACE + +The bug: both K and V caches go through the same turbo3 quantize (which rotates) +and dequant (which returns rotated values without inverse rotation). + +For K cache: this is correct because Q is also rotated → = ✓ +For V cache: this is WRONG because the attention output = attn_weights @ V +needs to be in ORIGINAL space, not rotated space. + +The fix: V cache must NOT be rotated during quantize, or the attention output +must be inverse-rotated after the V multiplication. + +Options: +1. Skip rotation for V in quantize — only rotate K +2. Add inverse rotation after attention output: out = R^T @ (attn @ R*V) = attn @ V ✓ +3. Use f16 for V cache (only compress K) — mixed types need support + +Option 1 is simplest: in the SET_ROWS kernel, check if this is K or V and +skip rotation for V. The dequant already skips rotation. +Problem: the SET_ROWS kernel doesn't know if it's writing K or V. + +Option 2 is cleanest: add R^T multiplication to attention output in the graph. +This is one ggml_mul_mat per layer (same as Q rotation). Cost is negligible. + +### Root Cause #1: V cache dequant returns rotated-space values +Python verified: cosine(input, dequant_output) = 0.02 (garbage) +After inverse rotation: cosine = 0.987 (correct) + +### Root Cause #2: dynamic_cast to llama_kv_cache_context FAILS +The Qwen 3.5 MoE uses llama_memory_hybrid_context (not llama_kv_cache_context). +The dynamic_cast returns null, so the V inverse rotation never executes. +This also means the Q pre-rotation never executed — explaining why the +chat server seemed to work (no rotation applied = raw quantize/dequant, which +happens to produce plausible-looking text at low context but garbage perplexity). + +### Fix needed +1. Store turbo rotation tensors in llm_graph_context (not KV cache) + OR access them through the hybrid memory interface +2. Apply V inverse rotation through a non-cast path +3. Verify Q rotation also executes (it was also using the same cast) + +### The "coherent text" mystery explained +The model produced grammatical text because: +- Without Q rotation, attention scores are computed in wrong space but still produce + some attention pattern (just not the right one) +- Without V inverse rotation, the output is in rotated space which is orthogonal to + the correct space — but each layer compounds the error +- Short conversations look plausible but perplexity reveals the content is wrong + +### Bisect: block 128 → STILL PPL=165.6 +Block size is NOT the issue. PPL identical at block 128 and block 32. + +### Bisect: disable Q rotation → STILL PPL=165.6 +Pre-rotate-queries is NOT the issue. PPL identical with/without Q rotation. + +### Root cause #2 confirmed: dynamic_cast fails for MoE hybrid memory +The Q rotation and V inverse rotation NEVER execute because mctx is +llama_memory_hybrid (not llama_kv_cache). The kv_cache IS inside the +hybrid object but we can't reach it. + +### What this means +Without Q rotation and without V inverse rotation: +- K cache is quantized in ROTATED space, dequanted in ROTATED space +- Q is in ORIGINAL space (no pre-rotation) +- = GARBAGE (cosine ≈ 0.02) +- V cache similarly in rotated space, output never inverse-rotated + +The entire time, K/V values were in the wrong coordinate system. +The "coherent text" was an illusion from the language model's resilience. + +### Fix approach +Need to get rotation tensors through the hybrid memory hierarchy. +Virtual methods on llama_memory_context_i are the cleanest path. + +## FIX: Inverse rotation restored in dequant + +### Perplexity (wikitext-2, 512 context, 8 chunks) + +| Cache Type | Perplexity | vs f16 | vs q8_0 | +|------------|------------|--------|---------| +| f16 | 6.121 | baseline | — | +| q8_0 | 6.111 | -0.16% | baseline | +| q4_0 | 6.142 | +0.34% | +0.51% | +| **turbo3** | **6.194** | **+1.19%** | **+1.36%** | + +**turbo3 within 1.4% of q8_0. Quality target MET.** + +### Root cause of previous failure +Pre-rotate-queries never executed because: +1. Q tensor ne[0]=256 (GQA concatenated heads) vs rotation ne[0]=128 +2. MoE hybrid memory context fails dynamic_cast to llama_kv_cache_context + +### Current state +- Inverse rotation restored in dequant with fp16 WHT butterfly (correct quality, faster) +- Pre-rotate-queries abandoned: WHT and RoPE don't commute (Gemini 3 Pro insight) +- Speed optimized via fp16 WHT + SIMD cooperative dequant + +## Top-of-Tree Quality and Speed (2026-03-25) + +Model: Qwen3.5-35B-A3B-Q8_0 | Hardware: Apple M5 Max 128GB | Flash Attention ON + +### Quality (wikitext-2, 512 context, 8 chunks) + +| Cache Type | Bits/val | Compression | Perplexity | vs q8_0 | +|------------|----------|-------------|------------|---------| +| f16 | 16 | 1.0x | 6.121 | -0.16% | +| q8_0 | 8 | 2.0x | 6.111 | baseline | +| q4_0 | 4 | 4.0x | 6.142 | +0.51% | +| **turbo3** | **3.25** | **4.9x** | **6.195** | **+1.4%** | + +### Speed (wikitext-2, 512 context, 32 chunks prefill) + +| Cache Type | Prefill tok/s | vs q8_0 | Notes | +|------------|--------------|---------|-------| +| q8_0 | 2694 | 1.00x | baseline | +| turbo3 (no rot) | 1577 | 0.59x | rotation stripped (wrong quality) | +| turbo3 (fp32 WHT) | 739 | 0.27x | old, before fp16 optimization | +| **turbo3 (fp16 WHT)** | **1074** | **0.40x** | **current top-of-tree** | + +### Summary +- **4.9x compression, 1.4% quality loss, 40% of q8_0 prefill speed** +- fp16 WHT gave 45% speedup over fp32 WHT +- Remaining speed gap: 2.51x vs q8_0 (from redundant full-block dequant in non-vec FA path) + +## Current State Summary (after quality fix) + +### What works: +- turbo3 KV cache compression: 4.6× compression ratio +- Perplexity: 6.194 (1.4% of q8_0) — QUALITY TARGET MET +- MSE-only mode: 3-bit PolarQuant, no QJL — better quality than paper's Algorithm 2 +- WHT rotation in dequant: correct inverse rotation restoring original space +- Block size 128: works on Metal for both MoE and Dense models +- Works for MoE (hybrid memory), Dense, and ISWA models + +### Speed: +- ~10.7 tok/s MoE gen (8× slower than q8_0 85.5) +- Bottleneck: inverse WHT rotation in dequant (called per block per attention) + +### What doesn't work: +- Pre-rotate-queries optimization: never executed because + 1. Q tensor ne[0]=256 (GQA concatenated heads) vs rotation ne[0]=128 + 2. Need per-head reshape before rotation +- Previous speed numbers (51-77 tok/s) were INVALID — measured garbage output speed + +### Lessons learned: +1. ALWAYS run perplexity before claiming quality +2. "Coherent text" is NOT quality validation +3. dynamic_cast fails silently for MoE hybrid memory +4. GQA models concatenate heads in ne[0] — can't apply per-head ops without reshape +5. #include in ggml-metal.metal causes silent CPU fallback +6. Block size matters for speed but NOT for quality (both 32 and 128 had same PPL) diff --git a/docs/turbo-speed-investigation.md b/docs/turbo-speed-investigation.md new file mode 100644 index 0000000..409cff0 --- /dev/null +++ b/docs/turbo-speed-investigation.md @@ -0,0 +1,400 @@ +# TurboQuant Speed Investigation + +## Problem Statement + +TurboQuant generates at 2.4 tok/s vs 85.0 tok/s for q8_0 on Qwen 3.5 35B-A3B MoE (M5 Max). +That's a 35× regression. Compression target (4.9×) is met, but speed makes it impractical. + +## Root Cause Analysis + +### Why it's slow + +The flash attention kernel calls `dequantize_turbo3_0(block_ptr, il, ®)` once per chunk: +- `type4x4` path: 16 elements per call, 128/16 = **8 calls per block** +- `type4` path (flash_attn_ext_vec): 4 elements per call, 128/4 = **32 calls per block** + +Each call does the FULL 128-element dequantize: +1. Unpack 128 indices from packed bits → centroid lookup +2. Inverse WHT rotation (7 butterfly passes × 128 = 896 ops) +3. Unpack 128 QJL signs +4. Inverse WHT rotation on signs (another 896 ops) +5. Combine and scale + +Total per call: ~2000 ops + 256 bytes stack allocation +Total per block: 8-32 × 2000 = **16,000-64,000 ops** (should be 2000) + +### Comparison with q8_0 + +q8_0 dequantize: `x[i] = qs[i] * d` — 1 multiply per element, no stack, no rotation. +Per block (32 elements): 32 ops. Per 128 elements: 128 ops. + +TurboQuant is doing 64,000 ops where q8_0 does 128. That's **500× more compute per block**. + +Even with WHT (18× less than dense matvec), we're at 64,000/18 ≈ 3,500 ops vs 128. +That's still **27× more compute**, close to the measured 35× regression. + +The extra ~8× gap is from stack allocation overhead (128-float arrays created/destroyed +32× per block) and memory bandwidth (reading the constant sign arrays 32× per block). + +--- + +## Investigation Plan + +### Approach A: Modify flash attention to dequantize once per block +- [ ] **A1**: Study the flash attention kernel template structure + - Find where `deq_k` and `deq_v` are called + - Understand the iteration pattern: which blocks, which chunks + - Identify where to insert a "pre-dequantize" step +- [ ] **A2**: Add threadgroup memory buffer for pre-dequantized blocks + - Allocate `threadgroup float turbo_deq_buf[128]` in flash attention kernel + - Before the per-chunk loop, call `turbo3_dequantize_full_block()` once into this buffer + - Replace per-chunk dequantize calls with reads from the buffer +- [ ] **A3**: Test with turbo-specific flash attention kernel instantiation + - The generic template uses function pointers for dequantize + - May need a specialized kernel that bypasses the per-chunk pattern +- [ ] **A4**: Benchmark after each change + +### Approach B: Custom flash attention kernel for turbo types +- [ ] **B1**: Fork `kernel_flash_attn_ext_vec` into `kernel_flash_attn_ext_vec_turbo` + - Remove the generic dequantize function pointer + - Inline turbo-specific dequantize at the block level + - Use threadgroup memory for the dequantized block +- [ ] **B2**: Replace per-chunk reads with direct buffer indexing + - Instead of `deq_k(pk4x4 + block_idx, chunk_idx, tmp)`, do `tmp = buf[chunk_idx * 16 : (chunk_idx+1) * 16]` +- [ ] **B3**: Register the custom kernel in pipeline lookup +- [ ] **B4**: Benchmark + +### Approach C: Restructure dequantize to amortize across chunks +- [ ] **C1**: Change the dequantize API to accept a pre-allocated buffer + - `dequantize_turbo3_0(block_ptr, il, ®, thread float * shared_buf)` + - On first call (il==0), fill the buffer; on subsequent calls, read from it + - Problem: can't change the function signature without changing all callers +- [ ] **C2**: Use the `nl_k` parameter differently + - Currently nl_k=32 for turbo3 (128/4). What if we set nl_k=1 and return all 128 at once? + - Would need the caller to handle 128-element chunks + - Probably breaks the template assumptions + +### Approach D: Reduce per-call overhead without architecture changes +- [ ] **D1**: Precompute the WHT butterfly as a lookup table + - Instead of 7 butterfly passes, use a precomputed 128-element permutation + - Trade memory for compute: 128 × 4 bytes = 512 bytes constant + - Won't help much since WHT is already fast +- [ ] **D2**: Pack the dequantize tighter — reduce stack allocations + - Merge the centroid lookup + WHT into a single pass + - Avoid allocating separate `recon[128]` and `signs_f[128]` — interleave +- [ ] **D3**: Use half precision for intermediate calculations + - `half` arithmetic is 2× faster on Apple Silicon + - May reduce quality slightly but worth testing + +--- + +## Expected Outcomes + +| Approach | Expected Speedup | Effort | Risk | +|----------|-----------------|--------|------| +| A (modify kernel) | 8-32× (eliminate redundant calls) | Medium | Medium — need to understand kernel internals | +| B (custom kernel) | 8-32× + optimal memory access | High | Low — clean separation | +| C (restructure API) | 8-32× | Low | High — may break template | +| D (reduce overhead) | 2-3× | Low | Low | + +**Recommended order**: D first (quick wins), then A or B (the real fix). + +Target: D → 5-8 tok/s, then A/B → 20-40 tok/s. + +--- + +## Progress Log + +### 2026-03-25: Initial investigation +- Dense matvec: 2.4 tok/s (35× slower than q8_0) +- WHT rotation: 2.4 tok/s (same — bottleneck is redundant calls, not per-call compute) +- Root cause confirmed: dequantize called 8-32× per block by flash attention +- Codex + roast reviewed WHT implementation: correct, no bugs + +### Next: Start with Approach D (reduce per-call overhead) + +### 2026-03-25: simd_broadcast attempt +- Added simd_broadcast fast path for K and V dequant (nl_k==32 && DK==128) +- Thread 0 dequantizes, broadcasts 128 floats via simd_broadcast loop +- **Result: still 2.4 tok/s** — the 128-iteration simd_broadcast loop per cc iteration + is itself expensive. 32 cc iterations × 128 broadcasts = 4096 simd_broadcast calls per block. +- Codex review caught: DK>128 OOB bug (fixed), turbo4 using turbo3 dequant (fixed), + uninitialized turbo_buf on non-lane-0 (fixed with zero-init) +- **Conclusion**: simd_broadcast is wrong tool. Need threadgroup memory instead. + +### Next: try threadgroup memory approach +- Allocate extra threadgroup memory in FATTN_SMEM +- One thread writes 128 floats to threadgroup, barrier, all threads read +- This reduces to 1 dequant + 1 barrier per cc iteration instead of 128 broadcasts + +### 2026-03-25: threadgroup memory attempt +- Replaced simd_broadcast with threadgroup memory + simdgroup_barrier +- Thread 0 dequantizes into threadgroup, barrier, all threads read +- **Result: still 2.4 tok/s** +- Eliminating 31/32 redundant dequant calls had NO effect on speed +- This means the dequant cost itself (even 1× per block) is NOT the bottleneck +- Or the bottleneck is elsewhere entirely (SET_ROWS quantize? block size overhead?) + +### Hypothesis: block size 128 vs 32 causes structural overhead +- q8_0 block size = 32, turbo block size = 128 +- The flash attention kernel processes DK4/NL elements per thread per cc iteration +- For q8_0: DK4/NL = 32/8 = 4 iterations (inner ii loop runs 4×) +- For turbo: DK4/NL = 32/32 = 1 iteration (inner ii loop runs 1×) +- But NL = 32 for turbo vs NL = 4 for q8_0 (32/8=4, C=32, NE=1→NW/NE=32 for both) +- Actually NL = NW/NE = 32/1 = 32 for both... so DK4/NL should be the same? +- Wait: for q8_0, nl_k=8 (32 elements / 4 per t4 = 8 chunks). DK4 = 128/4 = 32. DK4/NL = 32/32 = 1. +- So BOTH q8_0 and turbo have DK4/NL = 1 iteration in the inner loop. +- The only difference is the dequant function itself. + +### Next: profile whether the bottleneck is in dequant or elsewhere +- Test with a no-op dequant (return zeros) to measure the kernel overhead +- If still slow → bottleneck is NOT dequant, it's structural + +### 2026-03-25: NO-OP DEQUANT TEST — CRITICAL FINDING +- Made dequantize_turbo3_0 and dequantize_turbo3_0_t4 return all zeros (no-op) +- **Result: STILL 2.4 tok/s** +- **The dequant function is NOT the bottleneck.** +- Even with zero-cost dequant, turbo3 is 35× slower than q8_0. +- The bottleneck is ELSEWHERE: + 1. SET_ROWS (quantize) path — 128×128 WHT per block during KV cache write + 2. Kernel dispatch overhead for 128-element blocks + 3. Memory layout/alignment issues with the larger block structure + 4. The quantize kernel (kernel_set_rows_turbo) doing full WHT per KV write + +### New investigation target: SET_ROWS (quantize path) +- The quantize_turbo3_0 function in the Metal shader does: + 1. Normalize (128 ops) + 2. Copy to working buffer (128 ops) + 3. Forward WHT rotation (896 ops) + 4. Centroid lookup + pack (128 ops) + 5. Inverse WHT rotation (896 ops) + 6. Residual (128 ops) + 7. Forward QJL WHT rotation (896 ops) + 8. Sign pack (128 ops) + Total: ~3200 ops per block in the QUANTIZE path +- This runs for every token at every layer at every head +- For Qwen 3.5 MoE: 28 layers × 8 heads = 224 quantize calls per generated token +- Each quantize does 3200 ops = 716,800 ops per token just for KV cache write + +### Next: profile SET_ROWS path / try no-op quantize test + +### 2026-03-25: THE BUG — #include caused CPU fallback!!! 🎉 +- `#include "turbo-wht.h"` in ggml-metal.metal causes Metal JIT to fail +- The model SILENTLY falls back to CPU for all operations +- ALL previous benchmarks (2.4 tok/s) were measuring CPU, not Metal GPU +- **After inlining turbo-wht.h:** + - MoE: 2.4 → 10.7 tok/s gen (4.5× improvement, now on Metal) + - Qwopus: 2.4 → 5.3 tok/s gen (2.2× improvement) + - Prompt speeds: 4× → 60-67 tok/s (MoE), 3.5 → 30 tok/s (Qwopus) +- Remaining gap vs q8_0: 8× (MoE), 3.3× (Qwopus) — down from fake 35× +- **Rule: NEVER use #include in ggml-metal.metal — always inline** +- This is the same bug we hit with turbo-matrices.h earlier and forgot to apply to WHT + +### Real benchmarks (proper Metal GPU): + +| Model | Cache | Prompt | Gen | Compression | vs q8_0 | +|-------|-------|--------|-----|-------------|---------| +| MoE | q8_0 | 222.8 | 85.5 | 2.0× | 1.00× | +| MoE | turbo3 | 67.3 | 10.7 | 4.9× | 0.13× | +| Qwopus | q8_0 | 83.1 | 17.6 | 2.0× | 1.00× | +| Qwopus | turbo3 | 29.8 | 5.3 | 4.9× | 0.30× | + +### Remaining optimization targets (ACTUAL, not fake): +- 8× gap on MoE gen (85.5 vs 10.7) +- 3.3× gap on Qwopus gen (17.6 vs 5.3) +- WHT rotation is O(d log d) per block — already optimized +- The per-chunk redundant call pattern still wastes ~8× on dequant +- BUT: no-op dequant test was on CPU, need to redo on Metal + +### 2026-03-25: Upstream research findings +- TurboQuant feature request exists upstream: ggml-org/llama.cpp#20977 +- mudler has experimental fork — should compare notes +- New quant type contribution guidelines require perplexity + KL divergence data +- The CPU fallback pattern is known: MXFP4 hit the same issue (PR #20609) +- Metal Tensor API optimization (PR #20962) could give us another ~26% on mul_mat +- Our current 8× gap on MoE / 3.3× on Dense is REAL Metal overhead, not a bug +- The WHT rotation cost per block is the genuine bottleneck now + +### Summary of investigation +| Test | Gen tok/s | What we learned | +|------|-----------|-----------------| +| q8_0 baseline | 85.5 | Target speed | +| Dense matvec rotation | 2.4 | Was actually CPU fallback | +| WHT rotation | 2.4 | Still CPU fallback | +| simd_broadcast opt | 2.4 | Still CPU fallback | +| threadgroup opt | 2.4 | Still CPU fallback | +| No-op dequant | 2.4 | Still CPU fallback | +| No-op quantize | 2.4 | Still CPU fallback | +| **Inlined WHT (proper Metal)** | **10.7** | **THE FIX — #include caused CPU fallback** | + +Rule: NEVER use #include in ggml-metal.metal. Always inline. + +### 2026-03-25: Upstream competitive intelligence — GAME CHANGERS + +**Finding 1: QJL residual kills quality when naively re-added** +Dejan.ai found cosine similarity drops to 0.69 with naive QJL add-back. +Must use MSE-only for drop-in replacement, or fused kernel for QJL. + +**Finding 2: Pre-rotate queries, not keys (THE BIG WIN)** +` = ` +Rotate query ONCE, then dequant is just a centroid table lookup. +No rotation in the hot dequant path. Would make turbo3 as fast as q4_0. + +**Finding 3: WHT abandoned by everyone** +Both Dejan.ai and mudler tried WHT, both switched to dense QR. +But with pre-rotate-queries, rotation method doesn't matter. + +**Finding 4: unixsysdev gets -0.8% speed loss (CUDA, block=32)** +Pure polar coordinate encoding, no random rotation, fused kernel. + +**Finding 5: We're the only Metal implementation** +Everyone else is CPU-only or CUDA. Our Metal kernels are unique. + +### Immediate action items: +- [ ] Test MSE-only (no QJL) — simplify dequant, may improve quality + speed +- [ ] Implement pre-rotate-queries — eliminates rotation from dequant entirely +- [ ] Consider block size 32 for better GPU parallelism + +### 2026-03-25: Speed ceiling test — no rotation in dequant (proper Metal) +- Removed turbo_rotate_inverse from dequant (quality broken but measures ceiling) +- **Result: 49.1 tok/s gen (4.6× improvement over 10.7)** +- Prompt: 162.6 tok/s (2.4× improvement over 67.3) +- This is 57% of q8_0 speed (85.5 tok/s) +- Confirms: removing rotation from dequant via pre-rotate-queries + would get us from 10.7 → ~49 tok/s +- Remaining gap (49 vs 85) is from block size 128 + QJL overhead +- REVERTED change after measurement + +### 2026-03-25: PRE-ROTATE-QUERIES — THE WIN 🎉 +- Implemented: store R^T in KV cache, apply ggml_mul_mat(R_T, q) in build_attn_mha +- Stripped turbo_rotate_inverse from Metal dequant +- Codex caught: buffer_clear zeroing rotation after init (fixed) +- **MoE gen: 10.7 → 51.4 tok/s (4.8× speedup, 60% of q8_0)** +- **Qwopus gen: 5.3 → 14.6 tok/s (2.8× speedup, 83% of q8_0)** +- Prompt: 67 → 160 tok/s (MoE), 30 → 68 tok/s (Qwopus) + +### FINAL SUMMARY + +| Stage | MoE tok/s | Qwopus tok/s | What happened | +|-------|-----------|-------------|---------------| +| First benchmark | 2.4 | 1.3 | CPU fallback (#include bug) | +| Fixed #include | 10.7 | 5.3 | Real Metal, WHT rotation overhead | +| Speed ceiling test | 49.1 | — | Confirmed rotation is the bottleneck | +| **Pre-rotate-queries** | **51.4** | **14.6** | **Rotation moved from dequant to Q** | +| q8_0 baseline | 85.5 | 17.6 | Target | + +**Total improvement: 21× from first benchmark to final (2.4 → 51.4 tok/s)** +**Compression: 4.9× throughout** + +### Key lessons learned +1. NEVER use #include in ggml-metal.metal — causes silent CPU fallback +2. Always verify Metal library loads successfully before benchmarking +3. The no-op test (set function to return zeros) is the fastest way to isolate bottlenecks +4. Pre-rotate-queries (from Dejan.ai) is the right architectural approach +5. Codex + roast reviews catch real bugs (buffer clear ordering, stale code, MSL limitations) +6. The ggml_mul_mat approach for Q rotation is clean and correct + +### 2026-03-25: CHANGE 1 — Drop QJL (MSE-only) ✅ +- Python validation: cosine 0.9508 → 0.9831, better on 99.3% of real KV vectors +- MoE gen: 51.4 → 62.2 tok/s (73% of q8_0) +- Qwopus gen: 14.6 → 15.5 tok/s (88% of q8_0) +- Qwopus prompt: 83.1 tok/s = 100% of q8_0! + +### 2026-03-25: CHANGE 2 — WHT for Q rotation — NOT NEEDED +- Speed ceiling test: removing Q rotation entirely gives 61.3 tok/s (vs 62.2 with it) +- The dense 128×128 ggml_mul_mat adds <1% overhead — negligible on Metal +- The remaining gap (62.2 vs 85.5 on MoE) is structural: + - Block size 128 (4 blocks per head) vs q8_0 block 32 (4 blocks per head too, but simpler dequant) + - 128 centroid lookups (3-bit unpack + table) vs 32 byte * scale (q8_0) + - This is the irreducible cost of the algorithm + +### FINAL PERFORMANCE + +| Model | q8_0 | turbo3 MSE-only | Speed | Compression | +|-------|------|-----------------|-------|-------------| +| MoE 35B | 85.5 tok/s | **62.2 tok/s** | **73%** | **4.9×** | +| Qwopus 27B | 17.6 tok/s | **15.5 tok/s** | **88%** | **4.9×** | + +### 2026-03-25: Block size diagnostic — THE ANSWER + +| Cache | Block | Gen tok/s | vs q8_0 | +|-------|-------|-----------|---------| +| q8_0 | 32 | 84.1 | 1.00× | +| q4_0 | 32 | 84.2 | 1.00× | +| q4_1 | 32 | 84.6 | 1.01× | +| q5_0 | 32 | 80.6 | 0.96× | +| turbo3 | 128 | 62.5 | 0.74× | +| turbo4 | 128 | 42.8 | 0.51× | + +CONCLUSION: q4_0 (block 32, 4-bit) runs at 100% of q8_0 speed. +The 26% gap on turbo3 is ENTIRELY from block size 128. +The quantization math (centroid lookup) adds zero overhead at block 32. + +Next: implement block size 32 variant of turbo3. +Challenge: rotation operates on head_dim=128, but blocks are 32. +Solution: quantize 128 elements with rotation, store as 4×32 blocks. +Dequant reads 32-element blocks without rotation (pre-rotate-queries). + +### 2026-03-25: BLOCK SIZE 32 — TARGET EXCEEDED 🎉 + +| Model | q8_0 | turbo3 block32 | vs q8_0 | Compression | +|-------|------|----------------|---------|-------------| +| MoE prompt | 222.8 | **218.5** | **98%** | 4.6× | +| MoE gen | 85.5 | **77.7** | **91%** | 4.6× | +| Qwopus prompt | 83.1 | **89.5** | **108%** | 4.6× | +| Qwopus gen | 17.6 | **17.0** | **97%** | 4.6× | + +Target was 75+ tok/s on MoE. Hit 77.7. Qwopus prompt is FASTER than q8_0. + +Codex post-commit review found 3 bugs (filed #29): +- TURBO_D was QK_TURBO3 (broke turbo4 C code) — fixed +- SET_ROWS kernel turbo3-specific but instantiated for turbo4 — tracked +- Tail block drop for non-128 head dims — tracked + +### COMPLETE JOURNEY + +| Optimization | MoE gen | vs q8_0 | +|-------------|---------|---------| +| CPU fallback (bug) | 2.4 | 3% | +| Real Metal | 10.7 | 13% | +| Pre-rotate-queries | 51.4 | 60% | +| MSE-only (drop QJL) | 62.2 | 73% | +| **Block size 32** | **77.7** | **91%** | +| q8_0 baseline | 85.5 | 100% | + +**32× total improvement. 4.6× compression at 91-97% speed.** + +## Pre-Rotate-Queries Investigation (2026-03-25) + +Attempted to move WHT inverse rotation from dequant to graph-level ops. Goal: eliminate per-block O(128) rotation, reclaiming speed from ~10.7 to ~77 tok/s. + +**Result: FAILED.** Graph-side rotation gives PPL 23.5 vs dequant-side 6.19. Root cause unknown despite extensive debugging. See `docs/pre-rotate-queries-investigation.md` for full test matrix. + +**Key findings:** +- ggml_mul_mat(A, x) with row-major stored A gives A @ x (verified with 2x2 test) +- Both Q rotation and V inverse work mechanically (verified via isolation tests) +- Rotation matrices verified correct (R^T @ R = I, matches Metal WHT signs) +- Model head_dim = 256 with 128-element rotation groups (2 groups per head) +- Codex caught: `clear(true)` zeros rotation tensors without reinit (fixed) +- Codex caught: header comments had wrong matrix orientation (fixed) + +**Current state:** Dequant inverse rotation restored. PPL = 6.194. Speed = ~10.7 tok/s. + +**Root cause found (Gemini 3 Pro):** WHT and RoPE don't commute. KV stores WHT(RoPE(K)) but graph rotation gives RoPE(WHT(Q)). Fix: apply WHT after RoPE inside attention layer, not in build_attn_mha. + +## FP16 WHT Optimization (2026-03-25) + +Switched WHT inverse rotation from fp32 to fp16 (Metal native half precision). Centroids fit in fp16 (max |val| = 0.19). Butterfly add/sub stays within fp16 range. + +Also implemented cooperative SIMD dequant for flash_attn_ext_vec (t4) path: each of 32 SIMD lanes unpacks only its 4 elements, then WHT butterfly runs across lanes via simd_shuffle. + +### Prefill Speed Results (Qwen3.5-35B-A3B, wikitext-2, 32 chunks) + +| Config | Prefill tok/s | vs q8_0 | PPL | +|--------|-------------|---------|-----| +| q8_0 baseline | 2694 | 1.00x | 5.41 | +| turbo3 no rotation | 1577 | 0.59x | - | +| turbo3 fp32 WHT (old) | 739 | 0.27x | 5.46 | +| **turbo3 fp16 WHT** | **1074** | **0.40x** | **5.47** | + +**45% speedup from fp16 WHT.** Gap to q8_0 narrowed from 3.65x to 2.51x. PPL within 0.06 of baseline. diff --git a/scripts/turbo-quality-gate.sh b/scripts/turbo-quality-gate.sh new file mode 100755 index 0000000..c5f55f4 --- /dev/null +++ b/scripts/turbo-quality-gate.sh @@ -0,0 +1,49 @@ +#!/bin/bash +# TurboQuant quality gate — run BEFORE pushing any changes +# Checks perplexity against q8_0 baseline. Fails if >5% degradation. +# +# Usage: bash scripts/turbo-quality-gate.sh + +set -e + +LLAMA=~/local_llms/llama.cpp/build-turbo/bin +MODEL=~/local_llms/models/Qwen3.5-35B-A3B-Q8_0.gguf +WIKI=~/local_llms/llama.cpp/wikitext-2-raw/wiki.test.raw + +if [ ! -f "$WIKI" ]; then + echo "Downloading wikitext-2..." + bash ~/local_llms/llama.cpp/scripts/get-wikitext-2.sh +fi + +echo "=== TurboQuant Quality Gate ===" +echo "" + +# Run turbo3 perplexity +echo "Running turbo3 perplexity..." +PPL_TURBO=$($LLAMA/llama-perplexity -m $MODEL -f $WIKI -c 512 -ctk turbo3 -ctv turbo3 -fa on --chunks 8 -ngl 99 2>&1 | grep "Final" | grep -oE 'PPL = [0-9.]+' | grep -oE '[0-9.]+') + +if [ -z "$PPL_TURBO" ]; then + echo "FAIL: Could not get turbo3 perplexity (crash or timeout)" + exit 1 +fi + +echo "turbo3 PPL = $PPL_TURBO" + +# q8_0 baseline: 6.111 (hardcoded from our measurements) +BASELINE=6.111 +MAX_ALLOWED=$(echo "$BASELINE * 1.05" | bc) # 5% threshold + +echo "q8_0 baseline = $BASELINE" +echo "Max allowed (5%) = $MAX_ALLOWED" + +PASS=$(echo "$PPL_TURBO < $MAX_ALLOWED" | bc) +if [ "$PASS" -eq 1 ]; then + echo "" + echo "PASS: turbo3 PPL $PPL_TURBO < $MAX_ALLOWED (within 5% of q8_0)" + exit 0 +else + echo "" + echo "FAIL: turbo3 PPL $PPL_TURBO > $MAX_ALLOWED (exceeds 5% threshold)" + echo "DO NOT PUSH. Fix quality before committing." + exit 1 +fi