Skip to content

fix: HIP/ROCm compatibility — check cudaMemcpyToSymbol errors, guard …#41

Draft
terrysimons wants to merge 157 commits intoTheTom:feature/turboquant-kv-cachefrom
terrysimons:feature/turboquant-kv-cache
Draft

fix: HIP/ROCm compatibility — check cudaMemcpyToSymbol errors, guard …#41
terrysimons wants to merge 157 commits intoTheTom:feature/turboquant-kv-cachefrom
terrysimons:feature/turboquant-kv-cache

Conversation

@terrysimons
Copy link
Copy Markdown

@terrysimons terrysimons commented Apr 1, 2026

…D>=576 MMA

Add CUDA_CHECK() to all cudaMemcpyToSymbol/cudaMemcpyFromSymbol calls in the InnerQ calibration path. On HIP, unchecked errors from these calls are sticky and poison the runtime, causing subsequent API calls to fail with 'no ROCm-capable device is detected'.

Also guard the D>=576 MMA flash attention dispatch and kernel selection with #ifndef GGML_USE_HIP, matching the existing D>=576 tile exclusion (these kernels exceed HIP's shared/local memory limits).

Tested on: ROCm 6.4.4, gfx1151 (AMD Ryzen AI Max+ 395 / Strix Halo)

Overview

Additional information

Requirements

  • I have read and agree with the contributing guidelines
  • AI usage disclosure:
    YES - Fixed cuda guards that weren't in place that caused failures.

TheTom and others added 28 commits April 2, 2026 13:07
@TheTom @claude
WIP: add TurboQuant KV cache types (turbo3, turbo4)
da7c83f 
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>
@TheTom @claude
feat: Metal kernels for TurboQuant KV cache (turbo3, turbo4) TheTom#21
42ff28f 
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>
@TheTom @claude
feat: full TurboQuant with rotation matrices in Metal kernels TheTom#21
57976f0 
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>
@TheTom @claude
feat: inline rotation matrices in Metal shader + C round-trip test Th…
95943b3 
…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>
@TheTom @claude
fix: remove thread static from Metal dequantize, fix stale code TheTo…
39792d8 
…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>
@TheTom @claude
feat: replace dense 128x128 matvec with Fast Walsh-Hadamard rotation T…
e35d74a 
…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>
@TheTom @claude
docs: detailed speed investigation plan for TurboQuant Metal shader T…
ff70c67 
…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>
@TheTom @claude
docs: log simd_broadcast attempt — no speed improvement TheTom#23
d97ec7f 
Co-Authored-By: tturney@psyguard.ai
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
@TheTom @claude
docs: log threadgroup attempt — no speed improvement, rethinking TheT…
9bbdc3b 
…om#23

