feat: ROCm/HIP support for turbo3 KV cache (gfx1100/RDNA3) by apollosenvy · Pull Request #5 · TheTom/llama-cpp-turboquant

apollosenvy · 2026-03-27T11:13:57Z

Summary

Adds complete ROCm/HIP backend support for turbo3 KV cache quantization, enabling 4.6x KV compression on AMD GPUs.

Tested on: AMD Radeon RX 7900 XTX (gfx1100), ROCm 7.1, Qwen3.5-27B Q4_K_M

What's included

GPU kernels: dequant, vec_dot, Walsh-Hadamard Transform, quantize, GET_ROWS, CPY
Flash attention: F16 dequant bridge (turbo3 -> F16 temp buffer -> F16 FA kernel). The native turbo3 FA template compiles but produces broken kernels on HIP/RDNA3, so we route through the existing need_f16_K/V conversion path in launch_fattn.
CPU fallback: vec_dot + type traits for graph scheduler fallback paths
All dispatch wiring: supports_op for MUL_MAT, GET_ROWS, SET_ROWS, CPY, FLASH_ATTN_EXT, TURBO_WHT

Critical bugfix: K/V rotation before cache write

Found and fixed a quality issue: K and V vectors were being quantized to turbo3 without WHT rotation. The graph rotated Q (forward) and inverse-rotated the attention output, but the K/V write path in cpy_k()/cpy_v() was missing the forward rotation step.

Before fix: <accepts high pressure "Itself" - ount>& & nsp; a- xii'
After fix: 7 x 8 = 56. This is one of the most common multiplication facts.

This fix is in src/llama-kv-cache.cpp and may also benefit the Metal path if it has the same omission.

Performance

Metric	FP16 KV	turbo3 KV	Ratio
Prompt	90 tok/s	82 tok/s	91%
Generation	21 tok/s	12 tok/s	57%*
KV cache (32K ctx)	2048 MB	448 MB	4.6x

*Generation speed regression is from AMD_SERIALIZE_KERNEL=3 requirement (see below).

Known limitation

Requires AMD_SERIALIZE_KERNEL=3 environment variable. Without it, the multi-stream graph scheduler has a race condition between KV cache write and FA read ops. Per-op cudaStreamSynchronize is not sufficient. This is a ROCm runtime issue, not a turbo3 bug.

Files changed (24 files, +595 lines)

New files (5):

dequantize-turbo.cuh - Device dequant (centroid LUT)
vecdot-turbo.cuh - MMVQ vec_dot
turbo-wht.cu/cuh - WHT butterfly kernel
fattn-vec-instance-turbo3_0-turbo3_0.cu - FA template (D=256)

Modified (19): common.cuh, mmvq.cu, fattn-common.cuh, fattn-vec.cuh, fattn.cu, convert.cu, cpy-utils.cuh, cpy.cu, set-rows.cu, getrows.cu, ggml-cuda.cu, CMakeLists.txt (cuda+hip), ggml-turbo-quant.c, ggml-quants.h, ggml-cpu.c, quants.h, llama-kv-cache.cpp

Test plan

Standalone GPU kernel test (quantize/dequant round-trip)
Server starts with --cache-type-k turbo3 --cache-type-v turbo3
KV cache allocates at expected compression ratio
Inference produces coherent output
Math reasoning correct (7*8=56)
Perplexity benchmark vs upstream Metal numbers
Multi-turn conversation stability

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>

…m#27 Full investigation log with all tests, results, and the root cause. Upstream TurboQuant activity tracked in TheTom#27. 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>

apollosenvy · 2026-03-30T00:21:07Z

Split per your suggestion:

PR feat: HIP/ROCm support for turbo3/turbo2 (7900 XTX) #31 — ROCm/HIP port only. Single commit, minimal. Rebased on your latest (includes Boundary V). Also excludes D>=576 fattn-tile instances that exceed HIP's 64KB local memory limit.
Temporal decay — Will open as separate PR once Metal kernel or supports_op guard is added. The decay op shouldn't dispatch to backends that don't implement it.
head_dim fallback — Your commit 4c4511c14 already handles this. No separate PR needed from our side.

Re: modes 5/6/7 — we didn't remove them. Our branch was based on a commit before PR #30. The new PR #31 is rebased on your latest which includes Boundary V, so modes 5/6/7 are intact.

Complete experiment log: #1 4-mag LUT: 15.1 at 8K (BEST, +38%) #2 Batched extract: 13.7 (+25%) #3 Inline FA block: 13.5 (I-cache pressure) #4 Deferred norm: 12.9 (loses ILP) #5 2-pair half2: 12.0 (ternary overhead) #6 Select chain: 11.9 (branches kill) #7 Bit-arithmetic: 11.6 (ALU too heavy) #8 FMA branchless: 11.4 (ALU still too heavy) #9 Named-reg ternary: 10.3 (branches worst) #10 Main (8-LUT): 10.95 (baseline) #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

Complete experiment log: TheTom#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

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

Complete experiment log: #1 4-mag LUT: 15.1 at 8K (BEST, +38%) #2 Batched extract: 13.7 (+25%) #3 Inline FA block: 13.5 (I-cache pressure) #4 Deferred norm: 12.9 (loses ILP) #5 2-pair half2: 12.0 (ternary overhead) #6 Select chain: 11.9 (branches kill) #7 Bit-arithmetic: 11.6 (ALU too heavy) #8 FMA branchless: 11.4 (ALU still too heavy) #9 Named-reg ternary: 10.3 (branches worst) #10 Main (8-LUT): 10.95 (baseline) #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

apollosenvy · 2026-04-17T23:14:41Z

Closing — this got split into PR #31 (ROCm/HIP, merged) and PR #33 (Vulkan turbo3, merged), plus follow-up cleanup in PR #66. Nothing unique left on rocm-turbo3-pr. Thanks for the reviews on all of those.

Complete experiment log: TheTom#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

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

…0 findings:0

@2048

