Hi @elusznik, thanks for your contribution!

Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:

• Large PR: Large changes require prior discussion (e.g. an issue or RFC) and maintainers may not be able to review this PR as-is. Consider splitting it into smaller, focused PRs.

Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below.

Issue #20977

I've been working on extending unixsysdev's tq3_0 implementation with V cache support and flash attention. Repo here: https://github.com/animehacker/llama-turboquant

What this adds on top of unixsysdev's work:

• Normalization fix (1/32 → 1/√32 for the asymmetric K-side WHT)
• V cache compression (non-transposed storage + graph-side dequant to work around GGML's element-scatter path)
• Flash attention with tq3_0 (dequant tq3_0 → F32 → F16 in the attention graph, then use existing FA kernel)
• CPU backend F32 dequant path for pipeline parallelism

Tested on Llama-3.3-70B-Instruct-Q4_K_M, 2x RTX 3090:

• 72K context with tq3_0 K+V (4.57x compression)
• WikiText-2 PPL: 4.40 vs 4.09 baseline (+7.6%)

To be clear: this implements PolarQuant (Stage 1) only — WHT rotation + 3-bit Lloyd-Max. QJL residual correction is not included.

Paper with implementation details: https://oliverchurch.com/turboquant-for-ggml-achieving-4.57x-kv-cache-compression-in-llama.cpp.html

I've been working on a TurboQuant implementation in llama.cpp's GGML framework (CUDA backend, tested on Llama-3.3-70B with 2x RTX 3090s). A few findings that might be useful for the vLLM implementation:

1. The normalization factor for the WHT needs to be asymmetric: 1/√32 on the K-side during quantization, unnormalized on the Q-side. Using 1/32 (symmetric) produces plausible-looking but semantically broken output that's hard to catch without perplexity benchmarks.
2. V cache compression is essentially free in terms of quality. In my WikiText-2 benchmarks, K-only compression cost +6.6% PPL while adding V compression on top only added another +1% for 4.57x total compression.
3. For long context, dequanting to F16 and feeding into flash attention works well. Memory goes from O(n²) to O(n), which is what got us from 16K to 72K context.

Paper with implementation details: https://oliverchurch.com/turboquant-for-ggml-achieving-4.57x-kv-cache-compression-in-llama.cpp.html
Repo: https://github.com/animehacker/llama-turboquant

Happy to compare notes.

Addressed the actionable points from the Copilot review in 0aae7d78c:

• fixed TBQ fallback symbol remaps for ARM/RISC-V in ggml-cpu/arch-fallback.h
• fixed TBQ flash-attn AUTO validation to use merged GQA row widths in llama-context.cpp
• removed the extra temp-buffer path for quantized -> F32 dup in ggml-cpu/ops.cpp
• trimmed the redundant CPY test expansion while keeping the new quantized coverage in test-backend-ops.cpp

Re-ran:

• cmake --build build-cpu-pr --target test-quantize-fns test-backend-ops llama-cli -j4
• ./build-cpu-pr/bin/test-quantize-fns
• ./build-cpu-pr/bin/test-backend-ops test -b CPU -o GET_ROWS,SET_ROWS,CPY,FLASH_ATTN_EXT -p 'tbq'
• llama-bench smoke run with tbq4_0/tbq4_0

Hi @elusznik — great work on this PR. I've been running the turboquant-cpu-tbq-pr branch on an Apple Silicon M4 Mac mini (16GB unified memory) and implemented the NEON optimization listed in your follow-up roadmap. Sharing results and code here — I've also pushed a working branch at CuriosityQuantified/llama.cpp:neon-arm-optimization if you want to diff directly.

Build fix: arch-fallback.h missing ARM64 defines (blocks build on aarch64)

Before anything else — the branch does not build on ARM64 without this fix. The __aarch64__ section of ggml/src/ggml-cpu/arch-fallback.h is missing the two TBQ _generic symbol aliases, causing a linker failure:

// ggml/src/ggml-cpu/arch-fallback.h  (around line 81, inside #elif defined(__aarch64__))
#define ggml_vec_dot_tbq3_0_q8_K_generic ggml_vec_dot_tbq3_0_q8_K
#define ggml_vec_dot_tbq4_0_q8_K_generic ggml_vec_dot_tbq4_0_q8_K

NEON kernels

ggml/src/ggml-cpu/quants.c — ggml_vec_dot_tbq3_0_q8_K_generic and ggml_vec_dot_tbq4_0_q8_K_generic (identical pattern for both):

int j = 0;
#if defined(__ARM_NEON)
float32x4_t acc0 = vdupq_n_f32(0.0f);
float32x4_t acc1 = vdupq_n_f32(0.0f);
for (; j + 7 < QK_K; j += 8) {
    const float32x4_t tv0 = vld1q_f32(tmp + idx + j);
    const float32x4_t tv1 = vld1q_f32(tmp + idx + j + 4);
    const int8x8_t qi = vld1_s8(y[i].qs + j);
    const int16x8_t qi16 = vmovl_s8(qi);
    const float32x4_t qf0 = vcvtq_f32_s32(vmovl_s16(vget_low_s16(qi16)));
    const float32x4_t qf1 = vcvtq_f32_s32(vmovl_s16(vget_high_s16(qi16)));
    acc0 = vfmaq_f32(acc0, tv0, qf0);
    acc1 = vfmaq_f32(acc1, tv1, qf1);
}
sumf += d * vaddvq_f32(vaddq_f32(acc0, acc1));
#endif
// scalar tail
for (; j < QK_K; j++) {
    sumf += tmp[idx + j] * (d * y[i].qs[j]);
}
idx += QK_K;

ggml/src/ggml-turboq.c — matvec_row and matvec_t (add after existing AVX2 block):

#elif defined(__ARM_NEON)
float32x4_t acc0 = vdupq_n_f32(0.0f);
float32x4_t acc1 = vdupq_n_f32(0.0f);
for (; j + 7 < d; j += 8) {
    acc0 = vfmaq_f32(acc0, vld1q_f32(row + j),     vld1q_f32(x + j));
    acc1 = vfmaq_f32(acc1, vld1q_f32(row + j + 4), vld1q_f32(x + j + 4));
}
for (; j + 3 < d; j += 4) {
    acc0 = vfmaq_f32(acc0, vld1q_f32(row + j), vld1q_f32(x + j));
}
sum += vaddvq_f32(vaddq_f32(acc0, acc1));

Benchmark results — 8K context, Qwen3.5-4B-Q4_K_M, M4 Mac mini, 4 threads

Build: -DCMAKE_BUILD_TYPE=Release, Metal enabled, flash attention on, -nkvo 1 (KV on CPU, model weights on Metal GPU)

The gap to q4_0 narrows from ~50 t/s → ~16 t/s after NEON. The residual cost is the 128×128 Hadamard rotation matmuls (2 dense matmuls per 256-element block per TBQ block) — closing that fully would require a structured butterfly/WHT transform at the quantization algorithm level, not a kernel change.

Apple Silicon note

Running with -ngl 99 -ctk tbq4_0 (Metal KV offload) crashes — Metal backend does not support SET_ROWS for TBQ types. -nkvo 1 is the workaround (KV stays on CPU, model layers on Metal). Not a blocker for this PR scope, just worth noting for anyone testing on Apple Silicon.

Happy to submit a follow-up PR with these changes once this lands — or fold them in here if you prefer. Let me know what works best.

— @CuriosityQuantified

Hello @CuriosityQuantified, thanks for your input. Unfortunately I do not have ARM64 experience so I couldn't really do it myself. When it comes to the PR, I think a separate request would be more in line with the contribution guidelines of this project

@elusznik — understood, separate PR it is. I've opened it against your fork: elusznik#1

It's scoped to just the two changes from my comment above — the arch-fallback.h ARM64 fix and the NEON kernels. No other modifications. Should be easy to review and fold in whenever this lands upstream.

The numbers for tbq3_0 look great, so I hope this will be merged.

fwiw there are those of us (@TheTom, myself, and many others) that have been, and continue to, work on GPU optimized implementations. The CPU one is by far going to be the most underwhelming if you're comparing benchmarks.

Yes, now upstream does the rotation which is a lot of the work of TurboQuant (actually PolarQuant if you don't include QJL which kills speed but whatever). This was a good call out by @Mushoz! However, as @TheTom mentions above there are many other improvements that make getting this started worthwhile. Asymmetric KV compression is actually massive and works regardless of TurboQuant or not. In mine, Tom's and other people's extensive testing it is holding up very well - massive improvement in quantization, improvement in prefill and decode speed, with very little to even improved PPL and KLD. Truly a thing of beauty that @TheTom discovered there last week.

I very much respect @elusznik's intent to stick to guidelines and make this a minimal change in order to get the ball rolling. It's intimidating coming into such a popular repo with so much churn so respect and appreciation where it's due. With all due respect @ggerganov (and I have an immense amount), shitting on it and just calling it slop without any actual critique is just rude and makes you look like a jerk. I understand there's a lot of AI shit PRs around but this ain't one of them.

My two cents, is I think it is in everyone's best interest (the OS AI community) to help ensure that we don't have a ton of forks with disparate but material performance improvements littered about.

Hoping this lands (or something similar) and some of us can layer on discreet upstream PRs for some of the further decode, prefill and memory improvements we've identified.

It looks good, waiting for code to be merged hungry.

No news so far? I maybe expect too much but seeing only %10 improvement made me sad. Although it looks like cpu only. I want to use that for 6800xt gpu. Currently using gemma-4-26B-A4B-it-UD-IQ4_XS.gguf and just want to get best result from that. And indeed local LLMs getting better each day which is great news.

@ekryski thanks for the endorsement and the kind words, the slop comment made me feel kinda bad after putting in a couple of days' serious work into this

Instead of looking at new quants, you can take a look at existing quants for kv cache here: #21551

Asymmetric K/V is critical for some models
This is the biggest practical finding from my testing. Symmetric turbo (same type for K and V) is catastrophic on certain model families with Q4_K_M weights:

While looking at existing quants, I found an outlier kv quant pair that seems to perform better than quant neighbors. q3_K for K and q2_K for V.

The K3V2 / softmax-error-floor finding in #21591 (NexusQuant, by @jagmarques) quantifies the K-side bitcount question that's open here, and complements @TheTom's q8_0-K asymmetric data with a different recovery path on the architecture-sensitive cases.

Diminishing returns on K bitcount (NexusQuant, Mistral-7B):

Going from 3-bit to 4-bit keys: +0.06pp PPL. Going from 4 to 5: rounding noise. The reported mechanism: K feeds softmax, so K-quantization noise propagates across all positions, but the softmax error floor is hit at ~3 bits; further K precision buys nothing. V is linearly combined so V noise stays proportional, and lower V bitcount remains cheap.

Architecture-dependent layer-boundary breakage:

NexusQuant reports the Qwen2.5-7B catastrophic break with symmetric KV quant recovers by protecting the first and last 2 transformer layers at fp16, leaving the rest of K at low bitcount. This is layer-position protection, distinct from @TheTom's token-position boundary-V in layer-aware-v-compression.md. Mistral and Phi-3 don't need the protection per their data.

Two implications for the asymmetric K/V configs this PR is being asked to support:

1. K8 is likely over-allocated. TheTom's q8_0-K / turbo3-V recovers Qwen2.5-7B from PPL 3,556 to 6.71, but per NexusQuant's K-bitcount curve the K side could drop to 3 bits at near-zero PPL cost and ~2.7x further K-side compression. Worth measuring asymmetric tbq3_0-K / tbq3_0-V (or q3_K / tbq3_0) against the K8 asymmetric baseline.

2. Layer-boundary fp16 is a complementary lever. On architectures where bulk-layer K3 still breaks, protecting layers 0-1 and N-1, N at fp16 may recover quality without raising bulk K bitcount across the rest of the stack. Cheaper at long context than q8_0 across all layers.

Refs:

• Finding: asymmetric K/V quantization (3b-K/2b-V) gives 6x better quality than symmetric — architecture dependent #21591 (NexusQuant findings, closed)
• Repo + numbers: https://github.com/jagmarques/nexusquant

On Qwen2.5-7B-Instruct, per-head fp16 masking of the lowest-2% KV-heads matches K3V2 boundary-protect on retrieval at 2.28 vs 3.93 K-bits/element. Not a quant type, so orthogonal to this PR, but a relevant knob for the K-bit/arch-sensitivity question.

So because the CPU-only realisation brings small improvement - the CUDA realisation, which brings better improvement is blocked by design, because CUDA-realisation is only allowed after CPU one is passed?

No. Its not a question of speed.

No, the reason TurboQuant quants are not being added is because:
-> nobody has yet proven that they're better than quants of similar bpw in terms of quality (KLD)
-> the main novelty with Hadamard rotations has already been merged into the core llama.cpp quants
-> there's otherwise no reason to add new quant types if there's no good justification for them

No, the reason TurboQuant quants are not being added is because: -> nobody has yet proven that they're better than quants of similar bpw in terms of quality (KLD) -> the main novelty with Hadamard rotations has already been merged into the core llama.cpp quants -> there's otherwise no reason to add new quant types if there's no good justification for them

A couple of specifics on @pwilkin's three points, since most of the data is already in this PR and the surrounding thread.

On KLD at similar bpw: the table in this PR body answers it directly. On Qwen3.5-4B-Q4_K_M:

• q4_0: PPL 13.8400, KLD 0.00912, 3.56x compression
• tbq4_0: PPL 13.8323 (lower than q4_0), KLD 0.00960 (within noise), 3.94x compression

So tbq4_0 lands at slightly better PPL than q4_0 AND higher compression AND comparable KLD, in the PR's own CPU sweep. For an independent third-party reproduction at equivalent 4-bit V budget, @unamedkr / quantumaikr ported the algorithm to single-header C in https://github.com/quantumaikr/quant.cpp and measured Llama 3.2 3B turbo_kv_4b at PPL 14.28 (+5.3% vs FP32) vs llama.cpp q4_0 KV at PPL ~14.99 (+10.6%) on the same model. Same nominal 4-bit V cache budget, TurboQuant lower PPL.

Cross-model data from my own testing (50+ model/hardware combos, asymmetric K/V configs that survive aggressive weight quantization) is in my earlier comment and detailed at https://github.com/TheTom/turboquant_plus/blob/main/docs/papers/asymmetric-kv-compression.md. KLD per-bit budget across MoE and dense variants: https://github.com/TheTom/turboquant_plus/blob/main/docs/kl-divergence-results.md.

On the rotation in upstream: PR #21038 implements rotation only. Quoting the PR body: "The implementation does not introduce new types and is compatible with all existing quantizations" and "I don't know what is the impact of the remaining techniques explained in TurboQuant (PolarQuant, QJL, etc.). They could be important and can potentially improve further on top of this."

That's the design space TBQ in this PR and TurboQuant+ in mine occupy. Per Zandieh et al. (ICLR 2026, arxiv 2504.19874), TurboQuant pairs WHT rotation with Lloyd-Max-optimal codebook quantization tuned for the post-rotation distribution. The codebook is not in upstream, and that codebook is what enables sub-4-bpw operation at the quality envelope in the PR table. Rotation alone wrapping q4_0 keeps q4_0's compression and quality profile, which is a different design point.

On intent: this is opt-in via the existing --cache-type-k/v knob, alongside q4_0, q8_0, etc. The target users are hardware-constrained deployments where the bpw reduction unlocks a model size or context length that wouldn't otherwise fit (consumer GPUs at long context, larger MoE on laptop-class VRAM, mobile/edge inference). Default behavior unchanged for anyone who doesn't opt in.

On intent: this is opt-in via the existing --cache-type-k/v knob, alongside q4_0, q8_0, etc. The target users are hardware-constrained deployments where the bpw reduction unlocks a model size or context length that wouldn't otherwise fit (consumer GPUs at long context, larger MoE on laptop-class VRAM, mobile/edge inference). Default behavior unchanged for anyone who doesn't opt in.

Yes, I definitely need more such knobs for exactly the reason mentioned above (unlock more model/context combination)

Someone did a little study here https://anbeeld.com/articles/kv-cache-quantization-benchmarks-for-long-context

There is a lot of detail in there.

Tests on Qwen 3.6 27B show why TurboQuant is overrated but saved by TCQ, q5 deserves more attention, and symmetric q8 might be a waste of VRAM.

Someone did a little study here https://anbeeld.com/articles/kv-cache-quantization-benchmarks-for-long-context

There is a lot of detail in there.

Tests on Qwen 3.6 27B show why TurboQuant is overrated but saved by TCQ, q5 deserves more attention, and symmetric q8 might be a waste of VRAM.

I read that and tried q5_0 and q4_0 symmetric: token generation collapsed to 18-20 tok/s respectively (I made sure they both fit comfortably in VRAM), while turbo4 symmetric (using buun-llama.cpp) gets me 36 tok/s. With that speed difference, I can more than make up for perplexity by running a query again.

As pointed out by other maintainers and also myself: so far no one has presented evidence that would justify the opportunity cost of adding TurboQuant support. Anyways, closing this particular PR since it conflicts with procedure.

As pointed out by other maintainers and also myself: so far no one has presented evidence that would justify the opportunity cost of adding TurboQuant support. Anyways, closing this particular PR since it conflicts with procedure.

On consumer GPU with Tom's fork, me and thousands of othersa are able to run Qwen models locally with full context without a loss of accuracy. I could run LLMF on 10 year old GTX 1650 with 4 GB VRAM at 50 t/s at full context, something unimaginable. If you need actual numbers, it's not hard to run an experiment, but I would be happy to provide benchmarks if it helps.

If someone has a feature showing demonstrable gains with some quants compared to equivalent quantization types already in the main repo, then please open a new PR with those changes and with full benchmarks as indicated in the contributing guidelines. Sorry, but we cannot keep up with all the discussions that are based on vague unsubstantiated "feels" and claims without evidence.

For reference: this is a model analysis of the type I expect when adding a new quantization type.