Co-Authored-By: tturney@psyguard.ai
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
@TheTom @claude
docs: CRITICAL — dequant is NOT the bottleneck, no-op still 2.4 tok/s T…
3894b00 
…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>
@TheTom @claude
fix: inline turbo-wht.h — was causing CPU fallback, not Metal! TheTom#23
4162eb9 


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>
@TheTom @claude
docs: real Metal benchmarks after #include fix — 8× gap not 35× TheTo…
d85a4a1 
…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>
@TheTom @claude
docs: final investigation summary + upstream tracking TheTom#23 TheTo…
6206ee1 
…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>
@TheTom @claude
docs: upstream competitive intel — pre-rotate-queries is the key TheT…
c18effa 
…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>
@TheTom @claude
docs: speed ceiling test — 49 tok/s without dequant rotation (4.6× ga…
d3eee2b 
…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>
@TheTom @claude
docs: pre-rotate-queries implementation plan + speed ceiling 49 tok/s
3fc66fe 
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>
@TheTom @claude
feat: pre-rotate-queries optimization — 51.4 tok/s (5× speedup) TheTo…
1932133 
…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>
@TheTom @claude
docs: final investigation summary — 2.4 → 51.4 tok/s journey complete T…
e51418b 
…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>
@TheTom @claude
feat: MSE-only mode — drop QJL, all 3 bits to PolarQuant TheTom#23
f989c5e 
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>
@TheTom @claude
docs: Change 2 not needed — Q rotation overhead is negligible
6f15755 
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>
@TheTom @claude
docs: block size is the bottleneck — q4_0 at block 32 = 100% of q8_0
3081fda 
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>
@TheTom @claude
feat: block size 32 — 77.7 tok/s MoE (91% of q8_0), 17.0 Qwopus (97%) 🎉
cde1733 
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>
@TheTom @claude
fix: TURBO_D=128 independent of QK_TURBO3, file turbo4 bugs TheTom#29
a91039b 
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>
@TheTom @claude
docs: final investigation log — 77.7 tok/s, 91% of q8_0
3918d41 
Co-Authored-By: tturney@psyguard.ai
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
@TheTom @claude
CRITICAL: turbo3 perplexity is 165.6 vs q8_0 6.1 — quality broken The…
846705d 
…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>
@TheTom @claude
CRITICAL: found TWO root causes for PPL=165 TheTom#30
8557825 
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>
@TheTom @claude
docs: bisect confirms block size innocent, rotation access is the bug T…
e20ce5d 
…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>
@TheTom @claude
fix: restore inverse rotation in dequant — PPL 6.19 (1.2% of q8_0) Th…
c52aef5 
…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>
@TheTom @claude
fix: remove redundant extern from GGML_API macro (GCC 13.3 hard error)
e9c54d5 
GCC 13.3 on Ubuntu 24.04 hard errors on the redundant `extern` in the
GGML_API macro when used inside `extern "C" {}` blocks. MSVC and Clang
accept it silently.

Reported independently by joemc1470 (RX 6600 HIP) and
christopheraleman1015 (WSL2 Ubuntu).

Co-Authored-By: tturney@psyguard.ai
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
@TheTom TheTom force-pushed the feature/turboquant-kv-cache branch from 63b832b to e9c54d5 Compare April 3, 2026 16:14
TheTom and others added 3 commits April 3, 2026 12:21
@TheTom
Merge remote-tracking branch 'origin/master' into feature/turboquant-…
bc05a68 
…kv-cache
@iamwavecut
Enhance Metal operations for TQ weights and concurrency handling for …
ab9990d 
…Gemma 4

- Introduced a new function to check if TQ source tensors mutate during operations.
- Updated matrix multiplication logic to handle TQ weights more effectively, ensuring correct concurrency behavior.
- Adjusted Metal kernel definitions to support TQ weights with improved dispatch parameters.
- Enhanced comments for clarity on concurrency issues related to TQ weights.
@iamwavecut
Update GGMLQuantizationType and LlamaFileType enums to include TQ3_1S…
883bf4a 
… and TQ4_1S quantization types with corresponding sizes in GGML_QUANT_SIZES.
@stragulus
Copy link
Copy Markdown

stragulus commented Apr 5, 2026

For those using hip/rocm, be warned that you will likely run into issues with vram memory leaking. After a lot of struggling I discovered that this also happens on the mater llama.cpp branch if you pick a cache type key/value other than f16. The exact same issue plagues the turbo3/4 quants as well sadly. For more information, check out the bug report which has reproduction steps: ggml-org#19979

All branches I've tested (including the one in this PR) suffer from this issue on my system.

signalnine and others added 15 commits April 6, 2026 08:10
@signalnine
perf: TQ4_1S native kernel 3.5× faster — 240 t/s (was 68), smaller VR…
cd8f9fc 
…AM than q8_0

Autoresearch-discovered optimizations for TQ4_1S weight mul_mat_vec kernel.
Native TQ4_1S at 5.0 bpv now runs 36% FASTER than the q8_0 load-time
conversion (240 vs 176 t/s) while using 1.7× LESS VRAM (4.5 vs 7.5 GiB).

Key optimizations (found via 86 automated experiments):
1. fp16 activation buffer — halves activation bandwidth (the bottleneck)
2. Shared-memory centroid LUT — eliminates constant memory serialization
   on divergent lane access (+89% single change)
3. Half2 arithmetic + strided block processing — 2× arithmetic density
4. Vectorized 128-bit loads — uint32×4 weights, int4 activations (+45%)
5. Register __byte_perm centroid decode — zero-memory centroid lookup
6. NWARPS 8→4

Also:
- Load-time q8_0 conversion now opt-in (GGML_TQ_CONVERT_Q8=1) instead
  of default. Native kernel is strictly better on both speed and VRAM.
- Autoresearch harness gains coherence testing (server API + factual Q&A)
  to catch silent corruption that PPL alone misses.

Benchmarks (RTX 5090, Qwen2.5-7B-Instruct TQ4_1S):
  Upstream V12 runtime:  67 t/s  (4.5 GiB VRAM)
  q8_0 conversion:     176 t/s  (7.5 GiB VRAM)
  Native optimized:    240 t/s  (4.5 GiB VRAM)  ← this PR

Quality (vs f16 baseline):
  PPL:      7.54 (f16: 7.18, q8_0 conv: 7.55)
  Mean KLD: 0.056 (q8_0 conv: 0.057, q4_0: 0.078)
  NIAH:     5/5
  Coherence: 4/4 (Paris, 4, print, Shakespeare)
@signalnine
perf: warp-cooperative TQ4_1S dequant (16× less compute per block)
67cb511 
Replace per-element generic dequant template (which repeats the full
32-element WHT butterfly 16 times per block) with a warp-cooperative
version using __shfl_xor_sync. One WHT per block instead of 16.

Note: this improves the dequant kernel itself but doesn't fix the
prefill gap (5.9K vs 13.3K). The bottleneck is cuBLAS fp32 GEMM vs
the q8_0 conversion path's native int8 tensor core GEMM. The dequant
was never the slow part — the GEMM dispatch is fundamentally different.

For prefill-heavy workloads, load-time q8_0 conversion remains the
recommended path (default ON). GGML_TQ_NATIVE=1 for decode-heavy
interactive chat where the +29% decode speed matters more.
@signalnine
feat: multi-token TQ4_1S dp4a kernel + multi-GPU fix + static build fix
94763cf 
- Multi-token dp4a kernel for ne[1]≤8 (speculative decoding, small batches)
  loads weight data once per block, reuses across all ncols_dst tokens
- Runtime TQ4_1S→fp16 dequant + cuBLAS for ne[1]>8 prefill
- Fix multi-GPU crash: replace static global CUDA buffers with per-device
  pool allocations from ctx.pool(id), matching mmvq.cu pattern
- Fix static build: TURBO_IQ_API wrapped in #ifdef GGML_BACKEND_SHARED
@signalnine
fix: replace __dp4a with ggml_cuda_dp4a for HIP/ROCm compatibility
097d765
@TheTom
Merge pull request TheTom#52 from iamwavecut/feature/turboquant-kv-cache
7433ad9 
Enhance Metal operations for TQ weights and concurrency handling
@TheTom @claude
fix: GCC double extern in ops.cpp turbo3_cpu_wht_group_size
dcf8430 
extern "C" GGML_API creates double extern on paths where GGML_API
expands to 'extern'. Wrap in extern "C" {} block instead.

Reported by Madreag on RTX 5090 WSL2.

Co-Authored-By: Tom Turney <tturney1@gmail.com>
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
@TheTom @claude
feat: add MoE expert count kernel instantiations + TQ4_1S backend tests
acad28d 
Metal MoE support:
- Add kernel_mul_mm_id_map0 instantiations for ne20 = 32, 60, 64, 128, 160, 256
- Covers Yuan, Qwen1.5-MoE, OLMoE, Qwen2/3-MoE, Mistral Small 4,
  Llama 4 Maverick, DeepSeek-V2/V3, Qwen3.5-35B/122B
- Note: ne02=256 (Qwen3.5-35B-A3B) hits shmem assert in llama-server
  with flash attention — needs chunked map0 dispatch (follow-up)

Backend tests:
- Add TQ3_1S and TQ4_1S to all_types array in test-backend-ops
- Enables GET_ROWS and MUL_MAT coverage for WHT-rotated weight types

Co-Authored-By: tturney@psyguard.ai
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
@TheTom @claude
fix: cap map0 kernel shmem for 256-expert MoE models
830eb54 
Graph reservation passes worst-case ne20=ne02 (256x256x2=128KB),
exceeding the 32KB threadgroup memory limit on Apple Silicon.
At runtime ne20 is the actual n_expert_used (e.g. 8), so shmem
= 256*8*2 = 4KB, well within limits.

Cap the reservation shmem to 32KB to prevent the assert from firing.
Tested on Qwen3.5-35B-A3B (256 experts) with llama-server + flash
attention — previously crashed during warmup, now runs at 22 t/s.

Fixes TheTom#58

Co-Authored-By: tturney@psyguard.ai
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
@TheTom @claude @signalnine
fix: AMD/RDNA4 arch dispatch — scalar half path for TQ4_1S on AMD GPUs
6862209 
The dp4a int8 kernel is optimized for NVIDIA Turing+ dp4a throughput
(240 t/s on 5090). On RDNA4, sudot4 has different throughput
characteristics and the q8_1 activation quantization adds overhead,
causing a regression vs the V12 float kernel (101 vs 135 t/s on
RX 9070 XT).

Fix: check GGML_CUDA_CC_IS_AMD(cc) at dispatch time and route AMD
GPUs to a scalar half-precision kernel (same pattern as TQ3_1S).
NVIDIA continues using the dp4a path.

Changes:
- Add mul_mat_tq4_1s_scalar_multi kernel: pre-rotated half activations,
  shmem centroid LUT, scalar dot product (no dp4a/byte_perm)
- Dispatch: use_dp4a = !AMD && TQ4_1S. AMD falls through to scalar path.
- LAUNCH_SCALAR macro unifies TQ4_1S/TQ3_1S scalar dispatch

Expected RDNA4 result: restore V12-level decode (135 t/s, 130% of Q8_0)
instead of dp4a regression (101 t/s, 60% of Q8_0).

Co-Authored-By: Tom Turney <tturney1@gmail.com>
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
@andyluo7
docs: add AMD Instinct MI300X (gfx942) ROCm test results
bf395f1 
TurboQuant KV cache compression (turbo2/turbo3/turbo4) builds and runs
correctly on AMD Instinct MI300X with ROCm 7.0.2. Zero code changes
required — existing CUDA kernels compile via HIP translation.

Test results (Qwen2.5-1.5B Q4_K_M, single MI300X):
- WHT roundtrip: PASS (max error 2.98e-07)
- turbo3 prefill: +3% vs f16 (25,200 vs 24,453 tok/s)
- turbo3 decode: 88% of f16 (160 vs 181 tok/s)
- turbo4 prefill: +4% vs f16 (25,427 vs 24,453 tok/s)
- turbo4 decode: 89% of f16 (161 vs 181 tok/s)

MI355X (gfx950) compiles but needs gfx950 added to llama.cpp's
MMQ kernel dispatch (upstream issue, not TurboQuant-specific).

Tested-by: Andy Luo <andyluo7@users.noreply.github.com>
@andyluo7
feat: add CDNA4 (gfx950/MI355X) support + test results
e88018d 
Add AMD Instinct MI355X (gfx950) architecture support:

Code changes:
- vendors/hip.h: Add CDNA4 define for __gfx950__, include in CDNA family
- common.cuh: Add GGML_CUDA_CC_CDNA4 constant and IS_CDNA4 macro
- mma.cuh: Route CDNA4 to compatible MFMA instructions
  * bf16: mfma_f32_16x16x16bf16_1k (same as CDNA3)
  * int8: mfma_i32_16x16x32_i8 (same as CDNA3)
  * f32: mfma_f32_16x16x4f32 (CDNA2 path, NOT xf32 which doesn't exist on gfx950)
- mmq.cuh: Include CDNA4 in stream-k dispatch
- common.cuh: Exclude CDNA4 from CDNA3-specific e4m3_fnuz FP8 path (gfx950 uses standard e4m3fn)

MI355X test results (Qwen2.5-1.5B Q4_K_M, single GPU):
- turbo3: 39,140 tok/s prefill (98% of f16), 162 tok/s decode (64%)
- turbo4: 39,232 tok/s prefill (98% of f16), 214 tok/s decode (84%)
- WHT roundtrip: PASS (max error 2.98e-07)

Note: non-FA MMQ path crashes on gfx950 (xf32 MFMA unsupported).
TurboQuant types force FA and work correctly.

Tested-by: Andy Luo <andyluo7@users.noreply.github.com>
@TheTom
perf: TQ4_1S native kernel 3.5× faster + AMD arch dispatch (TheTom#57)
a4e8af4 
perf: TQ4_1S native kernel 3.5× faster — 240 t/s, less VRAM than q8_0 conversion
@Tuklus @claude
feat: Vulkan compute shader support for turbo3 KV cache
c2a70f3 
Full turbo3 quantize/dequant pipeline for Vulkan backend:

- types.glsl: block_turbo3_0 struct (norm + qs[8] + signs[4])
- dequant_turbo3_0.comp: standalone dequant shader (3-bit index
  reconstruction from 2-bit qs + 1-bit signs, centroid lookup)
- dequant_funcs.glsl: inline dequant for get_rows/mul_mat paths
- dequant_funcs_cm2.glsl: cooperative matrix 2 FA path support
- copy_to_quant.comp: quantize function with norm correction
- vulkan-shaders-gen.cpp: turbo3_0 type registration
- ggml-vulkan.cpp: pipeline creation and supports_op dispatch

Tested on AMD 7900 XTX (RADV): 243 pp / 25.8 tg t/s with turbo3 KV.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
@TheTom
Merge pull request TheTom#33 from apollosenvy/pr/vulkan-turbo3
eea498c 
feat: Vulkan compute shader support for turbo3 (experimental)
@TheTom @claude
metal: add TurboFlash attention kernel for turbo3 KV cache decode
0d6b38a 
Two-pass block-parallel attention kernel optimized for turbo3 V cache
decode on Apple Silicon Metal. Supports both q8_0-K (asymmetric) and
turbo3-K (symmetric) configurations via compile-time function constant.

Architecture:
- Pass 1: 32-thread SIMD group per (query-head, block) pair
  - Each lane handles DK/32 interleaved dimensions
  - Q loaded to per-lane registers, K dequant via q8_0 or turbo3 path
  - K scoring via simd_sum dot product
  - turbo3 V unpack with register codebook (8 centroids)
  - Online softmax (m/l/o state) entirely in registers
  - Zero shared memory in pass 1
- Pass 2: merge partial results across blocks
  - Online softmax correction with global max/sum
  - Inverse WHT via simd_shuffle_xor (stages 0-4) + shared memory (stages 5-6)
  - Eliminates 5 of 7 threadgroup barriers vs naive butterfly

Auto-detection: activates for single-token decode (ne01==1) when V is
turbo3 and K is q8_0 or turbo3. Controllable via TURBO_FLASH env var
(0=off, 1=force). Block size B=64 (proven optimal on Apple Silicon).

Benchmarks (Qwen2.5-7B Q8_0, asymmetric q8_0-K/turbo3-V):
- M5 Max 128GB: +1.5% decode at 8K (56.82 vs 56.00 tok/s), 93% of q8_0
- M2 Pro 32GB:  +0.6% decode at 8K (20.55 vs 20.42 tok/s)
- Advantage scales with context (+7.3% at 32K)

Inspired by Eric Kryski's TurboFlash architecture (mlx-swift-lm).

Co-Authored-By: tturney@psyguard.ai
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
@TheTom TheTom force-pushed the feature/turboquant-kv-cache branch from 10cb187 to 0d6b38a Compare April 8, 2026 23:49
TheTom and others added 6 commits April 9, 2026 08:20
@TheTom
Merge pull request TheTom#61 from andyluo7/add-rocm-mi300x-support
157cb85 
feat: AMD Instinct MI300X + MI355X (gfx942/gfx950) ROCm support
@Titaniumtown
vulkan: fix and complete turbo3 KV cache support
1ebda4c
@Titaniumtown
vulkan: add turbo3 backend tests
6a29b58 
- test_turbo_wht: forward/inverse WHT, 18 configs. NMSE tolerance 1e-5
  (f32 SIMD reduction order varies across GPU backends).
- test_turbo_wht_roundtrip: forward then inverse recovers original, 9
  configs. NMSE tolerance 1e-5.
- test_set_rows_turbo3: full quantization round-trip at small and large
  tensor sizes. Large tensors exercise the 2D dispatch grid. 21 configs.
- Existing: test_turbo_wht (18), FA with turbo3 KV (528).
- Total: 576 tests.
@TheTom
Merge pull request TheTom#62 from Titaniumtown/pr/turboquant-vulkan-work
8590cbf 
vulkan: fix and complete turbo3 KV cache support
@terrysimons
fix: HIP/ROCm compatibility — CUDA_CHECK errors, guard D>=576 MMA
57daafb 
Re-applied from earlier fix; lost after branch rebase.

Add CUDA_CHECK() to all cudaMemcpyToSymbol/cudaMemcpyFromSymbol calls in
the InnerQ calibration path. On HIP, unchecked errors from these calls
are sticky and poison the runtime, causing subsequent API calls to fail
with 'no ROCm-capable device is detected'.

Also guard the D>=576 MMA flash attention dispatch and kernel selection
with #ifndef GGML_USE_HIP, matching the existing D>=576 tile exclusion
(these kernels exceed HIP's shared/local memory limits).

Tested on: ROCm 6.4.4, gfx1151 (AMD Ryzen AI Max+ 395 / Strix Halo)
@terrysimons
fix(cuda): set explicit cuBLAS/hipBLAS workspace to stop VRAM creep
e844a68 
Addresses upstream llama.cpp issue ggml-org#19979: on ROCm/HIP, VRAM creeps up
during long prompt processing when using non-f16 KV cache types (q4_0,
q5_1, q8_0, turbo2/3/4), eventually causing OOM. Does not occur with
f16/f16 KV cache.

Root cause: cublasCreate() is called but cublasSetWorkspace() is not,
so rocBLAS lazily allocates its own internal workspace on first GEMM
and grows it monotonically as it encounters new Tensile kernel shapes
across the varying batch sizes / KV cache lengths of prompt processing.
The growth is invisible to llama.cpp's memory accounting (hipMemGetInfo
snapshots it but the fit algorithm runs before the workspace exists).

Fix: after cublasCreate(), allocate a fixed 32 MiB buffer via cudaMalloc
and install it as the explicit workspace via cublasSetWorkspace (aliased
to hipblasSetWorkspace on HIP). rocBLAS then uses our workspace instead
of growing its own, and the workspace is freed deterministically in the
backend context destructor.

Added:
 - GGML_CUDA_CUBLAS_WORKSPACE_SIZE constant (32 MiB) in common.cuh
 - cublas_workspaces[] array in ggml_backend_cuda_context
 - cublasSetWorkspace -> hipblasSetWorkspace alias in vendors/hip.h
 - Workspace alloc + install in cublas_handle()
 - Workspace free in ~ggml_backend_cuda_context() (after cublasDestroy)

Upstream reference:
  ggml-org#19979

The 32 MiB size matches NVIDIA cuBLAS recommendations for typical
workloads. Bump if you see rocBLAS falling back to smaller kernels.

Affects all ROCm GPU archs, especially GFX11 (RDNA3/3.5) where Tensile
solution selection triggers more workspace growth.
@terrysimons terrysimons force-pushed the feature/turboquant-kv-cache branch from d17f1dc to e844a68 Compare April 11, 2026 21:23
@TheTom TheTom force-pushed the feature/turboquant-kv-cache branch from 45f8a06 to 1073622 Compare April 16, 2026 01:14
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

ggml Nvidia GPU

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

9 participants

@terrysimons @stragulus @TheTom @seanrasch @signalnine @Tuklus @iamwavecut @andyluo7 @Titaniumtown