mirror of
https://github.com/TheTom/turboquant_plus.git
synced 2026-10-22 02:53:13 -05:00
docs: move engineering docs here from llama.cpp fork
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) <noreply@anthropic.com> Co-Authored-By: tturney@psyguard.ai
This commit is contained in:
committed
1 parent
2d592a6f30
commit
445a3dc4c7
6 files changed
+975
-4
No files matched your search
@@ -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
|
||||
|
||||
|
||||
@@ -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
|
||||
@@ -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
|
||||
@@ -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 → <R*q, R*k> = <q, k> ✓
|
||||
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)
|
||||
- <original_Q, rotated_K> = 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)
|
||||
@@ -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)**
|
||||
`<q, R^T * c[idx]> = <R*q, c[idx]>`
|
||||
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.
|
||||
Executable
+49
@@ -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
|
||||
Reference in new issue
Block a user