commit bfe452a6d315b776c2e06f7f11c45bd6cd9766b4 Author: LogicDaemon <LogicDaemon@users.noreply.github.com> Date: Sat Sep 5 13:29:03 2026 +0400 Squashed commit of the following: commit 10f8aff Author: thecodacus <thecodacus@gmail.com> Date: Tue Aug 25 23:39:14 2026 +0530 perf: lower default --cpu-tp-margin to 512 MiB Measured on a 12GB RTX 3060 with a dense 27B Q4_K_M: tightening the margin raises both decode and prefill, since every MiB reclaimed moves weights off the CPU leg. 1024 -> 640 -> 384 MiB gave 7.4 -> 7.8 -> 8.1 t/s decode and 14.4 -> 15.2 -> 16.5 t/s prefill. 512 keeps headroom for larger contexts; raise it if the GPU runs out of memory. commit e406c8c Author: thecodacus <thecodacus@gmail.com> Date: Tue Aug 25 22:08:09 2026 +0530 feat: add --cpu-tp to use the CPU as a tensor-parallel device llama.cpp's tensor-parallel meta device already dispatches each member backend asynchronously before synchronising, but the CPU is excluded from the device list. Adding it lets the GPU and CPU compute row-slices of every weight matrix concurrently instead of running one after the other. The GPU share is sized from the measured free VRAM (minus --cpu-tp-margin) so VRAM is saturated, and the CPU takes only the remainder. The CPU is appended last so the GPU subgraph is already in flight when the CPU backend blocks on its own. commit d927e7d Author: Anirban Kar <thecodacus@gmail.com> Date: Tue Aug 25 15:58:16 2026 +0530 docs: add MoE expert cache usage and tuning guide to README (TheTom#8) * docs: add MoE expert cache usage and tuning guide to README * docs: document async CPU split overlap and VRAM headroom rule commit 0ac3d9b Author: Anirban Kar <thecodacus@gmail.com> Date: Sun Aug 23 04:08:21 2026 +0530 feat: overlap CPU cold-expert and GPU hot-expert computation in the MoE cache (TheTom#9) * feat: overlap CPU cold-expert chain with GPU hot chain (GGML_SCHED_ASYNC_CPU) The dual-path MoE cache ran its two chains serially: the scheduler blocks on CPU splits, and the cold chain's input copies drain the whole CUDA stream. Three changes make them concurrent, gated behind GGML_SCHED_ASYNC_CPU=1: - sched: CPU splits execute on a persistent worker thread; the main loop keeps launching later splits and joins lazily at the first split that is CPU or reads a CPU-resident input - graph: cold chain is built before the hot chain so the hot split follows it and carries no cross-backend inputs (launches while the worker computes) - graph: the merge add is pinned to CPU so it does not pull the hot chain into a split that waits on the cold result * feat: GGML_SCHED_ASYNC_CPU=2 diagnostic mode (worker without graph restructure) * fix: mark expert cache pack buffer as weights Without the weights usage flag the scheduler's assignment pass does not pin the pack tensors to their backend, so a CPU-assigned consumer (the pinned merge in async mode) propagates backward and pulls the hot matmuls onto the CPU, copying pack weights over PCIe every layer. * feat: make async CPU splits the default, expose --sched-async-cpu flag Replaces the GGML_SCHED_ASYNC_CPU env var: the overlap is now on by default and controlled through llama_context_params.sched_async_cpu. - ggml: ggml_backend_sched_set_async_cpu() setter, env read removed - llama: context param + cparams plumbing, graph gate reads cparams - common: --sched-async-cpu / --no-sched-async-cpu (server, cli) - llama-bench: sched_async_cpu test dimension (--sched-async-cpu 0,1 benches both in one run, column shown as 'sac' when non-default) * fix: restrict CPU-pinned merge to decode-size batches For prefill batches the CPU merge and its per-layer activation copies cost more than the overlap hides (pp256 regressed 307 -> 257 t/s). Decode and MTP verify batches (n_tokens <= 8) keep the overlap. commit 8e5bc28 Author: Anirban Kar <thecodacus@gmail.com> Date: Sat Jul 25 03:54:47 2026 +0530 feat: wire MoE expert cache into laguna graph (TheTom#6) commit ac743f8 Author: Anirban Kar <thecodacus@gmail.com> Date: Sat Jul 25 00:41:02 2026 +0530 feat: MoE expert cache — keep hot routed experts resident in VRAM (TheTom#5) * feat: llama-moe-trace tool + expert-cache policy simulator Phase 0 of the MoE expert-cache work: capture routed expert ids (ffn_moe_topk) per layer per decode step via the scheduler eval callback — backend-agnostic, so CPU-offloaded expert layers are traced too. Companion simulate.py replays traces against cache policies (static top-S, LRU, LFU-decay, periodic re-election) and prints hit-rate vs VRAM-budget curves plus routing-skew stats to decide whether expert caching is worth building at all. * fix: moe-trace — common_init_from_params returns ptr with accessors * fix: moe-trace — include arg.h for common_params_parse * fix: moe-trace — take output path from MOE_TRACE_OUT env, -o not in this arg set * fix: moe-trace — topk tensor is a strided view; copy full byte range and index by nb Sizing the host buffer as n_used*n_tokens under-allocated: ffn_moe_topk is a view over the argsort output, so nb[1] spans n_expert ids per row and ggml_backend_tensor_get writes the whole strided range (heap corruption, malloc corrupted-top-size abort mid-prefill). * feat: moe-trace simulator v2 — model upload bytes/token and miss/upload cost per policy * feat: mul_mat_id tolerates id=-1 (skip) on CPU and CUDA fallback path Groundwork for the hot/cold expert-pack split: an id of -1 means the routed expert is not owned by this pack. The op contributes a zeroed output row for that slot so hot-pack and cold-pack results merge additively. CUDA fast paths (mmvq_id/mmq_id/mmf_id) still assert non-negative ids and are handled next. * feat: id=-1 skip support in CUDA mul_mat_id fast paths + tests mmid.cu gains a generic kernel that zeroes the dst rows of (token, slot) pairs whose expert id is -1; mmq and mmf launch it after the shared mm_ids_helper (which already drops unmatched ids from its compact mapping), and mmvq launches it before its kernels, which now early-return instead of using a negative id as a channel index (out of bounds read otherwise). test-backend-ops gains skip_ids variants of MUL_MAT_ID covering the mmvq/mmq/mmf/fallback dispatch corners. * fix: skip_ids MUL_MAT_ID test cases were inside an #if 0 block * fix: mm_ids_helper compact map corruption when ids contain -1 Negative ids inflated nex_prev for every expert, opening a phantom region at the head of the compact arrays that nothing wrote; the mmq quantize preprocessing iterates the flat range and read garbage row indices (illegal memory access). Skipped ids now occupy no compact position, ids_src1 is sentinel-filled, and the quantize gather/scatter kernels skip negative entries. * fix: mmvf id=-1 skip guard (float mul_mat_id vec path) Same treatment as mmvq: kernels early-return on a negative expert id instead of using it as a channel index, and the host launcher pre-zeroes the dst rows of skipped slots. Fixes wrong results (and a latent out-of-bounds channel read) for f16/f32 MUL_MAT_ID with -1 ids at small batch sizes. * fix: zero-init ids_dst before mm_ids_helper — garbage tail entries in the compact map caused pool-layout-dependent illegal memory accesses when ids contain -1 (tile-padded reads consumed uninitialized row indices) * feat: route MUL_MAT_ID with skip-capable ids away from mmq/mmf Ops whose ids may contain -1 announce it via op_params[0]; the CUDA dispatch then uses only the paths with verified skip support (mmvq for small batches, the general gather/scatter path otherwise). Large-batch mmq with sparse ids shows a pool-layout-dependent fault that resists isolation (all indexing audits pass; sanitizer and cuda-gdb unusable on the target driver); since the expert-pack design only decodes through packs, the safe route costs nothing in practice. mmq/mmf guards from earlier commits remain as hardening. * fix: op_params write in test — ggml_set_op_params_i32 is internal-only * feat: MoE expert cache loader — hot-pack build from routing profile init_moe_expert_cache() runs post-load: reads a moe-trace profile (GGML_MOE_CACHE_PROFILE) and slot count (GGML_MOE_CACHE_SLOTS), and for each CPU-resident MoE layer builds GPU-resident hot packs of the top-S experts plus i32 remap tables (hot: global->slot or -1; cold: global id or -1). Graph wiring (build_moe_ffn dual path) is the next commit; these tensors are unused until then, so this is behavior-neutral. * feat: MoE expert cache graph wiring — dual hot/cold mul_mat_id in build_moe_ffn When a layer carries hot packs, routed ids are remapped through moe_map_hot/cold (get_rows) and each expert projection runs as two skip-flagged mul_mat_id ops — hot pack on GPU, original tensors for cold — whose zero-padded outputs sum to the exact single-tensor result. Engaged for qwen35moe (separate gate/up/down, no expert scales); layers with merged gate_up, scales, or no packs use the unchanged path. * fix: flatten routed ids before map get_rows (batched-gather shape rule) * fix: topk ids are a strided view - cont before flatten for map lookup * fix: sched expert-copy path skips id=-1 (hot/cold expert split) * fix: restructure MoE pack dual path into full FFN chains per residency side Per-projection hot/cold interleave let the scheduler's split minimizer assign hot up/down matmuls to CPU, copying the GPU pack weights over PCIe every layer. Building one unbroken gate-up-swiglu-down chain per side keeps the hot pack on GPU; skipped (-1) rows stay zero through swiglu so a single add of the two chain outputs is exact. * feat: wire MoE expert cache into deepseek2 graph (GLM-4.7-Flash) * feat: add --moe-cache-profile and --moe-cache-slots flags Plumbed through llama_model_params; env vars remain as fallback so existing deployments keep working. Flags allow per-model settings in multi-model serving. * docs: fix moe-trace usage comment (output path is MOE_TRACE_OUT env) commit cd72c28 Author: Anirban Kar <thecodacus@gmail.com> Date: Fri Jul 24 17:53:41 2026 +0530 fix: restore closing brace lost in dequantize.cuh merge resolution commit 34f893e Merge: 40f2036 8f5ab83 Author: Anirban Kar <thecodacus@gmail.com> Date: Fri Jul 24 17:52:26 2026 +0530 merge: sync upstream master into perf commit 8f5ab83 Author: Matt Thompson <111157855+boondocklabs@users.noreply.github.com> Date: Fri Jul 24 03:54:47 2026 -0700 cohere2 moe template parser: enforce JSON schema for text responses if a response schema is provided (ggml-org#26018) commit 0cea362 Author: Xuan-Son Nguyen <son@huggingface.co> Date: Fri Jul 24 08:02:23 2026 +0200 vendor: update subprocess.h (ggml-org#26061) commit 40f2036 Author: Anirban Kar <thecodacus@gmail.com> Date: Sun Jul 19 04:34:24 2026 +0530 feat: fused turbo MMA flash-attention decode (turbo4/3/2, GQA tensor-core path) (TheTom#4) * feat(fattn): fused turbo4 MMA decode path (opt-in) — based on the spiritbuun fork MMA approach, adapted to graph-rotated Q Port buun's fused turbo4 tensor-core (MMA) flash-attention DECODE kernel into the turbo fork as an OPT-IN path (GGML_TURBO_MMA_FUSED=1, default OFF), routing turbo4-K==turbo4-V, D in {128,256}, decode (Q->ne[1]<=4) onto the GQA-packed MMA path so KV is read once per head-group instead of once per query head. How it works: - fattn-mma-f16.cuh: add flash_attn_ext_turbo4_load_tile (dequant turbo4 blocks -> SRAM in the f16 row-major half2 layout) seeded with OUR Lloyd-Max centroids (turbo-quant.cuh:297, NOT buun's -0.2415xx table). Thread type_K/type_V through iter/process_tile/global flash_attn_ext_f16 (defaults F16 -> existing f16/q8 MMA byte-identical). For turbo: force nstages=0 (the cp.async multi-stage path would read raw turbo bytes as half2), pass RAW byte pitch nb11/nb21 (not /sizeof(half2)), and dequantize K/V in the load tiles. sizeof(block_turbo4_0)-driven pointer math (66B block, no rnorm) — never assumes a fixed offset. - fattn-mma-turbo.cuh + 14 instance files (7 reachable ncols pairs x D in {128,256}): host launcher reusing the f16 device kernel with type_K/type_V=TURBO4_0 and need_f16_K/V=false so launch_fattn does NOT pre-convert KV to f16. - fattn.cu: switch_ncols2 -> per-ncols2 dispatchers (only the 7 compiled instances), env latch, and the gate (Q untouched — our fork already rotates Q at the graph level, so NO inline FWHT / src swap, which would double-rotate Q). Results (Qwen3.6-35B-A3B turbo4, GQA8, D=256, 5090): depth MMA VEC+PDL buun 8192 213.8 188.5 200 (+6.9% vs buun) 16384 207.6 171.7 193 (+7.5% vs buun) 32768 191.3 141.1 187 (+2.3% vs buun) MMA holds flat at depth where VEC collapses; the gap vs VEC grows with KV length. Quality: Mean KLD vs f16 base = 0.008396 (== VEC baseline). The MMA path is NOT strictly token-identical to VEC — MMA and VEC accumulate the P*V (VKQ) reduction in f16 with different reduction trees, so a near-tie greedy token can flip (~1/25 on a hard tie; same irreducible f16-order effect that exists between base f16-MMA and f16-VEC, and within the VEC-vs-base "same-top-p" 96.3% noise floor). Because strict token-identity does not hold, the gate DEFAULTS OFF and ships as opt-in; VEC remains the untouched default/kill-switch. (cherry picked from commit 5450669510a5cd0a356a863923b057f2552c6fbb) (cherry picked from commit b3e51cf) * turbo4 MMA decode: default ON (GQA-packed tensor-core path, quality-neutral; GGML_TURBO_MMA_FUSED=0 kill-switch) (cherry picked from commit 726234b538194c644be4066e1413b91f9edf12e7) (cherry picked from commit 38dc665) * fattn: extend fused MMA decode to turbo3 + turbo2 (parity on turbo4; extends the MMA path to turbo3/2) The GQA-packed tensor-core decode path was turbo4-only; turbo3/turbo2 fell back to VEC and collapsed at depth (turbo3 tg32 224->124 from d2048->d32768), losing to spiritbuun there. Added flash_attn_ext_turbo3_load_tile (3-bit split index: 2 low bits from qs + 1 high bit from signs) and flash_attn_ext_turbo2_load_tile (plain 2-bit), type-dispatched both the K and V MMA load sites, and routed the gate by K->type. is_turbo_kv already matches any non-f16 KV, so the kernel needed only the per-type tile loaders + instances (14 instance files now emit all three types). Correctness: MMA decode is bit-exact to VEC (Mean KLD @2048 chunks=8: turbo3 0.020278==0.020278, turbo2 0.041490==0.041490, turbo4 0.009296 unchanged). Decode tg32 (qwen3.6-35B-A3B, RTX 5090), MMA now flat at depth: turbo3 2k/8k/16k/32k = 224/225/217/203 (buun 193/193/183/178; was VEC 224/187/160/124) turbo2 2k/8k/16k/32k = 228/224/217/203 (buun 213/210/204/200; was VEC 229/218/203/183) turbo3/turbo2 now hold flat at depth (neither fork had them on the MMA path before). GGML_TURBO_MMA_FUSED=0 restores VEC. (cherry picked from commit 57ad5edd53a7fced0b99fecb99326132cbdb65d0) (cherry picked from commit 4e223ee) * fix(fattn): gate turbo2 fused-MMA decode to head_dim 128 Tester @everson isolated a turbo2-only decode regression on RTX 5060 Ti (Gemma-12B, head_dim 256): turbo2/turbo2 -2.46% at 8K to -1.17% at 200K, fully restored by GGML_TURBO_MMA_FUSED=0 at every depth, while turbo3 was neutral and turbo4 exact parity. At 2-bit KV the fused path's GQA-pack saving is tiny while its dequant/no-pipeline overhead is unchanged, so on head_dim 256 it is neutral on high-bandwidth GPUs (Ornith-35B and a dense Gemma-4-12B on RTX 5090, both within noise) and only costs on bandwidth-limited cards. turbo2 fused remains a large depth win on dense head_dim 128 (Llama-3.1-8B on RTX 5090: +7% at 8K, +24% at 32K, +69% at 131K), so this gates turbo2 fused to head_dim 128 only; head_dim 256 turbo2 falls through to VEC, which is the exact baseline path GGML_TURBO_MMA_FUSED=0 restored. turbo3 and turbo4 fused are unchanged at both head dims. Also corrects the stale "default OFF" gate comments to match the default-ON code. (cherry picked from commit 539ce5d) --------- Co-authored-by: TheTom <tturney1@gmail.com> commit 4cd54c6 Author: Anirban Kar <thecodacus@gmail.com> Date: Sun Jul 19 04:14:41 2026 +0530 fix: allow mixed turbo/q8_0 KV cache types in CUDA flash-attention selection (TheTom#3) * fix: allow mixed turbo/q8_0 KV cache types in CUDA flash-attention selection Without GGML_CUDA_FA_ALL_QUANTS (the default), kernel selection rejected any K/V type mismatch. A turbo KV setup is almost always mixed (ctk q8_0 + ctv turbo2/3), so the flash-attention probe failed, FA was auto-disabled, and context creation either refused the quantized V cache outright or fell back to the non-FA path whose KQ buffer (~8.7 GiB at 64k ctx, ubatch 1024) OOMs at load. Allow exactly the mixed pairs that have FA template instances compiled in: turbo2/3/4 against each other, q8_0, and f16. Also gate turbo head-dim geometry (multiples of 64; 128 for turbo4) so shapes without a compiled instance report unsupported instead of asserting at dispatch. Forcing -fa on masked the bug (the probe is skipped); server deployments use fa auto and hit it on every load. * feat: accept turbo2/3/4 cache type names in llama-bench The common arg parser already accepts them for llama-completion and llama-server; llama-bench has its own name table and rejected -ctk/-ctv turbo values. commit 9338207 Author: Anirban Kar <thecodacus@gmail.com> Date: Sun Jul 19 01:59:58 2026 +0530 CUDA: add Q2_0 support (TheTom#2) Co-authored-by: Pasha Khosravi <khosravipasha@gmail.com> commit 170322f Author: Anirban Kar <thecodacus@gmail.com> Date: Sat Jul 18 15:27:13 2026 +0530 feat: TurboQuant KV-cache quantization — all backends (TheTom#1) * feat: TurboQuant KV-cache quantization — core types + CPU backend Port of the TurboQuant codec (arXiv 2504.19874, reference: TheTom/llama-cpp-turboquant) onto this fork's base: WHT-rotated polar quantization for the KV cache. - new types TURBO2_0/TURBO3_0/TURBO4_0 (KV, 2/3/4-bit blocks of 128) and TQ3_1S/TQ4_1S (WHT-rotated Lloyd-Max weight formats), renumbered for this tree's type table - GGML_OP_TURBO_WHT + ggml_turbo_wht() builder (forward/inverse Walsh-Hadamard rotation with sign arrays, InnerQ scale hook) - ggml-turbo-quant.c codec core, rotation data tables, CPU op + vec_dot paths, type traits, quantize-chunk wiring - CUDA/Metal/Vulkan kernel sources imported (wired in follow-ups) - verified: test-turbo-quant round-trip on CPU (cosine 0.986-0.996) * feat: turbo KV cache — graph wiring and CLI types - pre-rotate Q (forward WHT) at every cache-backed attention site when K is a turbo type; no-op otherwise - inverse WHT on attention output in both FA and non-FA paths of build_attn_mha (group size derived from K, contiguity enforced) - register turbo2/3/4 as --cache-type-k/--cache-type-v values Core scope: no InnerQ scaling, no head-dim padding — turbo KV requires 128-aligned (or 64-group) head dims * feat: turbo KV cache — Metal backend (WHT kernel + FA dispatch) Wire the TurboQuant Metal backend host side and kernels on top of the CPU core already on this branch. Metal kernels (ggml-metal.metal): - Inlined turbo-wht.h (WHT sign arrays incl. half4 variants + FWHT butterfly + rotate fwd/inverse) and the turbo2/3/4 centroids. - quantize/dequantize_turbo2_0/3_0/4_0 (+_t4) and full-block cache. - kernel_turbo_wht (O(d log d) Walsh-Hadamard rotation op) and kernel_turbo4_dequant_f16 bulk prefill dequant. - kernel_set_rows_turbo/2/4 cache-quantize kernels + instantiations. - flash_attn_ext / flash_attn_ext_vec instantiations for symmetric and asymmetric turbo2/3/4 and q8_0<->turbo pairs. - TurboFlash two-pass kernels (p1/p2), default-disabled via env. Host side: - ggml-metal-device.cpp: get_pipeline_turbo_wht; flash-attn pipeline base names now encode both K and V types (k%s_v%s) so asymmetric K/V pairs resolve unambiguously. All existing FA kernel host_names in the shader renamed to match. - ggml-metal-device.m: TurboQuant compile-time env prep (4-mag LUT auto select, sparse-V, profile mode); supports_op accepts turbo K types and asymmetric turbo/q8_0 FA pairs; GGML_OP_TURBO_WHT. - ggml-metal-ops.cpp: GGML_OP_TURBO_WHT encode case + impl; asymmetric K/V assert relaxation; TurboFlash gating (use_turbo_flash, extra_tmp sizing, force-nonvec) + two-pass dispatch; +128 float FA smem for the turbo block cache. - ggml-metal-impl.h / -device.h / -ops.h: kargs (turbo_wht, turbo_flash_p1/p2), FC offsets, pipeline/op declarations. Adaptations vs the reference tree (older upstream vintage): - Reference intertwines separate TQ3_1S/TQ4_1S weight-quant work; only the turbo KV-cache lines are ported here. - set_rows_turbo dispatch selection is not wired: our upstream refactored get_pipeline_set_rows/kernel_set_rows_q32 to a different signature, and it is out of scope for this KV port. supports_op therefore does not claim CPY/SET_ROWS turbo support, so cache writes fall back to CPU; the set_rows_turbo shader kernels are present for when dispatch lands. - No InnerQ scale tensors touched (graph passes nullptr scale). * feat: turbo KV cache — CUDA backend (set-rows quantize, fattn vec dispatch, dequant) Port the TurboQuant CUDA backend from the reference implementation, adapted to this newer upstream vintage: - set-rows.cu: k_set_rows_turbo2/3/4 kernels (WHT rotation + norm-corrected PolarQuant KV write, InnerQ calibration hooks) + dispatch. turbo4 rnorm write dropped (field removed from block_turbo4_0 on this branch). - dequantize.cuh: dequantize_turbo2/3/4_0 and dequantize_tq3_1s/tq4_1s device functions. - convert.cu: turbo + TQ dequant cases in the fp16/fp32 contiguous and non-contiguous dispatch tables; warp-cooperative TQ4_1S dequant kernel. - fattn-common.cuh: vec_dot_fattn_vec_KQ_turbo2/3/4_0, dequantize_V_turbo* and get_vec_dot_KQ / get_dequantize_V dispatch entries. - fattn-vec.cuh: turbo K types take the float-Q path (Q_q8_1 off), turbo V uses the ne=4 quantized dequant path; extern decls for the 21 compiled K/V type combinations. Reference-side perf extras (shared-mem centroid LUT, sparse-V skip, nthreads tuning) intentionally not ported. - fattn.cu: turbo types accepted in kv_type_supported, vec kernel cases registered. Decode goes through the VEC kernel; prefill converts KV to f16 via the ported nc dequant path. - fattn-mma-turbo.cuh: reduced to a stub launcher. No typed MMA kernel exists in the reference either; the pre-existing template instances now compile against the stub and the dispatcher never selects this path. - ggml-cuda.cu: GGML_OP_TURBO_WHT dispatch + supports_op, SET_ROWS turbo gates (head_dim % 64, turbo4 % 128). - CMakeLists.txt: turbo fattn-vec instances added to the non-FA_ALL_QUANTS build. Verified on RTX 3060 (sm_86) with Bonsai-27B-Q1_0, partial offload, fa on: coherent generation for turbo2/3/4; perplexity over 4 chunks: f16 KV 7.9486 +/- 0.73, turbo3 KV 8.0261 +/- 0.74 (CUDA) vs 8.0279 (CPU-only). * feat: turbo KV cache — Vulkan backend (set-rows quantize, turbo3 FA dispatch, WHT op) - register turbo_wht.comp and set_rows turbo2/3/4 quantize variants in vulkan-shaders-gen; add dedicated DATA_A_TURBO3_0 flash-attn SPIR-V (scalar + coopmat1) with turbo3-only K/V bindings - copy_to_quant: 128-thread per-block set_rows path (normalize, randomized WHT, Lloyd-Max centroid quantize, subgroup pack), shared WHT tables across the three turbo widths - flash_attn_dequant: turbo2/3/4 centroid dequant macros wired into the FaTypeK/FaTypeV runtime switch (ids 43/44/45); f16 K/V aliases guarded out under DATA_A_TURBO3_0 - ggml-vulkan: GGML_OP_TURBO_WHT pipeline + dispatch, turbo set_rows pipelines (required subgroup size 32), FA pipeline selection for turbo3, supports_op gates (128-elem blocks, turbo3 K/V symmetry, scalar/coopmat1 only — cm2 falls back to scalar) * fix: extern initialization warning and C linkage for wht group size CI builds with -Werror: defining the GGML_API-decorated global with an initializer warns on every platform. Plain definition in the C TU; the C++ consumer now declares it extern "C" (it previously referenced a mangled symbol). * feat: turbo2/turbo4 FA dequant for Vulkan - generalize the dedicated turbo flash-attn SPIR-V path from turbo3-only to all three tiers: DATA_A_TURBO_KV umbrella define in types.glsl gates out the f16/q4/q5/q8 K/V binding aliases in flash_attn.comp, flash_attn_cm1.comp and flash_attn_dequant.glsl for any turbo tier - flash_attn_dequant: per-compilation dequantize4 for DATA_A_TURBO2_0 (2-bit centroids, fp16 norm + 32B idx) and DATA_A_TURBO4_0 (4-bit nibble idx, fp16 norm + fp16 aux + 64B) alongside the turbo3 one - vulkan-shaders-gen: register flash_attn_f32_f16_turbo{2,4}_0 scalar + coopmat1 variants via a per-tier loop - ggml-vulkan: FA pipeline selection picks the matching turbo SPIR-V for turbo2/turbo4 k_type (scalar + cm1), supports_op requires symmetric K/V per tier (turbo2/turbo2, turbo3/turbo3, turbo4/turbo4) * fix: bump RPC proto patch version for GGML_OP_TURBO_WHT * fix: export WHT group size via GGML_API for Windows shared builds * fix: HIP/MUSA portability for turbo kernels - map cudaMemcpyToSymbol/FromSymbol and __ballot_sync in vendor headers - pass explicit width to __shfl*_sync (HIP macros take 4 args) - drop unused warp_id locals flagged by nvcc -Werror * fix: drop SYCL turbo fattn-vec instances (no turbo dequant support in SYCL fattn) * fix: size vec_dot buffer by vec_dot_type and handle turbo/tq types in quantize test - tmp_q2 overflowed for types with vec_dot_type F32 (4 bytes per element) - skip turbo KV types: they quantize in the rotated domain, round trip does not reproduce the input by design - tq3_1s uses the 3-bit and lowbit thresholds * fix: add turbo fattn-vec instances to HIP and MUSA source lists * fix: add turbo fattn vec kernels to HIP VGPR ignore list The gfx908 build of the ncols=2 turbo3 vec kernels lands at 259 total VGPRs, slightly over the 256 gate. Register them like the other known high-pressure fattn kernels. commit 0888217 Merge: 5e7f627 7d56da7 Author: Anirban Kar <thecodacus@gmail.com> Date: Fri Jul 17 20:02:48 2026 +0530 chore: sync perf with upstream master commit 5e7f627 Author: Anirban Kar <thecodacus@gmail.com> Date: Wed Jul 8 23:56:46 2026 +0530 docs(readme): add usage + benchmark instructions for the MoE-offload optimizations commit 5f83fbb Author: thecodacus <thecodacus@gmail.com> Date: Thu Jul 2 12:42:55 2026 +0000 ggml : size prefetch slots per layer and fix fallback use-after-free Profiling showed the 2-slot rotation stalls ~5.5ms per layer waiting for the down_exps upload: 3 tensors per MoE layer need 3 slots for a full layer of lookahead. Default is now 3 (one layer), configurable via the env var value, degrading to however many slots fit on OOM. Also fixes a use-after-free: on a failed slot regrow the repointed staging tensors kept dangling pointers into freed device memory, which graph reuse carried into later evals. Staging repoints are now restored right after kernel launch, slots are sized once from the graph max, and allocation happens before freeing. Qwen3.6-35B-A3B pp2048 on RTX 3060: 1643 -> 1880 t/s (mainline: 1143). commit 1163cb3 Author: thecodacus <thecodacus@gmail.com> Date: Thu Jul 2 12:16:39 2026 +0000 ggml : overlap offloaded expert weight uploads with compute At large batch sizes virtually every expert is used, so the per-layer routing-ids readback that mainline waits on buys nothing while forcing a full device sync per expert tensor (3x per MoE layer). Above a batch threshold, upload the full expert tensors through a second backend instance on the same device (own stream) with two event-ordered staging slots, so uploads for tensor N+1 overlap compute of tensor N. Qwen3.6-35B-A3B pp2048 on RTX 3060 (-ncmoe 26, ub 2048): 1383 -> 1663 t/s. Generation output verified token-identical; decode path unaffected. Opt-in via GGML_SCHED_PREFETCH_EXPERTS=1. commit 20f5994 Author: thecodacus <thecodacus@gmail.com> Date: Thu Jul 2 09:44:53 2026 +0000 llama : pin mmap-backed CPU weights for faster H2D uploads Wire the existing GGML_CUDA_REGISTER_HOST path back up: after model load, cudaHostRegister the mmap pages backing weights kept in system memory. Recovers pageable-copy losses when MoE experts are streamed to the GPU during prefill (n-cpu-moe): Qwen3.6-35B-A3B pp2048 1144 -> 1385 t/s on RTX 3060. Opt-in via GGML_CUDA_REGISTER_HOST=1, unchanged otherwise.

@2048

commit 10f8aff Author: thecodacus <thecodacus@gmail.com> Date: Tue Aug 25 23:39:14 2026 +0530 perf: lower default --cpu-tp-margin to 512 MiB Measured on a 12GB RTX 3060 with a dense 27B Q4_K_M: tightening the margin raises both decode and prefill, since every MiB reclaimed moves weights off the CPU leg. 1024 -> 640 -> 384 MiB gave 7.4 -> 7.8 -> 8.1 t/s decode and 14.4 -> 15.2 -> 16.5 t/s prefill. 512 keeps headroom for larger contexts; raise it if the GPU runs out of memory. commit e406c8c Author: thecodacus <thecodacus@gmail.com> Date: Tue Aug 25 22:08:09 2026 +0530 feat: add --cpu-tp to use the CPU as a tensor-parallel device llama.cpp's tensor-parallel meta device already dispatches each member backend asynchronously before synchronising, but the CPU is excluded from the device list. Adding it lets the GPU and CPU compute row-slices of every weight matrix concurrently instead of running one after the other. The GPU share is sized from the measured free VRAM (minus --cpu-tp-margin) so VRAM is saturated, and the CPU takes only the remainder. The CPU is appended last so the GPU subgraph is already in flight when the CPU backend blocks on its own. commit d927e7d Author: Anirban Kar <thecodacus@gmail.com> Date: Tue Aug 25 15:58:16 2026 +0530 docs: add MoE expert cache usage and tuning guide to README (TheTom#8) * docs: add MoE expert cache usage and tuning guide to README * docs: document async CPU split overlap and VRAM headroom rule commit 0ac3d9b Author: Anirban Kar <thecodacus@gmail.com> Date: Sun Aug 23 04:08:21 2026 +0530 feat: overlap CPU cold-expert and GPU hot-expert computation in the MoE cache (TheTom#9) * feat: overlap CPU cold-expert chain with GPU hot chain (GGML_SCHED_ASYNC_CPU) The dual-path MoE cache ran its two chains serially: the scheduler blocks on CPU splits, and the cold chain's input copies drain the whole CUDA stream. Three changes make them concurrent, gated behind GGML_SCHED_ASYNC_CPU=1: - sched: CPU splits execute on a persistent worker thread; the main loop keeps launching later splits and joins lazily at the first split that is CPU or reads a CPU-resident input - graph: cold chain is built before the hot chain so the hot split follows it and carries no cross-backend inputs (launches while the worker computes) - graph: the merge add is pinned to CPU so it does not pull the hot chain into a split that waits on the cold result * feat: GGML_SCHED_ASYNC_CPU=2 diagnostic mode (worker without graph restructure) * fix: mark expert cache pack buffer as weights Without the weights usage flag the scheduler's assignment pass does not pin the pack tensors to their backend, so a CPU-assigned consumer (the pinned merge in async mode) propagates backward and pulls the hot matmuls onto the CPU, copying pack weights over PCIe every layer. * feat: make async CPU splits the default, expose --sched-async-cpu flag Replaces the GGML_SCHED_ASYNC_CPU env var: the overlap is now on by default and controlled through llama_context_params.sched_async_cpu. - ggml: ggml_backend_sched_set_async_cpu() setter, env read removed - llama: context param + cparams plumbing, graph gate reads cparams - common: --sched-async-cpu / --no-sched-async-cpu (server, cli) - llama-bench: sched_async_cpu test dimension (--sched-async-cpu 0,1 benches both in one run, column shown as 'sac' when non-default) * fix: restrict CPU-pinned merge to decode-size batches For prefill batches the CPU merge and its per-layer activation copies cost more than the overlap hides (pp256 regressed 307 -> 257 t/s). Decode and MTP verify batches (n_tokens <= 8) keep the overlap. commit 8e5bc28 Author: Anirban Kar <thecodacus@gmail.com> Date: Sat Jul 25 03:54:47 2026 +0530 feat: wire MoE expert cache into laguna graph (TheTom#6) commit ac743f8 Author: Anirban Kar <thecodacus@gmail.com> Date: Sat Jul 25 00:41:02 2026 +0530 feat: MoE expert cache — keep hot routed experts resident in VRAM (TheTom#5) * feat: llama-moe-trace tool + expert-cache policy simulator Phase 0 of the MoE expert-cache work: capture routed expert ids (ffn_moe_topk) per layer per decode step via the scheduler eval callback — backend-agnostic, so CPU-offloaded expert layers are traced too. Companion simulate.py replays traces against cache policies (static top-S, LRU, LFU-decay, periodic re-election) and prints hit-rate vs VRAM-budget curves plus routing-skew stats to decide whether expert caching is worth building at all. * fix: moe-trace — common_init_from_params returns ptr with accessors * fix: moe-trace — include arg.h for common_params_parse * fix: moe-trace — take output path from MOE_TRACE_OUT env, -o not in this arg set * fix: moe-trace — topk tensor is a strided view; copy full byte range and index by nb Sizing the host buffer as n_used*n_tokens under-allocated: ffn_moe_topk is a view over the argsort output, so nb[1] spans n_expert ids per row and ggml_backend_tensor_get writes the whole strided range (heap corruption, malloc corrupted-top-size abort mid-prefill). * feat: moe-trace simulator v2 — model upload bytes/token and miss/upload cost per policy * feat: mul_mat_id tolerates id=-1 (skip) on CPU and CUDA fallback path Groundwork for the hot/cold expert-pack split: an id of -1 means the routed expert is not owned by this pack. The op contributes a zeroed output row for that slot so hot-pack and cold-pack results merge additively. CUDA fast paths (mmvq_id/mmq_id/mmf_id) still assert non-negative ids and are handled next. * feat: id=-1 skip support in CUDA mul_mat_id fast paths + tests mmid.cu gains a generic kernel that zeroes the dst rows of (token, slot) pairs whose expert id is -1; mmq and mmf launch it after the shared mm_ids_helper (which already drops unmatched ids from its compact mapping), and mmvq launches it before its kernels, which now early-return instead of using a negative id as a channel index (out of bounds read otherwise). test-backend-ops gains skip_ids variants of MUL_MAT_ID covering the mmvq/mmq/mmf/fallback dispatch corners. * fix: skip_ids MUL_MAT_ID test cases were inside an #if 0 block * fix: mm_ids_helper compact map corruption when ids contain -1 Negative ids inflated nex_prev for every expert, opening a phantom region at the head of the compact arrays that nothing wrote; the mmq quantize preprocessing iterates the flat range and read garbage row indices (illegal memory access). Skipped ids now occupy no compact position, ids_src1 is sentinel-filled, and the quantize gather/scatter kernels skip negative entries. * fix: mmvf id=-1 skip guard (float mul_mat_id vec path) Same treatment as mmvq: kernels early-return on a negative expert id instead of using it as a channel index, and the host launcher pre-zeroes the dst rows of skipped slots. Fixes wrong results (and a latent out-of-bounds channel read) for f16/f32 MUL_MAT_ID with -1 ids at small batch sizes. * fix: zero-init ids_dst before mm_ids_helper — garbage tail entries in the compact map caused pool-layout-dependent illegal memory accesses when ids contain -1 (tile-padded reads consumed uninitialized row indices) * feat: route MUL_MAT_ID with skip-capable ids away from mmq/mmf Ops whose ids may contain -1 announce it via op_params[0]; the CUDA dispatch then uses only the paths with verified skip support (mmvq for small batches, the general gather/scatter path otherwise). Large-batch mmq with sparse ids shows a pool-layout-dependent fault that resists isolation (all indexing audits pass; sanitizer and cuda-gdb unusable on the target driver); since the expert-pack design only decodes through packs, the safe route costs nothing in practice. mmq/mmf guards from earlier commits remain as hardening. * fix: op_params write in test — ggml_set_op_params_i32 is internal-only * feat: MoE expert cache loader — hot-pack build from routing profile init_moe_expert_cache() runs post-load: reads a moe-trace profile (GGML_MOE_CACHE_PROFILE) and slot count (GGML_MOE_CACHE_SLOTS), and for each CPU-resident MoE layer builds GPU-resident hot packs of the top-S experts plus i32 remap tables (hot: global->slot or -1; cold: global id or -1). Graph wiring (build_moe_ffn dual path) is the next commit; these tensors are unused until then, so this is behavior-neutral. * feat: MoE expert cache graph wiring — dual hot/cold mul_mat_id in build_moe_ffn When a layer carries hot packs, routed ids are remapped through moe_map_hot/cold (get_rows) and each expert projection runs as two skip-flagged mul_mat_id ops — hot pack on GPU, original tensors for cold — whose zero-padded outputs sum to the exact single-tensor result. Engaged for qwen35moe (separate gate/up/down, no expert scales); layers with merged gate_up, scales, or no packs use the unchanged path. * fix: flatten routed ids before map get_rows (batched-gather shape rule) * fix: topk ids are a strided view - cont before flatten for map lookup * fix: sched expert-copy path skips id=-1 (hot/cold expert split) * fix: restructure MoE pack dual path into full FFN chains per residency side Per-projection hot/cold interleave let the scheduler's split minimizer assign hot up/down matmuls to CPU, copying the GPU pack weights over PCIe every layer. Building one unbroken gate-up-swiglu-down chain per side keeps the hot pack on GPU; skipped (-1) rows stay zero through swiglu so a single add of the two chain outputs is exact. * feat: wire MoE expert cache into deepseek2 graph (GLM-4.7-Flash) * feat: add --moe-cache-profile and --moe-cache-slots flags Plumbed through llama_model_params; env vars remain as fallback so existing deployments keep working. Flags allow per-model settings in multi-model serving. * docs: fix moe-trace usage comment (output path is MOE_TRACE_OUT env) commit cd72c28 Author: Anirban Kar <thecodacus@gmail.com> Date: Fri Jul 24 17:53:41 2026 +0530 fix: restore closing brace lost in dequantize.cuh merge resolution commit 34f893e Merge: 40f2036 8f5ab83 Author: Anirban Kar <thecodacus@gmail.com> Date: Fri Jul 24 17:52:26 2026 +0530 merge: sync upstream master into perf commit 8f5ab83 Author: Matt Thompson <111157855+boondocklabs@users.noreply.github.com> Date: Fri Jul 24 03:54:47 2026 -0700 cohere2 moe template parser: enforce JSON schema for text responses if a response schema is provided (ggml-org#26018) commit 0cea362 Author: Xuan-Son Nguyen <son@huggingface.co> Date: Fri Jul 24 08:02:23 2026 +0200 vendor: update subprocess.h (ggml-org#26061) commit 40f2036 Author: Anirban Kar <thecodacus@gmail.com> Date: Sun Jul 19 04:34:24 2026 +0530 feat: fused turbo MMA flash-attention decode (turbo4/3/2, GQA tensor-core path) (TheTom#4) * feat(fattn): fused turbo4 MMA decode path (opt-in) — based on the spiritbuun fork MMA approach, adapted to graph-rotated Q Port buun's fused turbo4 tensor-core (MMA) flash-attention DECODE kernel into the turbo fork as an OPT-IN path (GGML_TURBO_MMA_FUSED=1, default OFF), routing turbo4-K==turbo4-V, D in {128,256}, decode (Q->ne[1]<=4) onto the GQA-packed MMA path so KV is read once per head-group instead of once per query head. How it works: - fattn-mma-f16.cuh: add flash_attn_ext_turbo4_load_tile (dequant turbo4 blocks -> SRAM in the f16 row-major half2 layout) seeded with OUR Lloyd-Max centroids (turbo-quant.cuh:297, NOT buun's -0.2415xx table). Thread type_K/type_V through iter/process_tile/global flash_attn_ext_f16 (defaults F16 -> existing f16/q8 MMA byte-identical). For turbo: force nstages=0 (the cp.async multi-stage path would read raw turbo bytes as half2), pass RAW byte pitch nb11/nb21 (not /sizeof(half2)), and dequantize K/V in the load tiles. sizeof(block_turbo4_0)-driven pointer math (66B block, no rnorm) — never assumes a fixed offset. - fattn-mma-turbo.cuh + 14 instance files (7 reachable ncols pairs x D in {128,256}): host launcher reusing the f16 device kernel with type_K/type_V=TURBO4_0 and need_f16_K/V=false so launch_fattn does NOT pre-convert KV to f16. - fattn.cu: switch_ncols2 -> per-ncols2 dispatchers (only the 7 compiled instances), env latch, and the gate (Q untouched — our fork already rotates Q at the graph level, so NO inline FWHT / src swap, which would double-rotate Q). Results (Qwen3.6-35B-A3B turbo4, GQA8, D=256, 5090): depth MMA VEC+PDL buun 8192 213.8 188.5 200 (+6.9% vs buun) 16384 207.6 171.7 193 (+7.5% vs buun) 32768 191.3 141.1 187 (+2.3% vs buun) MMA holds flat at depth where VEC collapses; the gap vs VEC grows with KV length. Quality: Mean KLD vs f16 base = 0.008396 (== VEC baseline). The MMA path is NOT strictly token-identical to VEC — MMA and VEC accumulate the P*V (VKQ) reduction in f16 with different reduction trees, so a near-tie greedy token can flip (~1/25 on a hard tie; same irreducible f16-order effect that exists between base f16-MMA and f16-VEC, and within the VEC-vs-base "same-top-p" 96.3% noise floor). Because strict token-identity does not hold, the gate DEFAULTS OFF and ships as opt-in; VEC remains the untouched default/kill-switch. (cherry picked from commit 5450669510a5cd0a356a863923b057f2552c6fbb) (cherry picked from commit b3e51cf) * turbo4 MMA decode: default ON (GQA-packed tensor-core path, quality-neutral; GGML_TURBO_MMA_FUSED=0 kill-switch) (cherry picked from commit 726234b538194c644be4066e1413b91f9edf12e7) (cherry picked from commit 38dc665) * fattn: extend fused MMA decode to turbo3 + turbo2 (parity on turbo4; extends the MMA path to turbo3/2) The GQA-packed tensor-core decode path was turbo4-only; turbo3/turbo2 fell back to VEC and collapsed at depth (turbo3 tg32 224->124 from d2048->d32768), losing to spiritbuun there. Added flash_attn_ext_turbo3_load_tile (3-bit split index: 2 low bits from qs + 1 high bit from signs) and flash_attn_ext_turbo2_load_tile (plain 2-bit), type-dispatched both the K and V MMA load sites, and routed the gate by K->type. is_turbo_kv already matches any non-f16 KV, so the kernel needed only the per-type tile loaders + instances (14 instance files now emit all three types). Correctness: MMA decode is bit-exact to VEC (Mean KLD @2048 chunks=8: turbo3 0.020278==0.020278, turbo2 0.041490==0.041490, turbo4 0.009296 unchanged). Decode tg32 (qwen3.6-35B-A3B, RTX 5090), MMA now flat at depth: turbo3 2k/8k/16k/32k = 224/225/217/203 (buun 193/193/183/178; was VEC 224/187/160/124) turbo2 2k/8k/16k/32k = 228/224/217/203 (buun 213/210/204/200; was VEC 229/218/203/183) turbo3/turbo2 now hold flat at depth (neither fork had them on the MMA path before). GGML_TURBO_MMA_FUSED=0 restores VEC. (cherry picked from commit 57ad5edd53a7fced0b99fecb99326132cbdb65d0) (cherry picked from commit 4e223ee) * fix(fattn): gate turbo2 fused-MMA decode to head_dim 128 Tester @everson isolated a turbo2-only decode regression on RTX 5060 Ti (Gemma-12B, head_dim 256): turbo2/turbo2 -2.46% at 8K to -1.17% at 200K, fully restored by GGML_TURBO_MMA_FUSED=0 at every depth, while turbo3 was neutral and turbo4 exact parity. At 2-bit KV the fused path's GQA-pack saving is tiny while its dequant/no-pipeline overhead is unchanged, so on head_dim 256 it is neutral on high-bandwidth GPUs (Ornith-35B and a dense Gemma-4-12B on RTX 5090, both within noise) and only costs on bandwidth-limited cards. turbo2 fused remains a large depth win on dense head_dim 128 (Llama-3.1-8B on RTX 5090: +7% at 8K, +24% at 32K, +69% at 131K), so this gates turbo2 fused to head_dim 128 only; head_dim 256 turbo2 falls through to VEC, which is the exact baseline path GGML_TURBO_MMA_FUSED=0 restored. turbo3 and turbo4 fused are unchanged at both head dims. Also corrects the stale "default OFF" gate comments to match the default-ON code. (cherry picked from commit 539ce5d) --------- Co-authored-by: TheTom <tturney1@gmail.com> commit 4cd54c6 Author: Anirban Kar <thecodacus@gmail.com> Date: Sun Jul 19 04:14:41 2026 +0530 fix: allow mixed turbo/q8_0 KV cache types in CUDA flash-attention selection (TheTom#3) * fix: allow mixed turbo/q8_0 KV cache types in CUDA flash-attention selection Without GGML_CUDA_FA_ALL_QUANTS (the default), kernel selection rejected any K/V type mismatch. A turbo KV setup is almost always mixed (ctk q8_0 + ctv turbo2/3), so the flash-attention probe failed, FA was auto-disabled, and context creation either refused the quantized V cache outright or fell back to the non-FA path whose KQ buffer (~8.7 GiB at 64k ctx, ubatch 1024) OOMs at load. Allow exactly the mixed pairs that have FA template instances compiled in: turbo2/3/4 against each other, q8_0, and f16. Also gate turbo head-dim geometry (multiples of 64; 128 for turbo4) so shapes without a compiled instance report unsupported instead of asserting at dispatch. Forcing -fa on masked the bug (the probe is skipped); server deployments use fa auto and hit it on every load. * feat: accept turbo2/3/4 cache type names in llama-bench The common arg parser already accepts them for llama-completion and llama-server; llama-bench has its own name table and rejected -ctk/-ctv turbo values. commit 9338207 Author: Anirban Kar <thecodacus@gmail.com> Date: Sun Jul 19 01:59:58 2026 +0530 CUDA: add Q2_0 support (TheTom#2) Co-authored-by: Pasha Khosravi <khosravipasha@gmail.com> commit 170322f Author: Anirban Kar <thecodacus@gmail.com> Date: Sat Jul 18 15:27:13 2026 +0530 feat: TurboQuant KV-cache quantization — all backends (TheTom#1) * feat: TurboQuant KV-cache quantization — core types + CPU backend Port of the TurboQuant codec (arXiv 2504.19874, reference: TheTom/llama-cpp-turboquant) onto this fork's base: WHT-rotated polar quantization for the KV cache. - new types TURBO2_0/TURBO3_0/TURBO4_0 (KV, 2/3/4-bit blocks of 128) and TQ3_1S/TQ4_1S (WHT-rotated Lloyd-Max weight formats), renumbered for this tree's type table - GGML_OP_TURBO_WHT + ggml_turbo_wht() builder (forward/inverse Walsh-Hadamard rotation with sign arrays, InnerQ scale hook) - ggml-turbo-quant.c codec core, rotation data tables, CPU op + vec_dot paths, type traits, quantize-chunk wiring - CUDA/Metal/Vulkan kernel sources imported (wired in follow-ups) - verified: test-turbo-quant round-trip on CPU (cosine 0.986-0.996) * feat: turbo KV cache — graph wiring and CLI types - pre-rotate Q (forward WHT) at every cache-backed attention site when K is a turbo type; no-op otherwise - inverse WHT on attention output in both FA and non-FA paths of build_attn_mha (group size derived from K, contiguity enforced) - register turbo2/3/4 as --cache-type-k/--cache-type-v values Core scope: no InnerQ scaling, no head-dim padding — turbo KV requires 128-aligned (or 64-group) head dims * feat: turbo KV cache — Metal backend (WHT kernel + FA dispatch) Wire the TurboQuant Metal backend host side and kernels on top of the CPU core already on this branch. Metal kernels (ggml-metal.metal): - Inlined turbo-wht.h (WHT sign arrays incl. half4 variants + FWHT butterfly + rotate fwd/inverse) and the turbo2/3/4 centroids. - quantize/dequantize_turbo2_0/3_0/4_0 (+_t4) and full-block cache. - kernel_turbo_wht (O(d log d) Walsh-Hadamard rotation op) and kernel_turbo4_dequant_f16 bulk prefill dequant. - kernel_set_rows_turbo/2/4 cache-quantize kernels + instantiations. - flash_attn_ext / flash_attn_ext_vec instantiations for symmetric and asymmetric turbo2/3/4 and q8_0<->turbo pairs. - TurboFlash two-pass kernels (p1/p2), default-disabled via env. Host side: - ggml-metal-device.cpp: get_pipeline_turbo_wht; flash-attn pipeline base names now encode both K and V types (k%s_v%s) so asymmetric K/V pairs resolve unambiguously. All existing FA kernel host_names in the shader renamed to match. - ggml-metal-device.m: TurboQuant compile-time env prep (4-mag LUT auto select, sparse-V, profile mode); supports_op accepts turbo K types and asymmetric turbo/q8_0 FA pairs; GGML_OP_TURBO_WHT. - ggml-metal-ops.cpp: GGML_OP_TURBO_WHT encode case + impl; asymmetric K/V assert relaxation; TurboFlash gating (use_turbo_flash, extra_tmp sizing, force-nonvec) + two-pass dispatch; +128 float FA smem for the turbo block cache. - ggml-metal-impl.h / -device.h / -ops.h: kargs (turbo_wht, turbo_flash_p1/p2), FC offsets, pipeline/op declarations. Adaptations vs the reference tree (older upstream vintage): - Reference intertwines separate TQ3_1S/TQ4_1S weight-quant work; only the turbo KV-cache lines are ported here. - set_rows_turbo dispatch selection is not wired: our upstream refactored get_pipeline_set_rows/kernel_set_rows_q32 to a different signature, and it is out of scope for this KV port. supports_op therefore does not claim CPY/SET_ROWS turbo support, so cache writes fall back to CPU; the set_rows_turbo shader kernels are present for when dispatch lands. - No InnerQ scale tensors touched (graph passes nullptr scale). * feat: turbo KV cache — CUDA backend (set-rows quantize, fattn vec dispatch, dequant) Port the TurboQuant CUDA backend from the reference implementation, adapted to this newer upstream vintage: - set-rows.cu: k_set_rows_turbo2/3/4 kernels (WHT rotation + norm-corrected PolarQuant KV write, InnerQ calibration hooks) + dispatch. turbo4 rnorm write dropped (field removed from block_turbo4_0 on this branch). - dequantize.cuh: dequantize_turbo2/3/4_0 and dequantize_tq3_1s/tq4_1s device functions. - convert.cu: turbo + TQ dequant cases in the fp16/fp32 contiguous and non-contiguous dispatch tables; warp-cooperative TQ4_1S dequant kernel. - fattn-common.cuh: vec_dot_fattn_vec_KQ_turbo2/3/4_0, dequantize_V_turbo* and get_vec_dot_KQ / get_dequantize_V dispatch entries. - fattn-vec.cuh: turbo K types take the float-Q path (Q_q8_1 off), turbo V uses the ne=4 quantized dequant path; extern decls for the 21 compiled K/V type combinations. Reference-side perf extras (shared-mem centroid LUT, sparse-V skip, nthreads tuning) intentionally not ported. - fattn.cu: turbo types accepted in kv_type_supported, vec kernel cases registered. Decode goes through the VEC kernel; prefill converts KV to f16 via the ported nc dequant path. - fattn-mma-turbo.cuh: reduced to a stub launcher. No typed MMA kernel exists in the reference either; the pre-existing template instances now compile against the stub and the dispatcher never selects this path. - ggml-cuda.cu: GGML_OP_TURBO_WHT dispatch + supports_op, SET_ROWS turbo gates (head_dim % 64, turbo4 % 128). - CMakeLists.txt: turbo fattn-vec instances added to the non-FA_ALL_QUANTS build. Verified on RTX 3060 (sm_86) with Bonsai-27B-Q1_0, partial offload, fa on: coherent generation for turbo2/3/4; perplexity over 4 chunks: f16 KV 7.9486 +/- 0.73, turbo3 KV 8.0261 +/- 0.74 (CUDA) vs 8.0279 (CPU-only). * feat: turbo KV cache — Vulkan backend (set-rows quantize, turbo3 FA dispatch, WHT op) - register turbo_wht.comp and set_rows turbo2/3/4 quantize variants in vulkan-shaders-gen; add dedicated DATA_A_TURBO3_0 flash-attn SPIR-V (scalar + coopmat1) with turbo3-only K/V bindings - copy_to_quant: 128-thread per-block set_rows path (normalize, randomized WHT, Lloyd-Max centroid quantize, subgroup pack), shared WHT tables across the three turbo widths - flash_attn_dequant: turbo2/3/4 centroid dequant macros wired into the FaTypeK/FaTypeV runtime switch (ids 43/44/45); f16 K/V aliases guarded out under DATA_A_TURBO3_0 - ggml-vulkan: GGML_OP_TURBO_WHT pipeline + dispatch, turbo set_rows pipelines (required subgroup size 32), FA pipeline selection for turbo3, supports_op gates (128-elem blocks, turbo3 K/V symmetry, scalar/coopmat1 only — cm2 falls back to scalar) * fix: extern initialization warning and C linkage for wht group size CI builds with -Werror: defining the GGML_API-decorated global with an initializer warns on every platform. Plain definition in the C TU; the C++ consumer now declares it extern "C" (it previously referenced a mangled symbol). * feat: turbo2/turbo4 FA dequant for Vulkan - generalize the dedicated turbo flash-attn SPIR-V path from turbo3-only to all three tiers: DATA_A_TURBO_KV umbrella define in types.glsl gates out the f16/q4/q5/q8 K/V binding aliases in flash_attn.comp, flash_attn_cm1.comp and flash_attn_dequant.glsl for any turbo tier - flash_attn_dequant: per-compilation dequantize4 for DATA_A_TURBO2_0 (2-bit centroids, fp16 norm + 32B idx) and DATA_A_TURBO4_0 (4-bit nibble idx, fp16 norm + fp16 aux + 64B) alongside the turbo3 one - vulkan-shaders-gen: register flash_attn_f32_f16_turbo{2,4}_0 scalar + coopmat1 variants via a per-tier loop - ggml-vulkan: FA pipeline selection picks the matching turbo SPIR-V for turbo2/turbo4 k_type (scalar + cm1), supports_op requires symmetric K/V per tier (turbo2/turbo2, turbo3/turbo3, turbo4/turbo4) * fix: bump RPC proto patch version for GGML_OP_TURBO_WHT * fix: export WHT group size via GGML_API for Windows shared builds * fix: HIP/MUSA portability for turbo kernels - map cudaMemcpyToSymbol/FromSymbol and __ballot_sync in vendor headers - pass explicit width to __shfl*_sync (HIP macros take 4 args) - drop unused warp_id locals flagged by nvcc -Werror * fix: drop SYCL turbo fattn-vec instances (no turbo dequant support in SYCL fattn) * fix: size vec_dot buffer by vec_dot_type and handle turbo/tq types in quantize test - tmp_q2 overflowed for types with vec_dot_type F32 (4 bytes per element) - skip turbo KV types: they quantize in the rotated domain, round trip does not reproduce the input by design - tq3_1s uses the 3-bit and lowbit thresholds * fix: add turbo fattn-vec instances to HIP and MUSA source lists * fix: add turbo fattn vec kernels to HIP VGPR ignore list The gfx908 build of the ncols=2 turbo3 vec kernels lands at 259 total VGPRs, slightly over the 256 gate. Register them like the other known high-pressure fattn kernels. commit 0888217 Merge: 5e7f627 7d56da7 Author: Anirban Kar <thecodacus@gmail.com> Date: Fri Jul 17 20:02:48 2026 +0530 chore: sync perf with upstream master commit 5e7f627 Author: Anirban Kar <thecodacus@gmail.com> Date: Wed Jul 8 23:56:46 2026 +0530 docs(readme): add usage + benchmark instructions for the MoE-offload optimizations commit 5f83fbb Author: thecodacus <thecodacus@gmail.com> Date: Thu Jul 2 12:42:55 2026 +0000 ggml : size prefetch slots per layer and fix fallback use-after-free Profiling showed the 2-slot rotation stalls ~5.5ms per layer waiting for the down_exps upload: 3 tensors per MoE layer need 3 slots for a full layer of lookahead. Default is now 3 (one layer), configurable via the env var value, degrading to however many slots fit on OOM. Also fixes a use-after-free: on a failed slot regrow the repointed staging tensors kept dangling pointers into freed device memory, which graph reuse carried into later evals. Staging repoints are now restored right after kernel launch, slots are sized once from the graph max, and allocation happens before freeing. Qwen3.6-35B-A3B pp2048 on RTX 3060: 1643 -> 1880 t/s (mainline: 1143). commit 1163cb3 Author: thecodacus <thecodacus@gmail.com> Date: Thu Jul 2 12:16:39 2026 +0000 ggml : overlap offloaded expert weight uploads with compute At large batch sizes virtually every expert is used, so the per-layer routing-ids readback that mainline waits on buys nothing while forcing a full device sync per expert tensor (3x per MoE layer). Above a batch threshold, upload the full expert tensors through a second backend instance on the same device (own stream) with two event-ordered staging slots, so uploads for tensor N+1 overlap compute of tensor N. Qwen3.6-35B-A3B pp2048 on RTX 3060 (-ncmoe 26, ub 2048): 1383 -> 1663 t/s. Generation output verified token-identical; decode path unaffected. Opt-in via GGML_SCHED_PREFETCH_EXPERTS=1. commit 20f5994 Author: thecodacus <thecodacus@gmail.com> Date: Thu Jul 2 09:44:53 2026 +0000 llama : pin mmap-backed CPU weights for faster H2D uploads Wire the existing GGML_CUDA_REGISTER_HOST path back up: after model load, cudaHostRegister the mmap pages backing weights kept in system memory. Recovers pageable-copy losses when MoE experts are streamed to the GPU during prefill (n-cpu-moe): Qwen3.6-35B-A3B pp2048 1144 -> 1385 t/s on RTX 3060. Opt-in via GGML_CUDA_REGISTER_HOST=1, unchanged otherwise.

@2048

commit 10f8aff Author: thecodacus <thecodacus@gmail.com> Date: Tue Aug 25 23:39:14 2026 +0530 perf: lower default --cpu-tp-margin to 512 MiB Measured on a 12GB RTX 3060 with a dense 27B Q4_K_M: tightening the margin raises both decode and prefill, since every MiB reclaimed moves weights off the CPU leg. 1024 -> 640 -> 384 MiB gave 7.4 -> 7.8 -> 8.1 t/s decode and 14.4 -> 15.2 -> 16.5 t/s prefill. 512 keeps headroom for larger contexts; raise it if the GPU runs out of memory. commit e406c8c Author: thecodacus <thecodacus@gmail.com> Date: Tue Aug 25 22:08:09 2026 +0530 feat: add --cpu-tp to use the CPU as a tensor-parallel device llama.cpp's tensor-parallel meta device already dispatches each member backend asynchronously before synchronising, but the CPU is excluded from the device list. Adding it lets the GPU and CPU compute row-slices of every weight matrix concurrently instead of running one after the other. The GPU share is sized from the measured free VRAM (minus --cpu-tp-margin) so VRAM is saturated, and the CPU takes only the remainder. The CPU is appended last so the GPU subgraph is already in flight when the CPU backend blocks on its own. commit d927e7d Author: Anirban Kar <thecodacus@gmail.com> Date: Tue Aug 25 15:58:16 2026 +0530 docs: add MoE expert cache usage and tuning guide to README (TheTom#8) * docs: add MoE expert cache usage and tuning guide to README * docs: document async CPU split overlap and VRAM headroom rule commit 0ac3d9b Author: Anirban Kar <thecodacus@gmail.com> Date: Sun Aug 23 04:08:21 2026 +0530 feat: overlap CPU cold-expert and GPU hot-expert computation in the MoE cache (TheTom#9) * feat: overlap CPU cold-expert chain with GPU hot chain (GGML_SCHED_ASYNC_CPU) The dual-path MoE cache ran its two chains serially: the scheduler blocks on CPU splits, and the cold chain's input copies drain the whole CUDA stream. Three changes make them concurrent, gated behind GGML_SCHED_ASYNC_CPU=1: - sched: CPU splits execute on a persistent worker thread; the main loop keeps launching later splits and joins lazily at the first split that is CPU or reads a CPU-resident input - graph: cold chain is built before the hot chain so the hot split follows it and carries no cross-backend inputs (launches while the worker computes) - graph: the merge add is pinned to CPU so it does not pull the hot chain into a split that waits on the cold result * feat: GGML_SCHED_ASYNC_CPU=2 diagnostic mode (worker without graph restructure) * fix: mark expert cache pack buffer as weights Without the weights usage flag the scheduler's assignment pass does not pin the pack tensors to their backend, so a CPU-assigned consumer (the pinned merge in async mode) propagates backward and pulls the hot matmuls onto the CPU, copying pack weights over PCIe every layer. * feat: make async CPU splits the default, expose --sched-async-cpu flag Replaces the GGML_SCHED_ASYNC_CPU env var: the overlap is now on by default and controlled through llama_context_params.sched_async_cpu. - ggml: ggml_backend_sched_set_async_cpu() setter, env read removed - llama: context param + cparams plumbing, graph gate reads cparams - common: --sched-async-cpu / --no-sched-async-cpu (server, cli) - llama-bench: sched_async_cpu test dimension (--sched-async-cpu 0,1 benches both in one run, column shown as 'sac' when non-default) * fix: restrict CPU-pinned merge to decode-size batches For prefill batches the CPU merge and its per-layer activation copies cost more than the overlap hides (pp256 regressed 307 -> 257 t/s). Decode and MTP verify batches (n_tokens <= 8) keep the overlap. commit 8e5bc28 Author: Anirban Kar <thecodacus@gmail.com> Date: Sat Jul 25 03:54:47 2026 +0530 feat: wire MoE expert cache into laguna graph (TheTom#6) commit ac743f8 Author: Anirban Kar <thecodacus@gmail.com> Date: Sat Jul 25 00:41:02 2026 +0530 feat: MoE expert cache — keep hot routed experts resident in VRAM (TheTom#5) * feat: llama-moe-trace tool + expert-cache policy simulator Phase 0 of the MoE expert-cache work: capture routed expert ids (ffn_moe_topk) per layer per decode step via the scheduler eval callback — backend-agnostic, so CPU-offloaded expert layers are traced too. Companion simulate.py replays traces against cache policies (static top-S, LRU, LFU-decay, periodic re-election) and prints hit-rate vs VRAM-budget curves plus routing-skew stats to decide whether expert caching is worth building at all. * fix: moe-trace — common_init_from_params returns ptr with accessors * fix: moe-trace — include arg.h for common_params_parse * fix: moe-trace — take output path from MOE_TRACE_OUT env, -o not in this arg set * fix: moe-trace — topk tensor is a strided view; copy full byte range and index by nb Sizing the host buffer as n_used*n_tokens under-allocated: ffn_moe_topk is a view over the argsort output, so nb[1] spans n_expert ids per row and ggml_backend_tensor_get writes the whole strided range (heap corruption, malloc corrupted-top-size abort mid-prefill). * feat: moe-trace simulator v2 — model upload bytes/token and miss/upload cost per policy * feat: mul_mat_id tolerates id=-1 (skip) on CPU and CUDA fallback path Groundwork for the hot/cold expert-pack split: an id of -1 means the routed expert is not owned by this pack. The op contributes a zeroed output row for that slot so hot-pack and cold-pack results merge additively. CUDA fast paths (mmvq_id/mmq_id/mmf_id) still assert non-negative ids and are handled next. * feat: id=-1 skip support in CUDA mul_mat_id fast paths + tests mmid.cu gains a generic kernel that zeroes the dst rows of (token, slot) pairs whose expert id is -1; mmq and mmf launch it after the shared mm_ids_helper (which already drops unmatched ids from its compact mapping), and mmvq launches it before its kernels, which now early-return instead of using a negative id as a channel index (out of bounds read otherwise). test-backend-ops gains skip_ids variants of MUL_MAT_ID covering the mmvq/mmq/mmf/fallback dispatch corners. * fix: skip_ids MUL_MAT_ID test cases were inside an #if 0 block * fix: mm_ids_helper compact map corruption when ids contain -1 Negative ids inflated nex_prev for every expert, opening a phantom region at the head of the compact arrays that nothing wrote; the mmq quantize preprocessing iterates the flat range and read garbage row indices (illegal memory access). Skipped ids now occupy no compact position, ids_src1 is sentinel-filled, and the quantize gather/scatter kernels skip negative entries. * fix: mmvf id=-1 skip guard (float mul_mat_id vec path) Same treatment as mmvq: kernels early-return on a negative expert id instead of using it as a channel index, and the host launcher pre-zeroes the dst rows of skipped slots. Fixes wrong results (and a latent out-of-bounds channel read) for f16/f32 MUL_MAT_ID with -1 ids at small batch sizes. * fix: zero-init ids_dst before mm_ids_helper — garbage tail entries in the compact map caused pool-layout-dependent illegal memory accesses when ids contain -1 (tile-padded reads consumed uninitialized row indices) * feat: route MUL_MAT_ID with skip-capable ids away from mmq/mmf Ops whose ids may contain -1 announce it via op_params[0]; the CUDA dispatch then uses only the paths with verified skip support (mmvq for small batches, the general gather/scatter path otherwise). Large-batch mmq with sparse ids shows a pool-layout-dependent fault that resists isolation (all indexing audits pass; sanitizer and cuda-gdb unusable on the target driver); since the expert-pack design only decodes through packs, the safe route costs nothing in practice. mmq/mmf guards from earlier commits remain as hardening. * fix: op_params write in test — ggml_set_op_params_i32 is internal-only * feat: MoE expert cache loader — hot-pack build from routing profile init_moe_expert_cache() runs post-load: reads a moe-trace profile (GGML_MOE_CACHE_PROFILE) and slot count (GGML_MOE_CACHE_SLOTS), and for each CPU-resident MoE layer builds GPU-resident hot packs of the top-S experts plus i32 remap tables (hot: global->slot or -1; cold: global id or -1). Graph wiring (build_moe_ffn dual path) is the next commit; these tensors are unused until then, so this is behavior-neutral. * feat: MoE expert cache graph wiring — dual hot/cold mul_mat_id in build_moe_ffn When a layer carries hot packs, routed ids are remapped through moe_map_hot/cold (get_rows) and each expert projection runs as two skip-flagged mul_mat_id ops — hot pack on GPU, original tensors for cold — whose zero-padded outputs sum to the exact single-tensor result. Engaged for qwen35moe (separate gate/up/down, no expert scales); layers with merged gate_up, scales, or no packs use the unchanged path. * fix: flatten routed ids before map get_rows (batched-gather shape rule) * fix: topk ids are a strided view - cont before flatten for map lookup * fix: sched expert-copy path skips id=-1 (hot/cold expert split) * fix: restructure MoE pack dual path into full FFN chains per residency side Per-projection hot/cold interleave let the scheduler's split minimizer assign hot up/down matmuls to CPU, copying the GPU pack weights over PCIe every layer. Building one unbroken gate-up-swiglu-down chain per side keeps the hot pack on GPU; skipped (-1) rows stay zero through swiglu so a single add of the two chain outputs is exact. * feat: wire MoE expert cache into deepseek2 graph (GLM-4.7-Flash) * feat: add --moe-cache-profile and --moe-cache-slots flags Plumbed through llama_model_params; env vars remain as fallback so existing deployments keep working. Flags allow per-model settings in multi-model serving. * docs: fix moe-trace usage comment (output path is MOE_TRACE_OUT env) commit cd72c28 Author: Anirban Kar <thecodacus@gmail.com> Date: Fri Jul 24 17:53:41 2026 +0530 fix: restore closing brace lost in dequantize.cuh merge resolution commit 34f893e Merge: 40f2036 8f5ab83 Author: Anirban Kar <thecodacus@gmail.com> Date: Fri Jul 24 17:52:26 2026 +0530 merge: sync upstream master into perf commit 8f5ab83 Author: Matt Thompson <111157855+boondocklabs@users.noreply.github.com> Date: Fri Jul 24 03:54:47 2026 -0700 cohere2 moe template parser: enforce JSON schema for text responses if a response schema is provided (ggml-org#26018) commit 0cea362 Author: Xuan-Son Nguyen <son@huggingface.co> Date: Fri Jul 24 08:02:23 2026 +0200 vendor: update subprocess.h (ggml-org#26061) commit 40f2036 Author: Anirban Kar <thecodacus@gmail.com> Date: Sun Jul 19 04:34:24 2026 +0530 feat: fused turbo MMA flash-attention decode (turbo4/3/2, GQA tensor-core path) (TheTom#4) * feat(fattn): fused turbo4 MMA decode path (opt-in) — based on the spiritbuun fork MMA approach, adapted to graph-rotated Q Port buun's fused turbo4 tensor-core (MMA) flash-attention DECODE kernel into the turbo fork as an OPT-IN path (GGML_TURBO_MMA_FUSED=1, default OFF), routing turbo4-K==turbo4-V, D in {128,256}, decode (Q->ne[1]<=4) onto the GQA-packed MMA path so KV is read once per head-group instead of once per query head. How it works: - fattn-mma-f16.cuh: add flash_attn_ext_turbo4_load_tile (dequant turbo4 blocks -> SRAM in the f16 row-major half2 layout) seeded with OUR Lloyd-Max centroids (turbo-quant.cuh:297, NOT buun's -0.2415xx table). Thread type_K/type_V through iter/process_tile/global flash_attn_ext_f16 (defaults F16 -> existing f16/q8 MMA byte-identical). For turbo: force nstages=0 (the cp.async multi-stage path would read raw turbo bytes as half2), pass RAW byte pitch nb11/nb21 (not /sizeof(half2)), and dequantize K/V in the load tiles. sizeof(block_turbo4_0)-driven pointer math (66B block, no rnorm) — never assumes a fixed offset. - fattn-mma-turbo.cuh + 14 instance files (7 reachable ncols pairs x D in {128,256}): host launcher reusing the f16 device kernel with type_K/type_V=TURBO4_0 and need_f16_K/V=false so launch_fattn does NOT pre-convert KV to f16. - fattn.cu: switch_ncols2 -> per-ncols2 dispatchers (only the 7 compiled instances), env latch, and the gate (Q untouched — our fork already rotates Q at the graph level, so NO inline FWHT / src swap, which would double-rotate Q). Results (Qwen3.6-35B-A3B turbo4, GQA8, D=256, 5090): depth MMA VEC+PDL buun 8192 213.8 188.5 200 (+6.9% vs buun) 16384 207.6 171.7 193 (+7.5% vs buun) 32768 191.3 141.1 187 (+2.3% vs buun) MMA holds flat at depth where VEC collapses; the gap vs VEC grows with KV length. Quality: Mean KLD vs f16 base = 0.008396 (== VEC baseline). The MMA path is NOT strictly token-identical to VEC — MMA and VEC accumulate the P*V (VKQ) reduction in f16 with different reduction trees, so a near-tie greedy token can flip (~1/25 on a hard tie; same irreducible f16-order effect that exists between base f16-MMA and f16-VEC, and within the VEC-vs-base "same-top-p" 96.3% noise floor). Because strict token-identity does not hold, the gate DEFAULTS OFF and ships as opt-in; VEC remains the untouched default/kill-switch. (cherry picked from commit 5450669510a5cd0a356a863923b057f2552c6fbb) (cherry picked from commit b3e51cf) * turbo4 MMA decode: default ON (GQA-packed tensor-core path, quality-neutral; GGML_TURBO_MMA_FUSED=0 kill-switch) (cherry picked from commit 726234b538194c644be4066e1413b91f9edf12e7) (cherry picked from commit 38dc665) * fattn: extend fused MMA decode to turbo3 + turbo2 (parity on turbo4; extends the MMA path to turbo3/2) The GQA-packed tensor-core decode path was turbo4-only; turbo3/turbo2 fell back to VEC and collapsed at depth (turbo3 tg32 224->124 from d2048->d32768), losing to spiritbuun there. Added flash_attn_ext_turbo3_load_tile (3-bit split index: 2 low bits from qs + 1 high bit from signs) and flash_attn_ext_turbo2_load_tile (plain 2-bit), type-dispatched both the K and V MMA load sites, and routed the gate by K->type. is_turbo_kv already matches any non-f16 KV, so the kernel needed only the per-type tile loaders + instances (14 instance files now emit all three types). Correctness: MMA decode is bit-exact to VEC (Mean KLD @2048 chunks=8: turbo3 0.020278==0.020278, turbo2 0.041490==0.041490, turbo4 0.009296 unchanged). Decode tg32 (qwen3.6-35B-A3B, RTX 5090), MMA now flat at depth: turbo3 2k/8k/16k/32k = 224/225/217/203 (buun 193/193/183/178; was VEC 224/187/160/124) turbo2 2k/8k/16k/32k = 228/224/217/203 (buun 213/210/204/200; was VEC 229/218/203/183) turbo3/turbo2 now hold flat at depth (neither fork had them on the MMA path before). GGML_TURBO_MMA_FUSED=0 restores VEC. (cherry picked from commit 57ad5edd53a7fced0b99fecb99326132cbdb65d0) (cherry picked from commit 4e223ee) * fix(fattn): gate turbo2 fused-MMA decode to head_dim 128 Tester @everson isolated a turbo2-only decode regression on RTX 5060 Ti (Gemma-12B, head_dim 256): turbo2/turbo2 -2.46% at 8K to -1.17% at 200K, fully restored by GGML_TURBO_MMA_FUSED=0 at every depth, while turbo3 was neutral and turbo4 exact parity. At 2-bit KV the fused path's GQA-pack saving is tiny while its dequant/no-pipeline overhead is unchanged, so on head_dim 256 it is neutral on high-bandwidth GPUs (Ornith-35B and a dense Gemma-4-12B on RTX 5090, both within noise) and only costs on bandwidth-limited cards. turbo2 fused remains a large depth win on dense head_dim 128 (Llama-3.1-8B on RTX 5090: +7% at 8K, +24% at 32K, +69% at 131K), so this gates turbo2 fused to head_dim 128 only; head_dim 256 turbo2 falls through to VEC, which is the exact baseline path GGML_TURBO_MMA_FUSED=0 restored. turbo3 and turbo4 fused are unchanged at both head dims. Also corrects the stale "default OFF" gate comments to match the default-ON code. (cherry picked from commit 539ce5d) --------- Co-authored-by: TheTom <tturney1@gmail.com> commit 4cd54c6 Author: Anirban Kar <thecodacus@gmail.com> Date: Sun Jul 19 04:14:41 2026 +0530 fix: allow mixed turbo/q8_0 KV cache types in CUDA flash-attention selection (TheTom#3) * fix: allow mixed turbo/q8_0 KV cache types in CUDA flash-attention selection Without GGML_CUDA_FA_ALL_QUANTS (the default), kernel selection rejected any K/V type mismatch. A turbo KV setup is almost always mixed (ctk q8_0 + ctv turbo2/3), so the flash-attention probe failed, FA was auto-disabled, and context creation either refused the quantized V cache outright or fell back to the non-FA path whose KQ buffer (~8.7 GiB at 64k ctx, ubatch 1024) OOMs at load. Allow exactly the mixed pairs that have FA template instances compiled in: turbo2/3/4 against each other, q8_0, and f16. Also gate turbo head-dim geometry (multiples of 64; 128 for turbo4) so shapes without a compiled instance report unsupported instead of asserting at dispatch. Forcing -fa on masked the bug (the probe is skipped); server deployments use fa auto and hit it on every load. * feat: accept turbo2/3/4 cache type names in llama-bench The common arg parser already accepts them for llama-completion and llama-server; llama-bench has its own name table and rejected -ctk/-ctv turbo values. commit 9338207 Author: Anirban Kar <thecodacus@gmail.com> Date: Sun Jul 19 01:59:58 2026 +0530 CUDA: add Q2_0 support (TheTom#2) Co-authored-by: Pasha Khosravi <khosravipasha@gmail.com> commit 170322f Author: Anirban Kar <thecodacus@gmail.com> Date: Sat Jul 18 15:27:13 2026 +0530 feat: TurboQuant KV-cache quantization — all backends (TheTom#1) * feat: TurboQuant KV-cache quantization — core types + CPU backend Port of the TurboQuant codec (arXiv 2504.19874, reference: TheTom/llama-cpp-turboquant) onto this fork's base: WHT-rotated polar quantization for the KV cache. - new types TURBO2_0/TURBO3_0/TURBO4_0 (KV, 2/3/4-bit blocks of 128) and TQ3_1S/TQ4_1S (WHT-rotated Lloyd-Max weight formats), renumbered for this tree's type table - GGML_OP_TURBO_WHT + ggml_turbo_wht() builder (forward/inverse Walsh-Hadamard rotation with sign arrays, InnerQ scale hook) - ggml-turbo-quant.c codec core, rotation data tables, CPU op + vec_dot paths, type traits, quantize-chunk wiring - CUDA/Metal/Vulkan kernel sources imported (wired in follow-ups) - verified: test-turbo-quant round-trip on CPU (cosine 0.986-0.996) * feat: turbo KV cache — graph wiring and CLI types - pre-rotate Q (forward WHT) at every cache-backed attention site when K is a turbo type; no-op otherwise - inverse WHT on attention output in both FA and non-FA paths of build_attn_mha (group size derived from K, contiguity enforced) - register turbo2/3/4 as --cache-type-k/--cache-type-v values Core scope: no InnerQ scaling, no head-dim padding — turbo KV requires 128-aligned (or 64-group) head dims * feat: turbo KV cache — Metal backend (WHT kernel + FA dispatch) Wire the TurboQuant Metal backend host side and kernels on top of the CPU core already on this branch. Metal kernels (ggml-metal.metal): - Inlined turbo-wht.h (WHT sign arrays incl. half4 variants + FWHT butterfly + rotate fwd/inverse) and the turbo2/3/4 centroids. - quantize/dequantize_turbo2_0/3_0/4_0 (+_t4) and full-block cache. - kernel_turbo_wht (O(d log d) Walsh-Hadamard rotation op) and kernel_turbo4_dequant_f16 bulk prefill dequant. - kernel_set_rows_turbo/2/4 cache-quantize kernels + instantiations. - flash_attn_ext / flash_attn_ext_vec instantiations for symmetric and asymmetric turbo2/3/4 and q8_0<->turbo pairs. - TurboFlash two-pass kernels (p1/p2), default-disabled via env. Host side: - ggml-metal-device.cpp: get_pipeline_turbo_wht; flash-attn pipeline base names now encode both K and V types (k%s_v%s) so asymmetric K/V pairs resolve unambiguously. All existing FA kernel host_names in the shader renamed to match. - ggml-metal-device.m: TurboQuant compile-time env prep (4-mag LUT auto select, sparse-V, profile mode); supports_op accepts turbo K types and asymmetric turbo/q8_0 FA pairs; GGML_OP_TURBO_WHT. - ggml-metal-ops.cpp: GGML_OP_TURBO_WHT encode case + impl; asymmetric K/V assert relaxation; TurboFlash gating (use_turbo_flash, extra_tmp sizing, force-nonvec) + two-pass dispatch; +128 float FA smem for the turbo block cache. - ggml-metal-impl.h / -device.h / -ops.h: kargs (turbo_wht, turbo_flash_p1/p2), FC offsets, pipeline/op declarations. Adaptations vs the reference tree (older upstream vintage): - Reference intertwines separate TQ3_1S/TQ4_1S weight-quant work; only the turbo KV-cache lines are ported here. - set_rows_turbo dispatch selection is not wired: our upstream refactored get_pipeline_set_rows/kernel_set_rows_q32 to a different signature, and it is out of scope for this KV port. supports_op therefore does not claim CPY/SET_ROWS turbo support, so cache writes fall back to CPU; the set_rows_turbo shader kernels are present for when dispatch lands. - No InnerQ scale tensors touched (graph passes nullptr scale). * feat: turbo KV cache — CUDA backend (set-rows quantize, fattn vec dispatch, dequant) Port the TurboQuant CUDA backend from the reference implementation, adapted to this newer upstream vintage: - set-rows.cu: k_set_rows_turbo2/3/4 kernels (WHT rotation + norm-corrected PolarQuant KV write, InnerQ calibration hooks) + dispatch. turbo4 rnorm write dropped (field removed from block_turbo4_0 on this branch). - dequantize.cuh: dequantize_turbo2/3/4_0 and dequantize_tq3_1s/tq4_1s device functions. - convert.cu: turbo + TQ dequant cases in the fp16/fp32 contiguous and non-contiguous dispatch tables; warp-cooperative TQ4_1S dequant kernel. - fattn-common.cuh: vec_dot_fattn_vec_KQ_turbo2/3/4_0, dequantize_V_turbo* and get_vec_dot_KQ / get_dequantize_V dispatch entries. - fattn-vec.cuh: turbo K types take the float-Q path (Q_q8_1 off), turbo V uses the ne=4 quantized dequant path; extern decls for the 21 compiled K/V type combinations. Reference-side perf extras (shared-mem centroid LUT, sparse-V skip, nthreads tuning) intentionally not ported. - fattn.cu: turbo types accepted in kv_type_supported, vec kernel cases registered. Decode goes through the VEC kernel; prefill converts KV to f16 via the ported nc dequant path. - fattn-mma-turbo.cuh: reduced to a stub launcher. No typed MMA kernel exists in the reference either; the pre-existing template instances now compile against the stub and the dispatcher never selects this path. - ggml-cuda.cu: GGML_OP_TURBO_WHT dispatch + supports_op, SET_ROWS turbo gates (head_dim % 64, turbo4 % 128). - CMakeLists.txt: turbo fattn-vec instances added to the non-FA_ALL_QUANTS build. Verified on RTX 3060 (sm_86) with Bonsai-27B-Q1_0, partial offload, fa on: coherent generation for turbo2/3/4; perplexity over 4 chunks: f16 KV 7.9486 +/- 0.73, turbo3 KV 8.0261 +/- 0.74 (CUDA) vs 8.0279 (CPU-only). * feat: turbo KV cache — Vulkan backend (set-rows quantize, turbo3 FA dispatch, WHT op) - register turbo_wht.comp and set_rows turbo2/3/4 quantize variants in vulkan-shaders-gen; add dedicated DATA_A_TURBO3_0 flash-attn SPIR-V (scalar + coopmat1) with turbo3-only K/V bindings - copy_to_quant: 128-thread per-block set_rows path (normalize, randomized WHT, Lloyd-Max centroid quantize, subgroup pack), shared WHT tables across the three turbo widths - flash_attn_dequant: turbo2/3/4 centroid dequant macros wired into the FaTypeK/FaTypeV runtime switch (ids 43/44/45); f16 K/V aliases guarded out under DATA_A_TURBO3_0 - ggml-vulkan: GGML_OP_TURBO_WHT pipeline + dispatch, turbo set_rows pipelines (required subgroup size 32), FA pipeline selection for turbo3, supports_op gates (128-elem blocks, turbo3 K/V symmetry, scalar/coopmat1 only — cm2 falls back to scalar) * fix: extern initialization warning and C linkage for wht group size CI builds with -Werror: defining the GGML_API-decorated global with an initializer warns on every platform. Plain definition in the C TU; the C++ consumer now declares it extern "C" (it previously referenced a mangled symbol). * feat: turbo2/turbo4 FA dequant for Vulkan - generalize the dedicated turbo flash-attn SPIR-V path from turbo3-only to all three tiers: DATA_A_TURBO_KV umbrella define in types.glsl gates out the f16/q4/q5/q8 K/V binding aliases in flash_attn.comp, flash_attn_cm1.comp and flash_attn_dequant.glsl for any turbo tier - flash_attn_dequant: per-compilation dequantize4 for DATA_A_TURBO2_0 (2-bit centroids, fp16 norm + 32B idx) and DATA_A_TURBO4_0 (4-bit nibble idx, fp16 norm + fp16 aux + 64B) alongside the turbo3 one - vulkan-shaders-gen: register flash_attn_f32_f16_turbo{2,4}_0 scalar + coopmat1 variants via a per-tier loop - ggml-vulkan: FA pipeline selection picks the matching turbo SPIR-V for turbo2/turbo4 k_type (scalar + cm1), supports_op requires symmetric K/V per tier (turbo2/turbo2, turbo3/turbo3, turbo4/turbo4) * fix: bump RPC proto patch version for GGML_OP_TURBO_WHT * fix: export WHT group size via GGML_API for Windows shared builds * fix: HIP/MUSA portability for turbo kernels - map cudaMemcpyToSymbol/FromSymbol and __ballot_sync in vendor headers - pass explicit width to __shfl*_sync (HIP macros take 4 args) - drop unused warp_id locals flagged by nvcc -Werror * fix: drop SYCL turbo fattn-vec instances (no turbo dequant support in SYCL fattn) * fix: size vec_dot buffer by vec_dot_type and handle turbo/tq types in quantize test - tmp_q2 overflowed for types with vec_dot_type F32 (4 bytes per element) - skip turbo KV types: they quantize in the rotated domain, round trip does not reproduce the input by design - tq3_1s uses the 3-bit and lowbit thresholds * fix: add turbo fattn-vec instances to HIP and MUSA source lists * fix: add turbo fattn vec kernels to HIP VGPR ignore list The gfx908 build of the ncols=2 turbo3 vec kernels lands at 259 total VGPRs, slightly over the 256 gate. Register them like the other known high-pressure fattn kernels. commit 0888217 Merge: 5e7f627 7d56da7 Author: Anirban Kar <thecodacus@gmail.com> Date: Fri Jul 17 20:02:48 2026 +0530 chore: sync perf with upstream master commit 5e7f627 Author: Anirban Kar <thecodacus@gmail.com> Date: Wed Jul 8 23:56:46 2026 +0530 docs(readme): add usage + benchmark instructions for the MoE-offload optimizations commit 5f83fbb Author: thecodacus <thecodacus@gmail.com> Date: Thu Jul 2 12:42:55 2026 +0000 ggml : size prefetch slots per layer and fix fallback use-after-free Profiling showed the 2-slot rotation stalls ~5.5ms per layer waiting for the down_exps upload: 3 tensors per MoE layer need 3 slots for a full layer of lookahead. Default is now 3 (one layer), configurable via the env var value, degrading to however many slots fit on OOM. Also fixes a use-after-free: on a failed slot regrow the repointed staging tensors kept dangling pointers into freed device memory, which graph reuse carried into later evals. Staging repoints are now restored right after kernel launch, slots are sized once from the graph max, and allocation happens before freeing. Qwen3.6-35B-A3B pp2048 on RTX 3060: 1643 -> 1880 t/s (mainline: 1143). commit 1163cb3 Author: thecodacus <thecodacus@gmail.com> Date: Thu Jul 2 12:16:39 2026 +0000 ggml : overlap offloaded expert weight uploads with compute At large batch sizes virtually every expert is used, so the per-layer routing-ids readback that mainline waits on buys nothing while forcing a full device sync per expert tensor (3x per MoE layer). Above a batch threshold, upload the full expert tensors through a second backend instance on the same device (own stream) with two event-ordered staging slots, so uploads for tensor N+1 overlap compute of tensor N. Qwen3.6-35B-A3B pp2048 on RTX 3060 (-ncmoe 26, ub 2048): 1383 -> 1663 t/s. Generation output verified token-identical; decode path unaffected. Opt-in via GGML_SCHED_PREFETCH_EXPERTS=1. commit 20f5994 Author: thecodacus <thecodacus@gmail.com> Date: Thu Jul 2 09:44:53 2026 +0000 llama : pin mmap-backed CPU weights for faster H2D uploads Wire the existing GGML_CUDA_REGISTER_HOST path back up: after model load, cudaHostRegister the mmap pages backing weights kept in system memory. Recovers pageable-copy losses when MoE experts are streamed to the GPU during prefill (n-cpu-moe): Qwen3.6-35B-A3B pp2048 1144 -> 1385 t/s on RTX 3060. Opt-in via GGML_CUDA_REGISTER_HOST=1, unchanged otherwise.

…eTom#5) * feat: llama-moe-trace tool + expert-cache policy simulator Phase 0 of the MoE expert-cache work: capture routed expert ids (ffn_moe_topk) per layer per decode step via the scheduler eval callback — backend-agnostic, so CPU-offloaded expert layers are traced too. Companion simulate.py replays traces against cache policies (static top-S, LRU, LFU-decay, periodic re-election) and prints hit-rate vs VRAM-budget curves plus routing-skew stats to decide whether expert caching is worth building at all. * fix: moe-trace — common_init_from_params returns ptr with accessors * fix: moe-trace — include arg.h for common_params_parse * fix: moe-trace — take output path from MOE_TRACE_OUT env, -o not in this arg set * fix: moe-trace — topk tensor is a strided view; copy full byte range and index by nb Sizing the host buffer as n_used*n_tokens under-allocated: ffn_moe_topk is a view over the argsort output, so nb[1] spans n_expert ids per row and ggml_backend_tensor_get writes the whole strided range (heap corruption, malloc corrupted-top-size abort mid-prefill). * feat: moe-trace simulator v2 — model upload bytes/token and miss/upload cost per policy * feat: mul_mat_id tolerates id=-1 (skip) on CPU and CUDA fallback path Groundwork for the hot/cold expert-pack split: an id of -1 means the routed expert is not owned by this pack. The op contributes a zeroed output row for that slot so hot-pack and cold-pack results merge additively. CUDA fast paths (mmvq_id/mmq_id/mmf_id) still assert non-negative ids and are handled next. * feat: id=-1 skip support in CUDA mul_mat_id fast paths + tests mmid.cu gains a generic kernel that zeroes the dst rows of (token, slot) pairs whose expert id is -1; mmq and mmf launch it after the shared mm_ids_helper (which already drops unmatched ids from its compact mapping), and mmvq launches it before its kernels, which now early-return instead of using a negative id as a channel index (out of bounds read otherwise). test-backend-ops gains skip_ids variants of MUL_MAT_ID covering the mmvq/mmq/mmf/fallback dispatch corners. * fix: skip_ids MUL_MAT_ID test cases were inside an #if 0 block * fix: mm_ids_helper compact map corruption when ids contain -1 Negative ids inflated nex_prev for every expert, opening a phantom region at the head of the compact arrays that nothing wrote; the mmq quantize preprocessing iterates the flat range and read garbage row indices (illegal memory access). Skipped ids now occupy no compact position, ids_src1 is sentinel-filled, and the quantize gather/scatter kernels skip negative entries. * fix: mmvf id=-1 skip guard (float mul_mat_id vec path) Same treatment as mmvq: kernels early-return on a negative expert id instead of using it as a channel index, and the host launcher pre-zeroes the dst rows of skipped slots. Fixes wrong results (and a latent out-of-bounds channel read) for f16/f32 MUL_MAT_ID with -1 ids at small batch sizes. * fix: zero-init ids_dst before mm_ids_helper — garbage tail entries in the compact map caused pool-layout-dependent illegal memory accesses when ids contain -1 (tile-padded reads consumed uninitialized row indices) * feat: route MUL_MAT_ID with skip-capable ids away from mmq/mmf Ops whose ids may contain -1 announce it via op_params[0]; the CUDA dispatch then uses only the paths with verified skip support (mmvq for small batches, the general gather/scatter path otherwise). Large-batch mmq with sparse ids shows a pool-layout-dependent fault that resists isolation (all indexing audits pass; sanitizer and cuda-gdb unusable on the target driver); since the expert-pack design only decodes through packs, the safe route costs nothing in practice. mmq/mmf guards from earlier commits remain as hardening. * fix: op_params write in test — ggml_set_op_params_i32 is internal-only * feat: MoE expert cache loader — hot-pack build from routing profile init_moe_expert_cache() runs post-load: reads a moe-trace profile (GGML_MOE_CACHE_PROFILE) and slot count (GGML_MOE_CACHE_SLOTS), and for each CPU-resident MoE layer builds GPU-resident hot packs of the top-S experts plus i32 remap tables (hot: global->slot or -1; cold: global id or -1). Graph wiring (build_moe_ffn dual path) is the next commit; these tensors are unused until then, so this is behavior-neutral. * feat: MoE expert cache graph wiring — dual hot/cold mul_mat_id in build_moe_ffn When a layer carries hot packs, routed ids are remapped through moe_map_hot/cold (get_rows) and each expert projection runs as two skip-flagged mul_mat_id ops — hot pack on GPU, original tensors for cold — whose zero-padded outputs sum to the exact single-tensor result. Engaged for qwen35moe (separate gate/up/down, no expert scales); layers with merged gate_up, scales, or no packs use the unchanged path. * fix: flatten routed ids before map get_rows (batched-gather shape rule) * fix: topk ids are a strided view - cont before flatten for map lookup * fix: sched expert-copy path skips id=-1 (hot/cold expert split) * fix: restructure MoE pack dual path into full FFN chains per residency side Per-projection hot/cold interleave let the scheduler's split minimizer assign hot up/down matmuls to CPU, copying the GPU pack weights over PCIe every layer. Building one unbroken gate-up-swiglu-down chain per side keeps the hot pack on GPU; skipped (-1) rows stay zero through swiglu so a single add of the two chain outputs is exact. * feat: wire MoE expert cache into deepseek2 graph (GLM-4.7-Flash) * feat: add --moe-cache-profile and --moe-cache-slots flags Plumbed through llama_model_params; env vars remain as fallback so existing deployments keep working. Flags allow per-model settings in multi-model serving. * docs: fix moe-trace usage comment (output path is MOE_TRACE_OUT env)

TheTom and others added 30 commits March 26, 2026 12:15

docs: log simd_broadcast attempt — no speed improvement TheTom#23

283441c

Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

docs: log threadgroup attempt — no speed improvement, rethinking TheT…

aede1bb

…om#23 Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

docs: final investigation log — 77.7 tok/s, 91% of q8_0

c3a7afd

Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

docs: perplexity 6.194 confirmed — 1.4% of q8_0 TheTom#30

d9b9725

Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

apollosenvy mentioned this pull request Mar 30, 2026

feat: HIP/ROCm support for turbo3/turbo2 (7900 XTX) #31

Merged

TheTom force-pushed the feature/turboquant-kv-cache branch from 63b832b to e9c54d5 Compare April 3, 2026 16:14

TheTom force-pushed the feature/turboquant-kv-cache branch from 10cb187 to 0d6b38a Compare April 8, 2026 23:49

twobombs mentioned this pull request Apr 11, 2026

vulkan: TQ4_1s support for model weights #69

Merged

TheTom force-pushed the feature/turboquant-kv-cache branch from 45f8a06 to 1073622 Compare April 16, 2026 01:14

apollosenvy closed this Apr 17, 2026

aureliahazel93-commits added a commit to aureliahazel93-commits/llama-cpp-turboquant that referenced this pull request Jul 8, 2026

ledger:TheTom#5 pipeline:phase20-impl-001-pre-snapshot confidence:0.0…

8497a3c

…0 findings:0

brosequist mentioned this pull request Oct 28, 2026

Pr/rdna4 decode fusions v2 #410

Open

Navigation Menu

Repository navigation

feat: ROCm/HIP support for turbo3 KV cache (gfx1100/RDNA3) - #5

feat: ROCm/HIP support for turbo3 KV cache (gfx1100/RDNA3)#5
apollosenvy wants to merge 115 commits into
TheTom:feature/turboquant-kv-cachefrom
apollosenvy:rocm-turbo3-pr

apollosenvy commented Mar 27, 2026

Uh oh!

apollosenvy commented Mar 30, 2026

Uh oh!

apollosenvy commented Apr 17, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

Conversation

apollosenvy commented Mar 27, 2026

Summary

What's included

Critical bugfix: K/V rotation before cache write

Performance

Known limitation

Files changed (24 files, +595 lines)

Test plan

Uh oh!

apollosenvy commented Mar 30, 2026

Uh oh!

apollosenvy commented Apr 17, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants