vulkan: add TurboQuant KV cache support and optimized turbo mat-vec paths - #140
Open
Fenix46 wants to merge 216 commits into
Open
vulkan: add TurboQuant KV cache support and optimized turbo mat-vec paths#140Fenix46 wants to merge 216 commits into
Fenix46 wants to merge 216 commits into
Conversation
New types: GGML_TYPE_TURBO3_0 (3-bit) and GGML_TYPE_TURBO4_0 (4-bit) Implements PolarQuant + QJL compression per the ICLR 2026 paper. Block size = 128 (matching head_dim for optimal rotation Gaussianization) turbo3: 52 bytes per 128 values = 3.25 bits/value (4.9× vs fp16) turbo4: 68 bytes per 128 values = 4.25 bits/value (3.8× vs fp16) Status: - ✅ Type definitions in ggml.h - ✅ Block structures in ggml-common.h - ✅ Quantize/dequantize C implementation in ggml-turbo-quant.c - ✅ Registered in ggml.c type traits - ✅ Added to kv_cache_types in arg.cpp - ✅ Builds successfully - ✅ Shows in --help output - ❌ Metal SET_ROWS kernel not implemented (blocks GPU inference) - ❌ Needs Metal dequantize kernels for attention computation Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Added Metal shader implementations: - quantize_turbo3_0 / quantize_turbo4_0 (per-block quantization) - dequantize_turbo3_0 / dequantize_turbo4_0 (type4x4 and type4 variants) - kernel_set_rows_turbo template (128-element block size) - Flash attention instantiations for all dk/dv variants Added TURBO3_0/TURBO4_0 to Metal device SET_ROWS validation. Builds successfully. Testing with Qwen 3.5 35B-A3B MoE on M5 Max. Note: Initial version uses simplified quantization (no rotation matrix) for Metal compatibility. Full rotation requires custom kernel with extra buffer bindings — tracked for follow-up. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…g#21 Embedded pre-computed 128×128 rotation and QJL matrices (256KB constant memory) directly in the Metal shader. Both quantize and dequantize now perform the full TurboQuant algorithm: Quantize: normalize → rotate → codebook → inverse rotate → residual → QJL Dequantize: codebook → inverse rotate → QJL correction → rescale Previous version (no rotation) produced garbage. This should produce meaningful output since the rotation Gaussianizes the KV distribution. Note: dequantize does full 128-element rotation per chunk (8× work). Optimization possible with caching or restructured kernel in follow-up. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…ml-org#21 - Inlined turbo-matrices.h directly into ggml-metal.metal (256KB) to fix JIT compilation failure with #include - Added C round-trip test (test-turbo-quant.c): turbo3 cosine=0.906, turbo4 cosine=0.966 — matches Python prototype - Metal library loads successfully ("loaded in 5.9 sec") - Model runs on Metal but output quality needs debugging (Metal quantize/dequantize may have a bug vs the working C version) C round-trip PROVES the algorithm works in C. Metal shader needs debugging — likely an issue with the dequantize chunk addressing or the large constant arrays in thread-local memory. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…org#23 Codex review found: 1. Stale duplicate code in dequantize_turbo3_0_t4 (compile would fail) 2. thread static is risky/non-portable in MSL Fixed: removed thread static caching, using plain thread locals. Speed unchanged (2.4 tok/s) — the static caching wasn't actually working on Metal. True optimization needs architectural change in flash attention kernel to dequantize once per block, not per chunk. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…gml-org#26 Massive reduction in constant memory and compute: - 256KB of dense matrices → 512 bytes of sign arrays - O(d²) = 16,384 ops → O(d log d) = 896 ops per rotation - Metal shader file: 1.5MB → 432KB Speed: still 2.4 tok/s. WHT reduced per-rotation cost but the bottleneck is redundant calls (8-32× per block from flash attention). The dequantize function is called per 4/16-element chunk, each time doing the full 128-element WHT. Need to modify the flash attention kernel to dequantize once per block. Quality: WHT+signs gives BETTER quality than dense QR on real KV tensors (cosine 0.94 vs 0.79 at 2-bit). Sub-Gaussian distribution (kurtosis 1.53) means fewer outliers hitting extreme centroids. Reviewed by Codex: WHT butterfly correct, inverse order verified, QJL correction matches reference C implementation. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…gml-org#23 Root cause analysis: 8-32× redundant full-block dequantize per block from flash attention template. Four approaches documented with expected speedups and risk levels. Plan: D (reduce overhead) → A/B (eliminate redundant calls) Target: 2.4 tok/s → 20-40 tok/s Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…-org#23 Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…gml-org#23 No-op dequant test: even returning all zeros from dequantize, turbo3 runs at 2.4 tok/s (same as with full WHT rotation). The bottleneck is NOT in the attention dequantize path. New hypothesis: the SET_ROWS (quantize) path is the bottleneck. The Metal quantize_turbo3_0 function does 3 WHT rotations per KV write, totaling ~3200 ops per block × 224 blocks per token. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…rg#23 CRITICAL BUG: The #include "turbo-wht.h" caused Metal JIT compilation to fail at runtime. The model silently fell back to CPU for ALL ops. ALL previous benchmarks (2.4 tok/s) were measuring CPU, not Metal GPU. After inlining the header: - MoE gen: 2.4 → 10.7 tok/s (4.5× improvement, now actually on Metal) - MoE prompt: 4.2 → 60.9 tok/s (14.5× improvement) Remaining gap vs q8_0: 85 → 10.7 tok/s (8× slower, down from 35×) This is the SAME bug we hit with turbo-matrices.h earlier. Rule: NEVER use #include in ggml-metal.metal — always inline. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…org#23 Previous 2.4 tok/s was CPU fallback. Real Metal numbers: MoE: 10.7 tok/s gen (8× slower than q8_0, was thought to be 35×) Qwopus: 5.3 tok/s gen (3.3× slower than q8_0) Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…l-org#27 Full investigation log with all tests, results, and the root cause. Upstream TurboQuant activity tracked in ggml-org#27. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…-org#28 Key findings from Dejan.ai, unixsysdev, and mudler: 1. QJL naively added back destroys quality (cosine 0.69) 2. Pre-rotate queries eliminates rotation from dequant path 3. WHT abandoned by everyone — dense QR or no rotation preferred 4. unixsysdev gets -0.8% speed loss with fused CUDA kernel 5. We're the only Metal implementation Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…in) ggml-org#23 Removing WHT rotation from dequant (quality broken, speed test only): gen: 10.7 → 49.1 tok/s (4.6× improvement, 57% of q8_0) prompt: 67.3 → 162.6 tok/s Confirms pre-rotate-queries would deliver ~49 tok/s. Remaining gap (49 vs 85) is block size + QJL overhead. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Speed ceiling confirmed: stripping rotation from dequant gives 49.1 tok/s (vs 10.7 with rotation, vs 85.5 q8_0 baseline). Implementation plan: store rotation matrix in KV cache, apply to Q in graph builder, strip from Metal dequant. 6 files to modify. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…org#23 Instead of inverse-rotating every K during dequant, rotate Q once before attention. Math: <q, R^T*c[idx]> = <R*q, c[idx]>. Changes: - Store rotation matrix (R^T) in KV cache, filled after buffer clear - Apply ggml_mul_mat(R_T, q) in build_attn_mha after permute - Strip turbo_rotate_inverse from Metal dequant - Dynamic cast to access rotation from mctx Results: - MoE gen: 10.7 → 51.4 tok/s (4.8× speedup) - MoE prompt: 67.3 → 160.3 tok/s (2.4× speedup) - Now at 60% of q8_0 speed with 4.9× compression - Model produces coherent output Codex review: fixed buffer clear ordering (was zeroing rotation after init). Verified: rotation point is correct (after 4d reshape + permute, ne[0]=128). Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…gml-org#23 Full investigation log documenting every test, every dead end, and every breakthrough. 21× total improvement from CPU fallback to pre-rotate-queries. Key lessons: no #include in Metal, no-op testing, pre-rotate-queries, buffer clear ordering, codex+roast catch real bugs. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Validated on real Qwen3 KV tensors: cosine sim 0.9508 → 0.9831 (+3.2%) MSE-only better on 99.3% of vectors including p1 tails. 3-bit index split: lower 2 bits in qs[], upper 1 bit in signs[]. No QJL stage in quantize or dequant. Results: - MoE gen: 51.4 → 62.2 tok/s (73% of q8_0, was 60%) - MoE prompt: 160 → 200 tok/s (90% of q8_0) - Qwopus gen: 14.6 → 15.5 tok/s (88% of q8_0, was 83%) - Qwopus prompt: 67 → 83 tok/s (100% of q8_0!) Codex verified: bit packing correct, quantize/dequant consistent. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Speed ceiling without Q rotation: 61.3 tok/s (vs 62.2 with it). The 128×128 ggml_mul_mat adds <1% overhead on Metal. Remaining gap is structural (block size + dequant complexity). Final: MoE 62.2 tok/s (73%), Qwopus 15.5 tok/s (88%). Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Diagnostic benchmark proves the 26% gap is entirely from block size 128. q4_0 (block 32, 4-bit quantization) runs at 84.2 tok/s = identical to q8_0. Next: turbo3 with block size 32. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Changed QK_TURBO3 from 128 to 32 (storage block size). Rotation still operates on 128-element groups (QK_TURBO3_GROUP=128). SET_ROWS kernel processes 4 blocks per rotation group. Flash attention nl_k changed from 32 to 8 (matching q4_0). Block struct: 14 bytes per 32 values = 3.5 bits/val → 4.6× compression. Results: - MoE gen: 62.2 → 77.7 tok/s (91% of q8_0 at 85.5) - MoE prompt: 200 → 218.5 tok/s (98% of q8_0) - Qwopus gen: 15.5 → 17.0 tok/s (97% of q8_0 at 17.6) - Qwopus prompt: 83 → 89.5 tok/s (108% of q8_0 — FASTER) Target was 75+ tok/s. Exceeded. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
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 ggml-org#3 (TURBO_D). ggml-org#1 and ggml-org#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>
Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…l-org#30 Perplexity benchmarking reveals catastrophic quality failure: - f16: 6.121, q8_0: 6.111, q4_0: 6.142 - turbo3: 165.6 (27× worse) Speed benchmarks were meaningless — fast garbage. Root cause investigation needed before any quality claims. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
1. V cache returns rotated-space values (cosine=0.02 vs correct 0.987) 2. dynamic_cast to llama_kv_cache_context fails for MoE models (uses llama_memory_hybrid_context, not kv_cache_context) → Q rotation and V inverse rotation NEVER executed Fix: store rotation tensors in llm_graph_context, not KV cache. Or access through hybrid memory interface. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…gml-org#31 Block 128: PPL=165.6 (same as block 32) Disabled Q rotation: PPL=165.6 (same) Root cause: dynamic_cast fails for MoE hybrid memory context. Q rotation and V inverse rotation never execute. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…ml-org#31 ggml-org#30 ROOT CAUSE: pre-rotate-queries never executed because: 1. Q ne[0]=256 (GQA concatenated heads), rotation matrix ne[0]=128 2. mctx dynamic_cast failed for MoE hybrid memory FIX: put inverse WHT rotation back in dequantize_full_block. This is slower (10.7 tok/s vs 77.7) but produces CORRECT results. PERPLEXITY RESULTS: - f16: 6.121 - q8_0: 6.111 - q4_0: 6.142 - turbo3: 6.194 (+1.2% vs q8_0) ✅ The speed optimization (pre-rotate-queries) needs to be reimplemented to work with GQA head layout and hybrid memory types. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Quality confirmed: PPL 6.194 (+1.4% of q8_0) Speed: 10.7 tok/s (inverse rotation in dequant, no pre-rotate-queries) Previous speed claims (51-77 tok/s) were invalid — measured garbage output speed. Key lessons documented for future reference. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…disable=false) Previous state: db3595a added LLAMA_ATTN_ROT_K_OVERRIDE / _V_OVERRIDE per-side opt-in knobs but kept attn_rot_disable defaulting to TRUE for legacy LLAMA_ATTN_ROT_DISABLE compatibility. The override branches included `&& !attn_rot_disable` guards, so when LLAMA_ATTN_ROT_DISABLE is unset (default true) the per-side env knobs were silently no-ops. Users could not opt into rotation without also setting LLAMA_ATTN_ROT_DISABLE=0. Fix: flip attn_rot_disable default to false. Rotation is still OFF by default because attn_rot_k/v default to false. LLAMA_ATTN_ROT_DISABLE=1 still acts as a hard lock-out that blocks the per-side overrides for users who want a single switch to guarantee no rotation. Caught while running the cross-format KLD matrix for the rotation/PPL investigation paper — V-only override appeared to silently fail. Confirmed with logs that attn_rot_v stayed 0 even with LLAMA_ATTN_ROT_V_OVERRIDE=1 until this default flip.
Optimizations found via automated kernel optimization (33 experiments): - nthreads_KQ=1 + nthreads_V/=8 for better occupancy - Warp shuffle KQ scores (eliminates shared memory for reduction) - Precomputed scaled V centroids per block - __expf fast-math softmax - __launch_bounds__ occupancy 2 - Shmem KQ LUT: precompute Q×centroid in shared memory Also includes: - Auto-asymmetric KV: detect GQA ratio ≥6:1, upgrade K to q8_0 (fixes catastrophic PPL on Qwen2.5 symmetric turbo3) - HIP -Wnodiscard fix: (void) casts on cudaMemcpyToSymbol/FromSymbol
…nt (ggml-org#78) Post-attention V-padded reshape in build_attn was using hparams.n_head_kv(il), but cur returned from build_attn_mha has shape (n_embd_head * n_head, n_tokens) — n_head is the Q-head count. On GQA models where n_head != n_head_kv (e.g. Qwen2.5-0.5B with head_dim=64 padded → 128, n_head=14, n_head_kv=2), the reshape element count fails the assertion in ggml_reshape_3d and the process aborts. Symptom: GGML_ASSERT(ggml_nelements(a) == ne0*ne1*ne2) at ggml.c:3656. Reported and diagnosed by @bingh0 in TheTom#78. Verified locally on Qwen2.5-7B (head_dim=128, no padding, regression check passes) and on AMD MI300X with Qwen2.5-0.5B (head_dim=64, was crashing pre-fix). Three sites fixed (lines 2285, 2412, 2532 — same idiom in three build_attn overloads). Closes ggml-org#78. Likely also closes ggml-org#108 (speculative decoding hits the same assertion). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
fix: build_attn V-padded reshape uses Q-head count, not KV-head count (ggml-org#78)
Cherry-pick signalnine PR ggml-org#53: auto-asymmetric GQA + turbo VEC FA opts
…fault fix(kv-cache): per-side env-knob control for upstream attn rotation (default OFF)
Two well-meaning fixes both added `spirv-headers` to the package.nix
function pattern arglist (line 19 and line 22 on current HEAD), causing
a hard parse-time failure on any nix evaluation:
error: duplicate formal function argument 'spirv-headers'
at .devops/nix/package.nix:22:3:
21| shaderc,
22| spirv-headers,
| ^
23| useBlas ?
Drops the second occurrence. The remaining single declaration is what
the rest of the file actually references (line 19 binds the input;
`vulkanBuildInputs` and `nativeBuildInputs` consume it once each).
Reported by @cguentherTUChemnitz in ggml-org#81 (originally), then re-confirmed
by @alanscodelog on the current tip after two prior fix attempts both
landed the same line.
Verified: fixed file parses clean via nix-instantiate; injecting the
duplicate back reproduces the exact error message above.
Closes ggml-org#81.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
fix(nix): remove duplicate spirv-headers function arg (ggml-org#81)
Mirror of @apollosenvy's turbo3_0 Vulkan SET_ROWS port (PR ggml-org#33 + ggml-org#87) to the other two turbo types. Reported by @dpblnt in ggml-org#50 with a clean matrix on RX 9060 XT showing turbo3 V works on Vulkan but turbo2/turbo4 V abort with: pre-allocated tensor (cache_v_l*) in a buffer (Vulkan0) that cannot run the operation (SET_ROWS) at llama_context::sched_reserve() time, before any compute runs. Mechanical port across 4 files: - vulkan-shaders/types.glsl: block_turbo2_0 + block_turbo4_0 struct declarations matching the C side (ggml-common.h). - vulkan-shaders/copy_to_quant.comp: SET_ROWS quantize main() blocks for turbo2 (4 centroids, 2-bit pack, no signs byte) and turbo4 (16 centroids, 4-bit nibble pack, no signs byte). WHT setup and reduction structure identical to turbo3 (QK = 128 across all three). Centroid + midpoint tables ported from CENTROIDS_2BIT and CENTROIDS_4BIT in ggml-turbo-quant.c. - vulkan-shaders/vulkan-shaders-gen.cpp: turbo2_0 and turbo4_0 added to the set_rows iteration list at line ~789. - ggml-vulkan.cpp: SET_ROWS pipeline registrations + supports_op switch + dispatch element-count all extended with TURBO2_0 and TURBO4_0 cases. ## Verified on llvmpipe Vulkan (CPU software, AMD MI300X cloud droplet) Patched ggml-vulkan.cpp temporarily during repro to allow llvmpipe (normally filtered out as eCpu); patch reverted before commit. The SET_ROWS abort is a backend-capability check at graph build time so it fires regardless of GPU vs CPU Vulkan backend. | ctk / ctv | tg16 (t/s) | status | |-------------------|-----------:|---------------| | q4_0 / q4_0 | 17.68 | baseline | | q4_0 / turbo3 | 5.91 | already worked| | q4_0 / turbo4 | 6.14 | was aborting | | q4_0 / turbo2 | 5.65 | was aborting | llvmpipe perf numbers are not meaningful (CPU-emulated Vulkan); they are reported here only to confirm the abort is gone and the kernels run end-to-end without divergence. ## Needs GPU validation Cannot validate GPU shader correctness on the droplet (MI300X SR-IOV VF does not expose itself to RADV/amdvlk on cloud). Specifically: - Subgroup shuffle / ballot behavior on real GPU subgroup sizes - Shader compilation under non-llvmpipe Vulkan drivers - PPL / quality on the actual quantization math @dpblnt @apollosenvy if either of you has cycles, would appreciate a quick rebuild on RDNA Vulkan (gfx1100/gfx1200) to confirm: 1. The SET_ROWS abort that triggered ggml-org#50 is gone 2. Output coherence on turbo4 V (not garbage tokens) 3. PPL stays in the expected ballpark vs the CUDA / Metal implementations of the same quants Closes ggml-org#50. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…ows-turbo24 vulkan: add SET_ROWS support for turbo2_0 and turbo4_0 (ggml-org#50)
…869) (ggml-org#22267) * server: clamp n_discard to non-negative at JSON parse boundary (CVE-2026-21869) A negative n_discard from client JSON causes heap-buffer-overflow in update_slots() context-shift loop (CWE-787, CVSS 8.8). Clamp to 0 at ingress; n_discard=0 already triggers auto-discard (n_left/2). Ref: GHSA-8947-pfff-2f3c * cont : cleaner * cont : cleanerer * cont : cleanest --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
…ard-clamp security: cherry-pick CVE-2026-21869 (n_discard heap-buffer-overflow in server)
When users set --cache-type-k turbo3 (or turbo2/4) without
explicitly setting --cache-type-v, V defaults to F16 and the
(TURBOx_0, F16) pair hits ggml_cuda_flash_attn_ext_vec without a
matching FATTN_VEC_CASE → GGML_ABORT("fatal error") at fattn.cu:348.
The reverse direction (F16 K + turbo V) was already instantiated
for all three turbo variants. This adds the matching
(turbo K + F16 V) pairs.
Reported on Radeon 8060S (gfx1151) with Qwen2.5-7B-Instruct-Q4_K_M
running --cache-type-k turbo3 --cache-type-v f16. Not GPU-specific
— same crash on any CUDA/HIP target with that flag combo.
Files added:
- fattn-vec-instance-turbo2_0-f16.cu
- fattn-vec-instance-turbo3_0-f16.cu
- fattn-vec-instance-turbo4_0-f16.cu
Files updated:
- fattn.cu (3 dispatch entries)
- fattn-vec.cuh (6 extern decls)
- CMakeLists.txt (3 entries in non-FA_ALL_QUANTS list)
Mac/Metal build verified. CUDA/HIP build needs validation on a
target with the toolchain (compile-only, no behavior change for
existing instantiated combos).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…g#138) When --kl-divergence-base is set, line 527 sizes the log_probs vector with int * int. For Qwen3-class models (vocab=151,936) at n_ctx=16384, the product is 16384 * 151940 = 2,489,610,240 which overflows INT32_MAX (2,147,483,647). The wrapped negative value sign-extends to a giant size_t when passed to vector::resize, exceeds vector::max_size, and throws std::length_error. Reproduced on M5 Max with Qwen3.6-35B-A3B Q8_0: llama-perplexity -m <model> -f wiki.test.raw -c 16384 -fa 1 \ --kl-divergence-base /tmp/baseline.dat perplexity: saving all logits to /tmp/baseline.dat perplexity: tokenizing the input .. perplexity: tokenization took 316 ms perplexity: calculating perplexity over 18 chunks ... libc++abi: terminating due to uncaught exception of type std::length_error: vector Smaller-vocab models (vocab <= ~128K) and shorter context (n_ctx <= 8192) do not trip the overflow, which is why this only surfaces on Qwen3 family at the standard 16K PPL bench depth. Same size_t cast already exists at line 514 for the parallel logits.reserve allocation; this brings line 527 to the same convention. After fix the same command runs cleanly to completion. Validated A/B on Qwen3.6-35B-A3B Q8_0 at 16K context against wikitext-2-raw: Cache | PPL | KL Div | Top-1 agree -------+------------------+-----------------+------------- f16 | 5.3513 +/- 0.032 | (baseline) | - q8_0 | 5.3508 +/- 0.032 | 0.0014 +/- 0e-5 | 98.708% turbo3 | 5.3907 +/- 0.032 | 0.0121 +/- 1e-4 | 95.293%
Commit e69af78 added 3 new dispatch entries to fattn.cu for the (turbo2/3/4, F16) mixed-KV combinations and the matching template-instance .cu files, but only updated ggml/src/ggml-cuda/CMakeLists.txt. The parallel list in ggml/src/ggml-hip/CMakeLists.txt was missed, so the HIP build links without those instantiations and fails: ld.lld: error: undefined symbol: void ggml_cuda_flash_attn_ext_vec_case<64, TURBO3_0, F16>(...) void ggml_cuda_flash_attn_ext_vec_case<128, TURBO3_0, F16>(...) void ggml_cuda_flash_attn_ext_vec_case<256, TURBO3_0, F16>(...) (and same for TURBO2_0, TURBO4_0) clang++: error: linker command failed with exit code 1 Surfaced first by LocalAI's hipblas-turboquant build job (mudler/LocalAI#9740 CI). Fix is mechanical: mirror the 3 new entries from the CUDA CMakeLists into the HIP CMakeLists, paired next to their existing f16-X counterparts.
- Add dequant shaders for turbo2_0, turbo4_0, tq3_1s with WHT/RHT - Add mul_mat_vec shader for tq3_1s - Add flash attention support for turbo2_0, turbo3_0, turbo4_0 - Fix copy_to_quant/copy_from_quant for TurboQuant types - Fix dequant_funcs_cm2.glsl typo (grid -> g2) - Fix vulkan-shaders-gen: use vulkan1.3 target for _cm1/_int8/q8_1 - Add turbo2_0/turbo4_0 to FA scalar and cm1 shader generation - Add pre-built ggml-vulkan-shaders.hpp with all new shader externs Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
The local_size_x = 128 WHT butterfly path in copy_to_quant.comp was gated on DATA_A_TURBO3_0 only. TURBO2_0 and TURBO4_0 fell through to the generic 32-thread path, producing an incomplete 32-of-128 WHT and corrupting the KV cache entries for those types. Fix: extend the condition to cover all three turbo types. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Hoist the norm/qs/signs loads out of the per-element loop in dequantize4 for all three turbo types. Since iqs is always a multiple of 4, all four elements within a dequantize4 call share the same qs byte (and the same signs byte for turbo3_0). This reduces the number of buffer loads from 4x2-4x3 per call down to 2-3, lowering cache pressure in the Q*K inner loop. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Wire up pipeline_dequant_mul_mat_vec_f32_f32, _f16_f32, and _id_f32 for GGML_TYPE_TURBO2_0/3_0/4_0 and add them to the switch in ggml_vk_get_dequantize_mul_mat_vec() and ggml_vk_get_dequantize_mul_mat_vec_id(). Without this the decode path fell back to a slower dequantize-then-matmul route instead of the dedicated quantized kernel. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
|
I've tested this on an RX 580 8GB and it seems to work while having 2-3x performance gain on text ingestion. Thanks. |
|
I've also verified this on an GTX 960M. Seems to work there as well. |
Author
Good, can you share the test? |
Not sure what exact test, but when I entered a prompt, it didn't corrupt itself after a while, it behaved as it should. If there's a more objective way to test this, please let me know. The sanity test that I ran on said card: Details |
Author
ok... Can you make a cross-test with baseline? By baseline, I mean non-quantized KV cache. or use llama-bench in this way ./build/bin/llama-bench then |
Here are the results: DetailsFirst note: Despite the Intel GPU being recognized as a Vulkan device, llama-cpp was not using it (I confirmed through a GPU monitoring software) Second note: I can confirm that All in all, this does seem to work, just not for |
TheTom
force-pushed
the
feature/turboquant-kv-cache
branch
from
August 2, 2026 03:02
8a891f4 to
28c68fe
Compare
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
17 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.
Summary
This PR adds Vulkan backend support for TurboQuant KV cache formats and wires the missing optimized execution paths for TurboQuant types.
The main goal is to make TurboQuant KV cache usable on Vulkan for long-context inference while avoiding unnecessary fallback paths during decode and flash-attention execution.
Included changes:
TURBO2_0TURBO3_0TURBO4_0TQ3_1STQ4_1STURBO2_0,TURBO3_0, andTURBO4_0copy_to_quant/copy_from_quanthandling for TurboQuant typesSET_ROWSworkgroup configuration forTURBO2_0andTURBO4_0dequantize4()inflash_attn_base.glslmul_mat_vec/mul_mat_vec_idVulkan pipelines for TurboQuant typesMotivation
Before this PR, the Vulkan TurboQuant path was incomplete for KV-cache usage.
In particular:
TURBO2_0andTURBO4_0could fall through to the wrongSET_ROWSworkgroup size, producing incomplete WHT output and corrupting KV cache entries;mul_mat_vecpipelines and fall back to a slower dequantize-then-matmul path;This PR makes the Vulkan TurboQuant KV-cache path more complete, fixes correctness issues, and reduces avoidable overhead in the attention/decode hot paths.
Details
TurboQuant KV cache support
Adds Vulkan shader support and generated shader declarations for the TurboQuant KV-cache path, including dequantization, copy-to-quant, copy-from-quant, flash-attention integration, and required shader registration.
Supported internal formats in this PR:
TURBO2_0TURBO3_0TURBO4_0TQ3_1STQ4_1SUser-facing CLI cache type names:
turbo2turbo3turbo4tq3_1stq4_1sCorrectness fix for
SET_ROWSThe WHT butterfly path in
copy_to_quant.compusedlocal_size_x = 128, but the condition was only enabled forDATA_A_TURBO3_0.As a result,
TURBO2_0andTURBO4_0could fall back to the generic 32-thread path, producing incomplete 32/128 WHT output and corrupting KV-cache entries.This PR extends the condition to cover all three TurboQuant types:
DATA_A_TURBO2_0DATA_A_TURBO3_0DATA_A_TURBO4_0Flash-attention dequant optimization
The TurboQuant
dequantize4()path inflash_attn_base.glslnow hoists shared loads out of the per-element loop.Since
iqsis always a multiple of 4, the four elements handled bydequantize4()share the same packedqsbyte, and forTURBO3_0also the same signs byte.This reduces repeated buffer loads in the Q*K inner loop and lowers cache pressure during flash attention.
Dedicated Vulkan mat-vec pipelines
This PR registers dedicated Vulkan
mul_mat_vecandmul_mat_vec_idpipelines for:TURBO2_0TURBO3_0TURBO4_0Without this, the decode path could miss the optimized quantized kernels and fall back to a slower dequantize-then-matmul route.
Testing
Tested locally with Vulkan backend using TurboQuant KV cache enabled on AMD RX570 8GB.
It works perfectly; I haven't noticed any token generation issues on any of the Turbo variants. However, without native FP16 on the card, I can't fully validate the Turbo4 tests.
What I can add is that Turbo2 is faster, even though everything is computed in FP32.
Suggested build test: