merge Feature/turboquant kv cache#1
Merged
Conversation
Added nearest_centroid_4bit() with 16 optimal midpoints. C quantize now uses 4-bit indices + norm correction when TURBO4_USE_4BIT=1. C dequant uses CENTROIDS_4BIT[16] table + inverse rotation. Legacy 3-bit+QJL path preserved in #else for CUDA. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The ISWA (interleaved sliding window attention) build_attn overload was missing turbo Q rotation and V inverse rotation. K/V were WHT-rotated in SET_ROWS but Q was unrotated, causing rotated_K × unrotated_Q = garbage. Gemma 2 27B: PPL 13.7 trillion → 3.80 (turbo3), 3.79 (turbo4). q8_0 baseline: 3.75. turbo4 within 0.9% of q8_0. This also fixes any other ISWA model (Cohere2, OLMo2, Gemma3N, etc.). Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Ports the Metal turbo3 implementation to NVIDIA CUDA. End-to-end working
on RTX 5090 with Qwen3.5 35B A3B Q4_K_M: 3.47x KV compression at 32K
context, ~4x max context extension (256K → 1M tokens on 32GB VRAM).
New files:
- ggml-cuda/turbo-quant.cuh — block_turbo3_0 layout, WHT sign arrays,
3-bit centroid LUT, dequant helpers,
quantize kernel (set_rows path)
- ggml-cuda/turbo-wht.cu/.cuh — GGML_OP_TURBO_WHT CUDA kernel; 128-thread
blocks, in-place butterfly in shared memory,
forward + inverse WHT via compile-time template
- ggml-cuda/template-instances/fattn-vec-instance-turbo3_0-turbo3_0.cu
— VEC flash-attention instance for D=64/128/256
Modified files:
- dequantize.cuh — dequantize_turbo3_0 (produces float2 pairs)
- convert.cu — all 5 to-fp16/fp32 dispatchers
- fattn-common.cuh — vec_dot_fattn_vec_KQ_turbo3_0, dequantize_V_turbo3_0,
dispatcher extensions
- fattn-vec.cuh — turbo3 treated as unquantized (f16-style nthreads_KQ)
- fattn.cu — route turbo3 exclusively to VEC kernel; add dispatch macro
- set-rows.cu — k_set_rows_turbo3 kernel: per-128-elem group quantization
with WHT rotation and norm correction
- ggml-cuda.cu — supports_op + compute dispatch for TURBO_WHT + SET_ROWS
- llama-kv-cache.cpp — +2 tensor overhead for rotation matrices
Benchmark (RTX 5090, Qwen3.5 35B A3B Q4_K_M, FA on):
KV memory @32k: 702 MiB (f16) → 202 MiB (turbo3) = 3.47x compression
Max context: ~256K (f16) → ~1M (turbo3) = ~4x extension
Decode @short: 233 t/s (q8_0) → 190 t/s (turbo3) = 0.82x
Prefill @32k: 6335 t/s (q8_0) → 1215 t/s (turbo3) = 0.19x
Note: prefill speed degrades significantly vs Metal (Metal: 0.99x q8_0 at all
contexts; CUDA: 0.19x at 32K). Root cause: turbo3 currently uses the VEC
flash-attention kernel; q8_0 uses the more efficient TILE/MMA kernels at long
context. TILE/MMA support for turbo3 is the next milestone.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Remove the 5-line early-return that forced turbo3 onto the VEC flash attention kernel. The VEC kernel is still used for decode (Q->ne[1]==1) via the existing dispatch logic, but prefill now goes through the Turing MMA kernel (RTX 5090 is SM 12.0 >> 7.5). launch_fattn already pre-dequantizes K/V to FP16 when need_f16_K/V are set (which TILE/MMA always pass as true). Our ggml_get_to_fp16_cuda and ggml_get_to_fp16_nc_cuda dispatchers for TURBO3_0 — added in the original CUDA port commit — provide that conversion automatically. Stride recalculation (nb11 = nb11*bs*sizeof(half)/ts) also works out correctly for turbo3 (bs=32, ts=14): nb11*32*2/14 = ne[0]*sizeof(half). ✓ Before (VEC only): After (MMA for prefill): 2K prefill: 5032 t/s (0.73× q8_0) 6734 t/s (0.98× q8_0) 8K prefill: 3110 t/s (0.46× q8_0) 6613 t/s (0.98× q8_0) 32K prefill: 1215 t/s (0.19× q8_0) 6168 t/s (0.97× q8_0) Matches Metal M5 Max result (0.99× q8_0 flat across all context sizes). Decode unchanged (VEC, ~0.64-0.82× q8_0 depending on context depth). Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
…e (71→94 t/s, 98.5% of F16)
k_set_rows_turbo3 was the decode bottleneck: 1 thread/group serial kernel
gave 3.1% GPU utilisation (36.5 µs × 80 calls/token = ~21% of decode budget).
Replace with a fully parallel kernel — 1 block per 128-element group,
128 threads per block (one thread per element):
• Shared-memory WHT butterfly (7 stages, no atomics)
• Warp-reduce L2 norm + inter-warp accumulate via smem
• qs packed with __shfl_sync (4-wide gather), signs with __ballot_sync
• Reconstruction norm same pattern; one write per sub-block (warp lane 0)
Also tighten flash-attention dequant paths (fattn-common.cuh):
• vec_dot_fattn_vec_KQ_turbo3_0: elem0/elem1 always share the same
turbo3 block — load qs/signs once instead of twice per pair
• dequantize_V_turbo3_0: ne==4 fast path — load one qs byte and one
signs byte for all 4 elements; unrolled float2 / half2 output pairs
Benchmark (Qwen3.5 35B, RTX 5090, tg128):
Before: 71.86 t/s (0.75× q8_0)
After: 94.04 t/s (0.985× q8_0, within measurement noise of parity)
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
The outer loop in vec_dot_fattn_vec_KQ_turbo3_0 stepped k_KQ_0 by `nthreads` (8), but the Q register file is loaded in blocks of `nthreads*cpy_ne` (32) elements per thread — the same pattern used by the f16/bf16 VEC kernels. This caused thread t>0 to pair K element (16s + 2t) against Q element (64*(s/4) + 8t + 2*(s%4)), a complete index mismatch. Every generated token had garbage attention scores. Fix: match the f16 kernel pattern — step by nthreads*cpy_ne, add an inner k_KQ_1 loop over cpy_ne pairs, and index Q_v as Q_v[k_KQ_0/nthreads + k_KQ_1]. Also clean up stale "PPL 23.5 vs 6.19" TODO comments in llama-graph.cpp that documented the symptom of this bug. Tested on RTX 5090, Qwen3.5-35B-A3B-Q4_K_M: - PPL (wikitext-2): 6.2023 → 6.2996 (+1.57%, within 5% target) - NIAH: 11/11 at 4K–256K (matches q8_0; was 0/11 before fix) Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
TheTom#13) Models with n_embd_head_k not divisible by QK_TURBO3_GROUP (128) — e.g. GLM-4.7 Flash / DeepSeek2 MLA with head_dim=576 — previously hard-crashed with GGML_ASSERT(ne00 % QK_TURBO3_GROUP == 0) in set-rows.cu. Fix: - supports_op for SET_ROWS: return false when turbo3 and ne00 % 128 != 0, so llama.cpp falls back to an unquantised KV path instead of asserting - get_best_fattn_kernel: return BEST_FATTN_KERNEL_NONE for turbo3 when K head dim is not 128-aligned (VEC kernel only instantiated for D∈{64,128,256}) Affected models can now load with -ctk turbo3 -ctv turbo3; the non-aligned heads silently use f16 KV while 128-aligned heads continue to use turbo3. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
…heTom#13) Models with n_embd_head_k not divisible by 128 (e.g. GLM-4.7 Flash, DeepSeek2 MLA with head_dim=192/576) previously crashed with GGML_ASSERT in set-rows.cu or segfaulted in CPU flash attention. Changes: - llama-kv-cache.cpp: detect incompatible head_dim at KV cache init, log a warning, and auto-fall back to q8_0 instead of crashing - ggml-cuda.cu supports_op: return false for turbo3 SET_ROWS when ne00 % 128 != 0, and for TURBO_WHT when ne[0] % 128 != 0 - fattn.cu: return BEST_FATTN_KERNEL_NONE for turbo3 at non-128 D - ggml-cpu.c: add turbo3 entry to type_traits_cpu (vec_dot + from_float) so CPU flash attention can handle turbo3 K/V for models where CUDA FA returns NONE (e.g. D=192 which has no CUDA FA kernel for any type) - set-rows.cu: add tail kernel for non-128 remainder (future use) - turbo-wht.cu: head-dim-aware processing with tail pass-through - ops.cpp: CPU WHT head-dim-aware with tail identity copy Tested: DeepSeek-Coder-V2 (head_dim=192) now loads with -ctk turbo3 and auto-falls back to q8_0 with clear warning. Qwen3.5 (head_dim=128) continues to use turbo3 at full speed (223 t/s). Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Models with head_dim divisible by 64 but not 128 (e.g. DeepSeek2 MLA
head_dim=192, GLM-4.7 Flash head_dim=576) now use 64-element WHT
groups instead of crashing or falling back to q8_0.
Key changes:
**64-element WHT groups:**
- turbo-quant.cuh: 64-element sign arrays + FWHT-64 + rotate_64
- set-rows.cu: templated on GROUP_SIZE {128,64}, reads group_size
from SET_ROWS op_params (set by llama-kv-cache based on head_dim)
- turbo-wht.cu: templated on group_size, head-dim-aware dispatch
- ggml.h/ggml.c: ggml_turbo_wht() now takes explicit group_size
parameter (0=auto) to handle MLA where output dim differs from
K head dim
- ops.cpp: CPU WHT parameterized on group_size
**MLA fix — missing Q rotation in K-only build_attn:**
- build_attn(inp_attn_k, ...) had no turbo Q pre-rotation, while
SET_ROWS was applying WHT to K. Result: <unrotated_Q, rotated_K>
= garbage. Now all three build_attn variants apply Q rotation.
- Inverse WHT moved inside build_attn_mha (before v_mla projection)
and also added to the non-FA attention path.
**Guards + fallback:**
- supports_op, fattn.cu, llama-graph.cpp: relaxed from %128 to %64
- llama-kv-cache.cpp: falls back to q8_0 only when head_dim%64!=0
- CPU type_traits_cpu turbo3 entry for CPU FA vec_dot
Tested: GLM-4.7 Flash (head_dim=576) 208 t/s, DeepSeek-V2 (192) 143 t/s,
Qwen3.5 (128) 223 t/s — all producing correct output with turbo3.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…e versa) Users can now mix turbo3 and q8_0 independently for K and V caches: -ctk turbo3 -ctv q8_0 # compress K more, keep V at higher quality -ctk q8_0 -ctv turbo3 # opposite tradeoff Changes: - New VEC FA template instances for turbo3/q8_0 cross-type pairs (fattn-vec-instance-turbo3_0-q8_0.cu, fattn-vec-instance-q8_0-turbo3_0.cu) - fattn-vec.cuh: extern declarations for mixed instances - fattn.cu: dispatch + allow turbo3/q8_0 mix through the K!=V type guard - llama-graph.cpp: inverse WHT guard now checks V type (not K type) — when K=turbo3 but V=q8_0, V values are NOT WHT-rotated so inverse WHT must not fire. For MLA, V is a view of K so v->type correctly reflects K's turbo type. Tested: - 16/16 coherence tests pass (4 models × 4 KV combos) - NIAH at 4K+32K: all combos within baseline variance (10-11/11) q8_0/q8_0: 10/11, q8_0/turbo3: 10/11, turbo3/turbo3: 10/11, turbo3/q8_0: 9/11 (1 extra miss, normal variance) Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
quantize_row_turbo3_0_ref was a stub that only stored the L2 norm and zeroed out qs/signs — any CPU-only path (-ngl 0) silently produced garbage output. Now implements full PolarQuant: L2 norm → normalize → forward WHT rotation (signs1 → butterfly → signs2) → 3-bit centroid quantize → pack qs/signs → corrected norm (grp_norm / recon_norm). WHT group size (128 or 64) is communicated from the CPU SET_ROWS handler via a global variable, read from the op_params that llama-kv-cache.cpp sets based on head_dim. This handles models with different K and V head dims (e.g. DeepSeek2 K=192→64, V=128→128). Tested CPU-only (-ngl 0): - Mixtral 8x7B (head_dim=128): correct output at 4.5 t/s - DeepSeek-Coder-V2 (head_dim=192): correct output at 18.2 t/s - GPU paths unchanged (Qwen 203 t/s, DeepSeek 164 t/s, GLM 202 t/s) Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Adds turbo2: 2-bit PolarQuant with WHT rotation, 2.5 bits/value, 6.4x compression vs f16. Same architecture as turbo3 but with a 4-centroid Lloyd-Max codebook instead of 8. Block format: block_turbo2_0 = 10 bytes per 32 values - norm (fp16, 2B): corrected L2 norm (grp_norm / recon_norm) - qs (8B): 2-bit centroid indices, 4 per byte Usage: -ctk turbo2 -ctv turbo2 Full stack: - ggml.h/ggml-common.h: GGML_TYPE_TURBO2_0 enum + block_turbo2_0 struct - ggml-turbo-quant.c: CPU quantize (with WHT) + dequantize - turbo-quant.cuh: CUDA centroids, midpoints, nearest-centroid, dequant - set-rows.cu: k_set_rows_turbo2 kernel (GROUP_SIZE-templated, parallel WHT) - dequantize.cuh + convert.cu: turbo2 to f16/f32 conversion - fattn-common.cuh: vec_dot_KQ_turbo2 + dequantize_V_turbo2 - fattn-vec.cuh + fattn.cu: VEC kernel dispatch + template instances - Mixed types: turbo2/q8_0 cross-type FA instances - common/arg.cpp: CLI --cache-type-k turbo2 - llama-graph.cpp + llama-kv-cache.cpp: graph + cache integration NIAH (Qwen3.5 35B, RTX 5090, 4K-512K): 4K-16K: 11/11 | 32K: 9/11 | 64K: 10/11 | 128K-512K: 11/11 Coherence: 31/31 across 4 models x all KV combos (GPU + CPU) Speed: 228 t/s decode on Qwen3.5 (comparable to turbo3 223 t/s) Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The mixed KV types commit changed the inverse WHT guard from checking k->type to v->type. For MLA, V is a view of K's first 512 elements (v->ne[0]=512, 512%128=0 → group=128). But K was quantized with 64-group WHT (k->ne[0]=576, 576%128≠0 → group=64). The mismatch produced garbage output on GLM-4.7 Flash. Fix: derive group_size from K when K is turbo (since K determines the rotation used during quantization), from V only when K is not turbo. Tested: 16/16 coherence (4 models × 4 KV combos), including GLM-4.7. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
InnerQ equalizes K channel variances before WHT rotation to reduce quantization error on models with anisotropic K distributions. Enabled via TURBO_INNERQ=N env var (N = calibration token count). Pipeline: calibrate per-channel K² stats → compute scales → apply scale to K before WHT (SET_ROWS) → apply 1/scale to Q and V output via ggml_turbo_wht op (scale passed as src[1]). Math: <Q/s, s*K> = <Q, K> preserves dot products. Key design: - d_innerq_active starts at 0 (CUDA zero-init) — kernel never multiplies by uninitialized scales - Auto-disables for 64-group models (group_size < 128) where channel mapping is not 1:1 across WHT groups - Auto-disables when max channel ratio < 1.2 (already balanced) - Cross-TU coordination via turbo-innerq.cu/cuh for scale_inv tensor updates between SET_ROWS and graph-side WHT - Scale_inv tensor stored in KV cache, initialized to identity Also: turbo2 (2-bit) now requires head_dim divisible by 128. At 64-group WHT, the 4-centroid codebook doesn't have enough resolution — DeepSeek (head_dim=192) produced garbage. turbo2 auto-falls back to q8_0 for non-128-aligned heads with a warning. Tested: 18/18 coherence (4 models × 4+ KV combos), InnerQ active on Qwen (222 t/s) and Mixtral (146 t/s), auto-disabled on GLM/DeepSeek. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
In the VEC flash-attention kernel, skip V dequantization for KV positions where ALL attention weights are below 1e-6. At long context most positions contribute noise, not signal — skipping their V dequant saves compute with zero quality impact. Both half2 and f32 paths are covered. The threshold (1e-6) matches the existing SOFTMAX_FTZ_THRESHOLD behavior where exp() values below this are flushed to zero. Benefit scales with context length — at short context nearly every position contributes so no skip occurs. At 512K+ most positions are skipped during generation. Tested: 10/11 NIAH at 512K (matches pre-sparse-V), 3/3 coherence. Generation at 512K: 25 t/s, prefill: 2535 t/s. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
PPL benchmarks revealed that 64-group WHT (for non-128-aligned heads) causes catastrophic quality loss on some models: - DeepSeek (head_dim=192): PPL 344,304 vs 9.9 baseline - GLM-4.7 (head_dim=576): PPL 22.79 vs 14.97 baseline (+52%) The 64-group WHT passed NIAH/coherence but the weaker decorrelation (6 stages, 3 groups per 192-dim head) doesn't preserve statistical quality. Straight PolarQuant (no WHT) is even worse. All turbo types now require head_dim divisible by 128. Models with non-128-aligned heads (DeepSeek2, GLM-4.7 Flash) auto-fall back to q8_0 with a warning. The 64-group WHT code remains for future investigation but is not used. PPL summary (Qwen3.5, head_dim=128, wikitext-2): f16: 6.20 | q8_0: 6.18 | turbo3: 6.31 (+2.2%) | turbo2: 6.69 (+8.3%) Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
4 centroids, no signs byte. Simpler than turbo3. PPL 6.507 (+6.48% vs q8_0) on Qwen3.5-35B-A3B. Metal shader: quantize, dequant (vec + non-vec), SET_ROWS kernel, FA template instantiations (non-vec nl=2, vec nl=8). Runtime: supports_op, vec/non-vec gating, llama-bench parser. Codex-reviewed: no bugs found, only known limitation is non-128 head_dim tail handling (same as turbo3, covered by q8_0 fallback). Closes TheTom#19 Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> Co-Authored-By: tturney@psyguard.ai
- Add mixed-type kernel instantiations for all turbo pairs (turbo2/3/4)
- Both vec and non-vec FA paths, all supported head dims
- Pipeline naming uses k/v prefix format (kturbo3_vturbo4) to avoid
underscore ambiguity in type names
- Relax symmetric K/V assertion to allow turbo mixed pairs
- Update pipeline name construction in both FA pipeline getters
- 90 new kernel instantiations (54 non-vec + 36 vec)
- All 335 FA kernels use unified k{type}_v{type} naming scheme
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Co-Authored-By: tturney@psyguard.ai
Add full asymmetric K/V quantization support for Metal flash attention:
- Pipeline naming uses k{type}_v{type} format for all FA kernels (335 total),
eliminating underscore ambiguity in type names
- 90 turbo × turbo asymmetric instantiations (turbo2/3/4 all combinations)
- 150 q8_0 × turbo asymmetric instantiations (both directions, all head dims)
- Gatekeeper and assertion updated to allow turbo × turbo and q8_0 × turbo pairs
- Zero regression on existing symmetric paths (validated across 4 models, 2 machines)
The q8_0 × turbo kernels fix a silent dispatch failure where mixed q8_0-K + turbo-V
configs would NaN (turbo4-V) or fall to undefined paths (turbo3-V). This enables
the asymmetric quality rescue: q8_0-K + turbo-V recovers near-baseline PPL on
low-bit models where symmetric turbo-K degrades.
Validated on Metal (M2 Pro + M5 Max):
- phi-4-Q8_0: symmetric turbo3 +4.2%, turbo4 +1.7% (no regression)
- Qwen2.5-7B Q4_K_M: q8_0-K + turbo4-V +1.0%, q8_0-K + turbo3-V +2.0% (rescued)
- Qwen3.5-35B MoE, 27B Dense, Mistral-24B: all healthy (no regression)
- Cross-hardware M2/M5 parity confirmed on all tested configs
Closes TheTom#27
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Co-Authored-By: tturney@psyguard.ai
…back)
Instead of falling back to q8_0 for models with non-128-aligned heads,
pad each head to the next multiple of 128 before WHT rotation. The
padded elements are zero and don't affect dot products since WHT
preserves inner products: <WHT(Q_pad), WHT(K_pad)> = <Q, K>.
This keeps turbo compression active on DeepSeek2, GLM-4.7 Flash, and
other non-128 head_dim models. Padding overhead is wasted bits on the
zero-padded quantized elements:
head_dim=192 → 256 (33% overhead, still 3.5x compression)
head_dim=576 → 640 (11% overhead, still 4.1x compression)
Changes:
- llama-kv-cache.cpp: allocate K/V cache at padded dimension, pad
k_cur/v_cur per-head via ggml_pad before set_rows, return padded
views from get_k/get_v
- llama-graph.cpp: pad Q per-head before turbo_wht, extract original
V head_dim from attention output after inverse WHT. All three
build_attn variants (KV, K-only/MLA, ISWA) updated.
PPL (wikitext-2, 512 context, 8 chunks):
DeepSeek (192→256): 11.61 vs 9.90 baseline (+17%)
Previously: 344,304 with 64-group WHT (catastrophic)
GLM-4.7 (576→640): 9.11 vs 14.97 baseline
Qwen3.5 (128): 6.31 unchanged
Tested: 16/16 coherence, NIAH 9-10/11 at 32K, CPU-only works.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
GLM-4.7 Flash with turbo3 KV cache pads head_dim 576→640 for WHT.
D=640 had no CUDA FA kernel, falling back to CPU FA at 37 t/s.
Added MMA flash attention configs for D=640 (DKQ=640, DV=512),
using identical tile sizes as D=576 (nbatch_K2=288, nbatch_V2=256).
ncols1 capped at 2 for D=640: at ncols1≥4 the Q shared memory
(ncols × (DKQ/2+4) × 4 = 83KB at ncols=64) exceeds the per-block
limit. ncols1≤2 keeps total shared memory under 80KB.
Changes:
- fattn-mma-f16.cuh: D=640 config entries + extern declarations
(ncols2=16 only, ncols1∈{1,2,4})
- fattn.cu: case 640 in kernel selection + simplified MMA dispatch
(always ncols2=16, ncols1 capped at 2)
- fattn-tile.cuh/cu: D=640 tile configs + dispatch
- 3 MMA template instances + 1 tile instance
Tested:
GLM-4.7 Flash: 192 t/s (was 37), PPL 17.05 (vs 14.97 f16)
Qwen3.5: 222 t/s, PPL 6.31 (unchanged)
DeepSeek: 143 t/s, PPL 11.61 (unchanged)
16/16 coherence, NIAH 9-10/11 at 32K
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…ug 2) Mixed turbo3-K/turbo2-V and turbo2-K/turbo3-V had no CUDA FA kernel instances, causing ~11x prefill regression (falling back to CPU FA). Added VEC template instances for both cross-type pairs at D=64/128/256. Updated the mixed-type guard in get_best_fattn_kernel to allow any combination of turbo2, turbo3, and q8_0. Tested: turbo3/turbo2 and turbo2/turbo3 both run at full CUDA VEC speed (~170 t/s prefill, ~221 t/s decode on Qwen3.5 35B). Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…Tom#25 bug 1 Ports GGML_TYPE_TURBO4_0 to CUDA using the 4-bit PolarQuant format (16 centroids, nibble-packed, no QJL). Previously turbo4 crashed on CUDA with "cannot run the operation (SET_ROWS)". Changes TURBO4_USE_4BIT default from Metal-only to all backends. The 4-bit format (16 centroids) has better quality than the legacy 3-bit+QJL format and is simpler to implement (no residual projection). Full CUDA stack: - turbo-quant.cuh: 4-bit centroids, midpoints, nearest-centroid, dequant element, per-block quantize - set-rows.cu: k_set_rows_turbo4 kernel (128 threads, WHT rotation, 4-bit quantize, nibble pack via warp shuffle, corrected norm) - dequantize.cuh + convert.cu: turbo4 to f16/f32 - fattn-common.cuh: vec_dot_KQ_turbo4 + dequantize_V_turbo4 - fattn-vec.cuh + fattn.cu: VEC dispatch + all cross-type instances (turbo4×turbo4, turbo4×q8_0, turbo4×turbo3, turbo4×turbo2) - ggml-cpu.c: CPU FA vec_dot for turbo4 PPL (Qwen3.5, wikitext-2): 6.23 (+0.8% vs q8_0) at 3.8× compression Speed: 217 t/s decode (comparable to turbo3 222 t/s) Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
… (issue TheTom#28) The block-size divisibility check in llama-context.cpp rejected turbo4 on GLM-4.7 Flash (head_dim=576, QK_TURBO4=128, 576%128≠0) before the KV cache zero-padding code could run. Fix: for turbo types, compute the padded head_dim (ceil to 128) before the divisibility check, matching what llama-kv-cache.cpp actually does. Tested: GLM-4.7 Flash turbo4 loads and runs at 193 t/s. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Add V-only layer-adaptive modes to TURBO_LAYER_ADAPTIVE env var:
- Mode 5: first2+last2 V=turbo4, rest V=turbo2
- Mode 6: last8 V=turbo4, rest V=turbo2
- Mode 7 (recommended): first2+last2 V=q8_0, rest V=turbo2
Mode 7 ("Boundary V") protects quality-sensitive boundary layers with
q8_0-V while aggressively compressing middle layers with turbo2-V.
K cache is unchanged (stays at whatever -ctk specifies).
Validated on Metal (M5 Max) across 4 models, 2 context lengths:
- phi-4-Q8_0: 4.784 PPL (vs turbo2 4.835, turbo3 4.742)
- Qwen2.5-7B Q4_K_M: 6.835 (vs turbo2 6.911, turbo3 6.707)
- Qwen3.5-35B MoE: 5.148 (vs turbo2 5.257, turbo3 5.137)
- Qwen3.5-27B Dense: 6.423 (vs turbo2 6.534, turbo3 6.273)
- 8K context: stable, no collapse
- NIAH retrieval: pass
- Speed: no penalty
Effective V compression between turbo2 and turbo3, closer to turbo2
on deeper models. Quality consistently better than uniform turbo2-V.
Usage: TURBO_LAYER_ADAPTIVE=7 llama-server -m model.gguf -ctk q8_0 -ctv turbo2 -fa 1
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Co-Authored-By: tturney@psyguard.ai
…follow-up) state_write_data and state_read_data used hparams.n_embd_k_gqa (576) for ggml_row_size, but turbo types zero-pad to 640. For turbo4 (QK=128), 576 % 128 != 0 → ggml_row_size assertion failure during prompt cache save on llama-server slot reuse. Fix: use k->ne[0] / v->ne[0] (actual padded tensor width) instead of hparams values in all four serialization paths (K write, K read, V write, V read). Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…rnels Port TheTom's warp-cooperative turbo3 SET_ROWS kernel and turbo2/turbo3 flash attention templates to HIP/ROCm (7900 XTX, gfx1100). HIP vendor header fixes: - Add cudaMemcpyToSymbol/FromSymbol -> hipMemcpyToSymbol/FromSymbol - Add cudaMemcpyHostToDevice/DeviceToHost mappings - Fix __shfl_sync, __shfl_xor_sync, __shfl_up_sync, __shfl_down_sync to support both 3-arg and 4-arg calls (CUDA allows defaulting width to warpSize, HIP macros required 4 args) - Add __ballot_sync -> __ballot with uint32_t cast (HIP returns 64-bit on wave64 platforms, turbo code expects 32-bit) HIP CMakeLists: - Add turbo3 and turbo2 flash attention template instances (same files as CUDA CMakeLists, were missing from HIP build) Tested: Mistral-Small-24B turbo3 PPL = 5.28 (+2.4% vs F16 baseline 5.16) Previously showed catastrophic PPL ~15000 due to CPU quantize stub bug (fixed by TheTom in 53f1298). Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
One norm per rotation group instead of four identical copies.
Eliminates 6 bytes of redundant storage per 128-element group.
turbo3: 3.50 -> 3.125 bits/value, 4.57x -> 5.12x compression
turbo2: 2.50 -> 2.125 bits/value, 6.4x -> 7.53x compression
Zero PPL regression validated across:
- Asymmetric q8_0-K + turbo{2,3}-V
- Symmetric turbo3/turbo3
- Boundary V (LA-V7)
- 3 architectures (dense, Qwen, MoE)
- 3 context lengths (512, 8K, 32K)
- 2 Apple Silicon platforms (M5 Max, M2 Pro)
- NIAH 3/3 pass
+3-7% decode on tested M2 Pro setup. No regression on M5.
Also adds derived NL_TURBO3/NL_TURBO2 macros replacing ~250
hardcoded FA template nl values. Block size is now a one-line
edit in ggml-common.h.
Credit to @AmesianX whose block_size=256 CUDA implementation
prompted this investigation.
Co-Authored-By: tturney@psyguard.ai
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
The block_size=128 change (adac2c6) broke CUDA quantization: with QK=128, blocks_per_group=1, but the warp-cooperative packing still used blk_base+warp_id, causing warps 1-3 to write OOB. Fix: compute elem_in_block = j % QK_TURBO_N and use it for block pointer (j / QK_TURBO_N) and byte offsets (elem_in_block / 4 for qs, elem_in_block / 8 for signs). Works for both QK=32 and QK=128. Validated on RTX 3090 (sm_86), llama3.1:8b Q4_K_M, q8_0/turbo3: PPL = 7.587 (matches QK=32 baseline exactly).
Sparse V: now enabled by default on all Metal (was M5+ only). Validated across 30+ testers with zero PPL impact. Opt-out: TURBO_SPARSE_V=0. Boundary V: auto-enabled (mode 7) when -ctv turbo2 is set. Protects first 2 + last 2 layers with q8_0-V, rest turbo2-V. 37-91% quality recovery across 4 tested models. Opt-out: TURBO_LAYER_ADAPTIVE=0. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
* nix: support unified apple-sdk * Impl roll op for Metal * Revert "nix: support unified apple-sdk" This reverts commit abfa473. * update ops.md * update op docs
* ggml: add graph_reused * use versioning instead of reuse flag * increment version with atomic * use top bits for split numbering * add assert * move counter to ggml.c * set uid in split_graph only * fix windows * address further review comments * get next_uid rather than doing bit manipulation * rename + add comment about uid
* fix NemotronH vocab loading by using trust_remote_code for unsupported config patterns * fix NemotronH tokenizer loading by overriding set_vocab with trust_remote_code
* support nvfp4 tensors for Gemma4 * add wo_s to build_attn * add wo_s to build_attn * fix glm4
…ers (ggml-org#21245) * model : refactor QKV into common build_qkv and create_tensor_qkv helpers * model : extend build_qkv to bert/mpt/dbrx/olmo/lfm2/nemotron-h/granite-hybrid/gemma3n-iswa/t5-dec and fix wqkv_s
…g#21962) (ggml-org#21980) * server: tests: fetch random media marker via /apply-template (ggml-org#21962 fix) * server: allow pinning media marker via LLAMA_MEDIA_MARKER env var get_media_marker() checks LLAMA_MEDIA_MARKER at first call and uses it as-is if set, falling back to the random marker otherwise. Tests no longer need to fetch the marker dynamically via /apply-template: the fixture sets LLAMA_MEDIA_MARKER=<__media__> so the hardcoded prompts work as before. Address review feedback from ngxson * server: make get_media_marker() thread-safe via magic statics Use a C++11 static local with a lambda initializer instead of a global static with an empty-check. The runtime guarantees initialization exactly once without explicit locking. Address review feedback from ggerganov * nits * nits
* model: using single llm_build per arch * fix merge * nits
* optimize hmx_mat_mul functions by calculating row and column tiles upfront * refactor core_dot_chunk_fp16 to use size_t for tile counts and improve readability * wip * set scale outside of loop * wip * refactor core_mma_chunk_fp16 and mat_mul_qk_0_d16a32 to use size_t for tile counts * wip * wip * refactor transfer_output_chunk_fp16_to_fp32 to use size_t for dimensions * refactor core_dot_chunk_fp16 to use size_t for tile row stride calculation * wip * refactor hmx_mat_mul functions to use hvx_vec_splat_f16 for column scales initialization * refactor hmx_mat_mul_permuted_w16a32_batched to streamline scale setting and locking * refactor core_dot_chunk_fp16 to improve tile stride calculations for output * refactor hmx_mat_mul functions to use Q6_V_vsplat_R for column scales initialization * fix compiling error * wip * optimize row and column tile indexing in core_mma_chunk_fp16 function * wip * Revert "wip" This reverts commit cde679e. * Add size limit check for HAP_mmap in htp_iface_mmap and drop_mmap functions * wip
…dreno (ggml-org#21938) * opencl: refactor q8_0 gemm/gemv Adreno dispatch * opencl: refactor q8_0 set_tensor * opencl: fix whitespace
* model : Gemma4 model type detection * model : Gemma4 model type detection
* cmake : allow libcommon to be shared * cmake : rename libcommon to libllama-common * cont : set -fPIC for httplib * cont : export all symbols * cont : fix build_info exports * libs : add libllama-common-base * log : add common_log_get_verbosity_thold()
* server: respect the ignore eos flag * ci: add android arm64 build and release * patch * pin android-setup actions to v4 * Apply suggestions from code review Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com> * lf in the suggestion --------- Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
* CUDA: use a ring-buffer for cuda graphs * bump limit to 128 * use LRU eviction * better naming * do periodic clean-up
…ng (ggml-org#21052) * Update workflows to remove dependence on llvmpipe * Try setting Dawn_DIR * remove c++20 initializers * Move to proper guid * Try avoiding segfaults on vulkan backend process exit * Remove compiler warnings on parameter casting * Fix soft_max and update reg_tile accumulation to f32 for better precision * Refactor flash_attn a bit * remove c++20 initializers and format * Increase div precision for NVIDIA * revert div precision and comment out ggml-ci node for now * Formatting * Try debugging on a failing CI node * Revert "Try debugging on a failing CI node" This reverts commit 1971e33.
…t-kv-cache # Conflicts: # ggml/src/ggml-vulkan/ggml-vulkan.cpp # ggml/src/ggml-vulkan/vulkan-shaders/copy_to_quant.comp # ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs_cm2.glsl # ggml/src/ggml-vulkan/vulkan-shaders/types.glsl # ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp # tests/test-backend-ops.cpp
richginsberg
pushed a commit
that referenced
this pull request
Apr 19, 2026
Codex post-commit review found: 1. TURBO_D was QK_TURBO3 (now 32) — broke turbo4 C array sizes 2. SET_ROWS kernel turbo3-specific but instantiated for turbo4 3. Tail block drop for non-128 head dims Fixed TheTom#3 (TURBO_D). #1 and TheTom#2 don't affect turbo3+dk128 path. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
richginsberg
pushed a commit
that referenced
this pull request
Apr 19, 2026
Complete experiment log: #1 4-mag LUT: 15.1 at 8K (BEST, +38%) TheTom#2 Batched extract: 13.7 (+25%) TheTom#3 Inline FA block: 13.5 (I-cache pressure) TheTom#4 Deferred norm: 12.9 (loses ILP) TheTom#5 2-pair half2: 12.0 (ternary overhead) TheTom#6 Select chain: 11.9 (branches kill) TheTom#7 Bit-arithmetic: 11.6 (ALU too heavy) TheTom#8 FMA branchless: 11.4 (ALU still too heavy) TheTom#9 Named-reg ternary: 10.3 (branches worst) TheTom#10 Main (8-LUT): 10.95 (baseline) TheTom#11 Non-vec FA: 10.2 (wrong kernel) Ceiling: 24.5 (no dequant) Apple8 hardware truth: 1 divergent constant read < 7 ALU ops (even with fma) Branches cost MORE than divergent constant reads Array indexing ALWAYS spills on Metal 4 constant addresses is the sweet spot The 4-mag LUT is the dequant-level ceiling on Apple Silicon. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> Co-Authored-By: tturney@psyguard.ai
zero4281
pushed a commit
to zero4281/llama-cpp-turboquant
that referenced
this pull request
May 23, 2026
Codex post-commit review found: 1. TURBO_D was QK_TURBO3 (now 32) — broke turbo4 C array sizes 2. SET_ROWS kernel turbo3-specific but instantiated for turbo4 3. Tail block drop for non-128 head dims Fixed TheTom#3 (TURBO_D). richginsberg#1 and TheTom#2 don't affect turbo3+dk128 path. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
zero4281
pushed a commit
to zero4281/llama-cpp-turboquant
that referenced
this pull request
May 23, 2026
Complete experiment log: richginsberg#1 4-mag LUT: 15.1 at 8K (BEST, +38%) TheTom#2 Batched extract: 13.7 (+25%) TheTom#3 Inline FA block: 13.5 (I-cache pressure) TheTom#4 Deferred norm: 12.9 (loses ILP) TheTom#5 2-pair half2: 12.0 (ternary overhead) TheTom#6 Select chain: 11.9 (branches kill) TheTom#7 Bit-arithmetic: 11.6 (ALU too heavy) TheTom#8 FMA branchless: 11.4 (ALU still too heavy) TheTom#9 Named-reg ternary: 10.3 (branches worst) TheTom#10 Main (8-LUT): 10.95 (baseline) TheTom#11 Non-vec FA: 10.2 (wrong kernel) Ceiling: 24.5 (no dequant) Apple8 hardware truth: 1 divergent constant read < 7 ALU ops (even with fma) Branches cost MORE than divergent constant reads Array indexing ALWAYS spills on Metal 4 constant addresses is the sweet spot The 4-mag LUT is the dequant-level ceiling on Apple Silicon. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> Co-Authored-By: tturney@psyguard.ai
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
20 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Overview
Merge TheTom's turboquant kv cache into master
Additional information
Requirements