SYCL Turboquant implementation attempt - #144
Open
cclecle wants to merge 230 commits into
Open
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>
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>
…eTom#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>
…m#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>
…heTom#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>
…heTom#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>
…om#23 Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…heTom#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>
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>
…m#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>
…om#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) TheTom#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>
…m#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>
…heTom#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 TheTom#3 (TURBO_D). TheTom#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>
Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…Tom#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>
…heTom#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>
…eTom#31 TheTom#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>
…s-turbo24 vulkan: add SET_ROWS support for turbo2_0 and turbo4_0 (TheTom#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>
…d-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>
) 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.
…ache Wholesale sync of 384 upstream commits since merge-base 7fc1c4e (2026-04-22). Headline upstream feature: MTP / Multi-Token Prediction (ggml-org#22673) + spec-decoding stack (ggml-org#22838 parallel drafting, ggml-org#22227 spec-simple checkpoints, ggml-org#19493 server spec checkpointing, plus 5 spec bug-fixes). 11 conflicts resolved across CUDA fattn / Metal / Vulkan / common: ggml/src/ggml-cuda/fattn-mma-f16.cuh RDNA config matrix: union TQ's (640, 512) entries with upstream's expanded (112..576) RDNA matrix. Took upstream's new sentinel fallback (no ampere fallback for RDNA). ggml/src/ggml-cuda/fattn.cu - Extended hoisted ncols2_max to include 640 head-dim. - Volta: dropped TQ's local ncols2_max redefinition in favor of upstream's hoisted version (with 640 added). - WMMA gate: union exclusions (40, 72, 192, 512, 576, 640). - Preserved TQ's RDNA4 vector-kernel branch for TurboQuant cache types (renamed inner gqa_ratio_eff_rdna4 to avoid shadowing); took upstream's restructured MFMA/CDNA path verbatim. ggml/src/ggml-cuda/ggml-cuda.cu Supported-op switch: union TQ's GGML_OP_TURBO_WHT case with upstream's GGML_OP_ADD/SUB/MUL/DIV FP16 cases. ggml/src/ggml-metal/ggml-metal-device.h Kept TQ's get_pipeline_turbo_wht declaration; took upstream's new get_pipeline_mul_mv_ext(lib, const ggml_tensor * op, ...) signature (replaces split tsrc0/tsrc1 args). ggml/src/ggml-metal/ggml-metal-device.cpp Kept TQ's get_pipeline_turbo_wht implementation; took upstream's new get_pipeline_mul_mv_ext signature — body already uses op-> for tsrc0/tsrc1/ne12/r2/r3. ggml/src/ggml-metal/ggml-metal-ops.cpp Preserved TQ's is_tq_weight rotate→matmul→unrotate path with original hardcoded dispatch shape. Updated non-TQ fallback to upstream's pipeline- param dispatch (pipeline.nr0 / nr1 / nsg + (ne11+nr1-1)/nr1 shape). ggml/src/ggml-vulkan/* (3 files) Upstream-wholesale via `git checkout --theirs`. Upstream architecturally refactored FA from compile-time DATA_A_* variants to runtime FaTypeK/FaTypeV spec-constant switches. TQ's TURBO3_0 GLSL path is DEFERRED — Vulkan TURBO3_0 support needs re-implementation against the new architecture in a follow-up PR. Mac mini + M5 Max have no Vulkan; no in-house validation path for an immediate re-adaptation. common/arg.cpp --spec-default: took upstream's new struct shape (params.speculative.types vector + params.speculative.ngram_mod.{n_match,n_min,n_max}). common/speculative.cpp Low-acceptance reset: took upstream's sinfo.n_low / sinfo.i_last (variables moved into sinfo struct). NOT-CONFLICTED upstream additions that touch TQ-adjacent surface (auto-merged clean, but worth eyes during review): - src/llama-memory-recurrent.{cpp,h} (MTP rollback API) - src/llama-memory-hybrid.{cpp,h} (recall feedback_llama_memory_types + feedback_layer0_hybrid_trap) - src/llama-graph.cpp, src/llama-kv-cache.cpp, src/llama-context.cpp - tools/server/server-context.cpp (+~1100 lines: MTP + parallel drafting + spec checkpointing) - src/models/qwen35*.cpp, qwen3next.cpp, delta-net-base.cpp (entirely new in upstream — MTP draft-head integration) Known-deferred follow-ups: 1. Vulkan TURBO3_0 re-implementation against runtime spec-constant FA arch 2. PR ggml-org#21245 QKV refactor helpers — landed; TQ models not migrated to use them. Migrate in a focused follow-up; do not bundle here. Validation gate (pending): M2 mini PPL/decode comparison @ Qwen2.5-7B-Q8_0 K=q8_0/V=turbo4 asymmetric ctx 2048 + 16384 — see PR body. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> Co-Authored-By: tturney@psyguard.ai
…ed pipeline
Regression introduced by the upstream b9190 merge. Upstream moved ne12/ne13/r2/r3
from kernel args to function constants (FC_MUL_MM + 2..5) in kernel_mul_mm.
get_pipeline_mul_mm was updated to set them; get_pipeline_mul_mm_tq_rotated was
not, so the TQ-rotated kernel templates (which reuse the same kernel_mul_mm
body) read those constants as zero, producing wrong tensor offsets → NaN/inf
outputs.
Symptoms caught by extended M2 mini validation:
- test-backend-ops MTL0: ALL MUL_MAT(type_a=tq3_1s|tq4_1s, ...) FAIL with
ERR = inf > 0.000500000 (and one NaN at index 136). 279 TQ MUL_MAT tests
pass on TQ tip 5aeb2fd.
- Qwen2.5-7B-TQ4_1S PPL: 146/146 chunks "nan" both dense and asym KV
(vs 6.7530 / 6.7887 baseline).
- Qwen2.5-7B-TQ4_1S decode bench tg128: 7.47 → 3.95 t/s = -47% regression,
pp128 variance ±0.27 → ±40.85 indicating dispatch chaos.
Fix:
- Compute ne12/ne13/r2/r3 in get_pipeline_mul_mm_tq_rotated identical to
get_pipeline_mul_mm.
- Set FC_MUL_MM + 2..5 alongside the existing bc_inp/bc_out constants.
- Include ne12/ne13/r2/r3 in the pipeline cache name so different tensor
shapes don't collide on a single compiled pipeline (cache poisoning).
The MUL_MAT_ID variant (get_pipeline_mul_mm_id_tq_rotated) mirrors
get_pipeline_mul_mm_id which only sets bc_inp — kernel_mul_mm_id is a
different template and doesn't need ne12/ne13/r2/r3, so no change there.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Co-Authored-By: tturney@psyguard.ai
CI failure (every -Werror build job — ubuntu-latest-cuda, ggml-ci-x64-cpu-*,
arm64-cpu-*, macOS-latest-{x64,arm64,arm64-webgpu}, ubuntu-22-{hip-quality-check,musa},
android-arm64, ubuntu-cpu x64-22.04):
ggml/src/ggml-turbo-quant.c:247:15: error: no previous prototype for
'turbo_cpu_fwht_inverse' [-Werror=missing-prototypes]
Pre-existing issue on the TQ branch — the function is defined GGML_API but
has no prototype declaration. M5 Max + M2 mini builds don't use -Werror so it
slid through local validation. Upstream CI does.
Fix: forward-declare the function near the top of the .c file. Matches the
extern declaration already used by tests/test-turbo-quant.c. No semantic change.
The other GGML_API symbol in this file (turbo3_cpu_wht_group_size) is a
variable, not a function — -Wmissing-prototypes does not apply.
Flagged by @pacak on PR TheTom#146 (ubuntu-latest-cuda CI). Restores green CI on
all -Werror build jobs without affecting any runtime path.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Co-Authored-By: tturney@psyguard.ai
Two CI build failures both pre-existing on TQ tip, exposed by upstream-policy CI on PR TheTom#146: 1. ubuntu-22-hip-quality-check — fattn-common.cuh:1312-1313 in TQ's HIP hip_f16_alloc destructor calls hipStreamSynchronize / hipFree without consuming the return value. HIP's recent runtime declares these [[nodiscard]] and the HIP quality build uses -Werror: error: ignoring return value of type 'hipError_t' declared with 'nodiscard' attribute [-Werror,-Wunused-value] Fix: (void) cast both calls. We're in a destructor and can't propagate errors anyway; intent is fire-and-forget cleanup. Matches the idiom used in upstream code for the same situation. 2. ubuntu-22-musa — turbo-quant.cuh uses cudaMemcpyToSymbol in InnerQ calibration; MUSA's vendor header (vendors/musa.h) aliases every other cudaMemcpy* variant but missed cudaMemcpyToSymbol. Result on MUSA build: error: use of undeclared identifier 'cudaMemcpyToSymbol'; did you mean 'musaMemcpyToSymbol'? Fix: add the missing alias next to the other cudaMemcpy* defines. Mirrors the same alias already present in vendors/hip.h:143. Both are TQ-only paths (HIP f16 alloc was added in 0757ff4 2026-04-18; MUSA was never built against TQ in-tree). M5 Max + M2 mini local Metal builds unaffected by either change. Flagged by @pacak on PR TheTom#146 (ubuntu-latest-cuda + cross-vendor CI fails). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> Co-Authored-By: tturney@psyguard.ai
Cherry-picked from upstream ggml-org/llama.cpp@87589042c (merged 2026-05-17). option(LLAMA_BUILD_WEBUI ... ON) always leaves the deprecated flag DEFINED, so the compat-block guard `AND NOT DEFINED LLAMA_BUILD_UI` never fires. tools/ui/CMakeLists.txt then ORs both flags, so passing only the new `-DLLAMA_BUILD_UI=OFF` was silently ignored. Removes the deprecated options and simplifies the compat block + UI gate to a single flag. Fixes the nix-sandbox build failure reported by @arch-fan and @pacak on PR TheTom#146 — both hit the resulting xxd.cmake crash when an empty tools/ui/dist/index.html was produced by failed npm + HF Bucket provisioning. After this cherry-pick, `-DLLAMA_BUILD_UI=OFF` alone works as documented. Co-Authored-By: TheTom <tturney@psyguard.ai> Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Round 2 of CI fixes addressing the remaining red jobs on the b9190 sync
PR. All were pre-existing TQ-tip bugs exposed by upstream CI's -Werror
policy (M5 Max + M2 mini local builds don't use -Werror).
1. ggml/src/ggml-cuda/fattn-mma-f16.cuh — fall back to ampere config
(not zero-sentinel) in get_config_rdna
----------------------------------------------------------------
Reverts the round-1 conflict choice. Round 1 took upstream's new
sentinel `fattn_mma_config(32, 1, 0, 0, 0, 0, 0, false)` for the
RDNA fallback. Template instances like
fattn-mma-f16-instance-ncols1_1-ncols2_16.cu do constexpr arithmetic
on the returned config (np = nwarps * cols_per_warp / ncols, etc).
nwarps=0 from the sentinel propagates to np=0, triggering compile-
time div/mod-by-zero at lines 1265/1371/1375/1512/1519/1572. HIP
quality build is -Werror,-Wdivision-by-zero so it errors out.
TQ-tip behavior (delegate to ampere) returns a valid config —
restore it. Keeps all (640, 512) RDNA entries unioned in round 1.
2. ggml/src/ggml-cuda/vendors/musa.h — add cudaMemcpyFromSymbol alias
----------------------------------------------------------------
turbo-quant.cuh InnerQ calibration uses both cudaMemcpyToSymbol AND
cudaMemcpyFromSymbol. Round-1 fix added _ToSymbol; _FromSymbol was
missed. Mirrors vendors/hip.h line 142.
3. src/llama-kv-cache.cpp — [[maybe_unused]] stubs + remove unused `il`
----------------------------------------------------------------
The non-CUDA stub block (g_innerq_finalized, g_innerq_scale_inv_host,
turbo_innerq_needs_tensor_update, turbo_innerq_mark_tensor_updated)
are declared static but every consumer is gated by #ifdef GGML_USE_CUDA,
so the file-local copies look unused on non-CUDA builds. Annotate
with [[maybe_unused]]. Also drops two `const uint32_t il = layer.il;`
locals in the state-save k/v writer loops where `il` was unreferenced —
dead-code from a removed logging pass.
4. scripts/xxd.cmake — defensive quote of ${hex_data}
----------------------------------------------------------------
Belt-and-suspenders for the LLAMA_BUILD_UI nix-sandbox failure. The
primary fix is the cherry-pick of upstream PR ggml-org#23190 (previous
commit), which makes -DLLAMA_BUILD_UI=OFF actually work. This patch
makes the underlying xxd.cmake robust: when an empty UI source file
slips through, produce a 0-length .hpp instead of crashing with
cmake's cryptic "string sub-command LENGTH requires two arguments"
error. Worth proposing upstream as a follow-up.
Local Metal build green on M5 Max with all four fixes applied.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Co-Authored-By: tturney@psyguard.ai
arch-fan's next nix-sandbox build (after PR ggml-org#23190 cherry-pick + earlier empty-input defensive quote) hit a different xxd.cmake failure: scripts/xxd.cmake:10 (file): file failed to open for reading (No such file or directory): /build/source/build/tools/ui/dist/bundle.js Empty-file case (LENGTH error) was already handled by quoting the variable. This is the sibling case: file READ itself fails when the UI provisioning flow leaves an asset missing entirely (npm absent AND HF Bucket download blocked → some assets created empty, some not created at all). Fix: early-return with a valid 0-byte symbol when ${INPUT} doesn't exist. Also unify the empty-content path to emit {0} instead of {} (zero-element array initializer is C++ extension, not portable). Verified end-to-end on M5 Max by reproducing arch-fan's exact conditions: build/tools/ui/dist/ removed, PATH stripped of npm, LLAMA_USE_PREBUILT_UI=OFF. Without the fix, build crashes on bundle.js.hpp generation. With the fix, all four .hpp files generate as 0-byte symbols, llama-ui target completes cleanly, server builds with LLAMA_UI_DEFAULT_ENABLED=0 (no embedded UI but no crash) — exactly upstream's intended graceful degradation. No effect on normal builds with UI assets present (regenerated all 4 .hpp files at original 26MB / 2.5MB / 34KB / 1.4KB sizes, byte- identical to pre-fix output). Worth proposing upstream as defensive hardening for the xxd helper. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> Co-Authored-By: tturney@psyguard.ai
sync: upstream master b9190 + MTP/spec stack (DO NOT MERGE — tester review)
Short README calling out what this fork adds vs upstream ggml-org/llama.cpp: TurboQuant KV-cache and weight quantization types (turbo3, turbo4, TQ3_1S, TQ4_1S) and their CUDA / HIP-ROCm / Metal / Vulkan kernel integrations. Points at TheTom/turboquant_plus for the codec design and papers, and TheTom/tqkit for cross-backend bench results. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> Co-Authored-By: tturney@psyguard.ai
Restore the upstream llama.cpp README so users have the standard build / model / bindings docs, with a short fork-specific intro at the top calling out the TurboQuant additions and pointing at TheTom/turboquant_plus for the codec design.
Cross-checked against commit log diff vs upstream + paper corpus: - add turbo2 KV-cache format (was already in codebase but missed) - add auto-asymmetric K/V compression policy (PR TheTom#53) - add Boundary V (layer-aware experimental V compression for turbo2-V) - add Sparse V dequantization (on all Metal targets) - add turbo VEC flash attention +9% decode on CUDA - add V2.1 fused Metal TQ kernels - add TurboFlash Apple10 known-limitation caveat - add Vulkan SET_ROWS support for turbo2/turbo4 - note turbo block size moved from 32 to 128 Direct paper links inline (TheTom/turboquant_plus papers) instead of vague references. Drop the tqkit cross-reference per scope.
Move fork README to a professional structure suitable for downstream projects evaluating the integration: - Title + tagline + badges (license / status / paper corpus) - Production deployments section (LocalAI, AtomicChat, others) - Status table (branch, commits ahead, upstream tracking) - Quantization types table with bits / notes / paper links - Compression policies (asymmetric K/V, Boundary V, sparse V) - Backend coverage matrix with per-backend notes + caveats - Model-family support + operational fixes the fork carries - Quick start + usage with concrete invocations - Citation pointer to the TurboQuant+ paper corpus Upstream llama.cpp README preserved verbatim below the fork section so users still have the standard build / model / bindings docs.
- Add Chronara.io (quantum-safe fintech infrastructure) as a production user of this fork. - Link AtomicChat (https://atomic.chat/).
Restructure the KV-cache usage section to make the asymmetric-turbo
pattern unmissable and to guide users to escalate compression instead
of starting at maximum and walking back:
- Add prominent "Start light, then compress" callout — some model
families are more delicate (small models, certain MoE configs,
quant-sensitive instruction-tuned variants).
- Reframe the recommendations table as a 5-step ladder, from safest
('f16'-K + 'turbo3'-V, first-contact) through recommended default
('q8_0'-K + 'turbo3'-V, asymmetric turbo sweet spot) to aggressive
V and MoE-aware aggressive, ending with the discouraged symmetric-K
config and its failure-mode citation.
- Add concrete CLI examples for steps 1, 2, 3.
- Add closing reminder that the frontier is per-model — walk back a
step if quality drops.
Two corrections:
1. Fix the turbo compression ladder. Higher turbo number = more bits
per element = less aggressive. The V-side ladder runs turbo4
(lightest) → turbo3 → turbo2 (heaviest), not the other way around.
- Step 1 (safest start) is now f16-K + turbo4-V (was incorrectly
listed as turbo3 V)
- Step 2 (conservative) adds q8_0-K
- Step 3 (recommended default, "asymmetric turbo") stays q8_0-K +
turbo3-V — the paper's sweet spot
- Steps 4-5 (aggressive V) reach turbo2-V with Boundary V
- Step 6 (discouraged) is symmetric K compression
- Updated CLI examples accordingly.
2. Inline the title of the asymmetric K/V paper ("Asymmetric K/V Cache
Compression: Why V is Free and K is Everything") and call out that
it documents the failure modes you'll hit if K is compressed
aggressively — required reading before step 6.
Add a 'Lineage — why the +' subsection making the relationship to Google's original TurboQuant paper (ICLR 2026) explicit. That paper introduced the Walsh-Hadamard-rotated polar codebook scheme and hit 4.6× compression at ~1% PPL loss. The '+' denotes the substantial extension work in this project: asymmetric K/V policy, Boundary V, attention-gated sparse V dequantization, TQ3_1S/TQ4_1S weights, turbo2/turbo4 tiers, cross- backend kernels (CUDA dp4a, HIP/ROCm, Vulkan coopmat, Metal TurboFlash/V2.1 fused), and model-family quality/operational fixes. The original TurboQuant codec remains the foundation; this is the acknowledgment of that lineage.
Owner
|
Thanks for this, really appreciate the SYCL port. We reviewed it and we are happy to merge, it is cleanly isolated to Before we pull it we just need some compilation and test evidence, since we don't have an Intel/oneAPI box on our side to verify it ourselves:
Our A couple of small things worth a look while you're at it: the new |
Owner
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.
Overview
SYSCL Implementation using Claude, tested a little bit on Intel A380 and oneapi 2025.2.
Additional information
Code not properly reviewed.
Requirements