Bug Report: kvarn6 KV-quantization causes massive performance regression on ROCm (RX 7900 XTX)

[raufaser]

Issue: Using KV-quantization kvarn6 results in extremely slow performance compared to q8_0 on AMD ROCm hardware.

Performance Comparison

Both tests used the same prompt (10,410 tokens) and identical settings except for KV quantization.

Note: The kvarn6 run was canceled early due to excessive runtime.

Metric	q8_0	kvarn6	Regression
Prompt Processing	821.19 tokens/sec	53.79 tokens/sec	~15.3x slower
Text Generation	29.52 tokens/sec	1.34 tokens/sec	~22x slower

This is a severe performance regression that makes kvarn6 practically unusable.

Steps to Reproduce

1. Use llama.cpp build 11192 (ee8754c57) with ROCm backend on Windows
2. Load a large model (tested with Qwen3.6-27B-Fable-Fusion-Q4_K_M)
3. Run with full offload (-ngl 99) and flash attention (-fa 1)
4. Compare these two configurations:

q8_0 (baseline):

llama-server.exe -m <model> -ngl 99 -fa 1 -ctk q8_0 -ctv q8_0 -c 40000

kvarn6 (regression):

llama-server.exe -m <model> -ngl 99 -fa 1 -ctk kvarn6 -ctv kvarn6 -c 40000

System Information

• GPU: AMD Radeon RX 7900 XTX (24 GB VRAM)
• CPU: AMD Ryzen 9 7900X 12-Core
• RAM: 64 GB
• OS: Windows
• Backend: ROCm (HIP), NO_VMM = 1
• Build: 11192 (ee8754c) with Clang 20.1.8 for Windows x86_64
• Model: Qwen3.6-27B-Fable-Fusion (Q4_K_M, 65 layers, 27.32B params, 17.22 GiB)

Relevant Log Excerpts

q8_0 performance:

prompt processing, n_tokens = 10410, progress = 1.00, t = 11.93 s / 872.56 tokens per second
prompt eval time = 12681.67 ms / 10414 tokens (1.22 ms per token, 821.19 tokens per second)
eval time = 73468.79 ms / 2169 tokens (33.87 ms per token, 29.52 tokens per second)

kvarn6 performance (run canceled after ~5.5 minutes):

prompt processing, n_tokens = 10410, progress = 1.00, t = 193.54 s / 53.79 tokens per second
n_decoded = 144, tg = 1.34 t/s

Expected Behavior

kvarn6 should provide competitive or better performance than q8_0 while using less VRAM, or at least not be an order of magnitude slower.

Actual Behavior

kvarn6 is 15-22x slower than q8_0 on ROCm hardware, making it impractical for use.

Additional Notes

• This appears to be ROCm-specific; the kvarn6 implementation may not be optimized for AMD GPUs
• The issue affects both prompt processing and token generation equally severely
• The kvarn6 run was manually canceled due to excessive runtime; actual degradation may be even more severe over longer generations
• Full log level 5 logs for both q8_0 and kvarn6 runs are attached to this issue

log_kvarn6_lv5.txt
log_q8_0_lv5.txt

[raufaser]

The problem still exists. kvarn is unusable with ROCm.

[Anbeeld]

Sorry, I do keep this in mind. It's just the past week or so was very saturated, and also I was trying to fix tensor split instability on CUDA but unfortunately just wasted a lot of time... for now

[gashap0n]

• This appears to be ROCm-specific; the kvarn6 implementation may not be optimized for AMD GPUs

Doesn't seem to be an exclusive ROCm issue. I managed to achieve the same behavior using the same GPU but with Vulkan as backend. In my case, I was running Qwen 3.8 27B:

[Qwen3.8-27B]
model = files/Qwen3.8-27B-UD-Q4_K_XL.gguf
port = 8080
ctx-size = 131072
batch-size = 2048
ubatch-size = 512

n-gpu-layers = 999
flash-attn = on
reasoning = on

spec-type = draft-mtp
spec-draft-n-max = 3

cache-type-k = kvarn5
cache-type-v = kvarn5

no-mmap = 1
parallel = 1
cont-batching = 1
kv-unified = 1

temp = 1.0
top-p = 0.95
top-k = 20
min-p = 0.0
presence-penalty = 1.5
repeat-penalty = 1.0

reasoning-format = deepseek
jinja = true

With q8_0/q8_0, tokens per second hover around 45-50 regardless of the context window, nonetheless, kvarn quantization drops the generation to 15. I tested both with and without draft model, and I'm still experiencing the drop.

Edit: I'm on Linux

[raufaser]

Chances are high that we encounter two different problems which cause similar outcomes. Your decoding speed was 3x slower, mine was 22x slower.

[raufaser]

I did some investigation assisted by AI and I want to share my findings with you. Perhaps you already know all of that and since I do not really no what that all means, the findings can be completely bogus.

---

Testing this on RX 7900 XTX (gfx1100, Windows x64 HIP release build) — found the cause of the slowdown, plus a separate crash worth tracking.

1. The 15-22x slowdown is the expected portable-native fallback, not a broken kernel

The model here is head-dim 256 (n_embd=5120, n_head=24, n_head_kv=4, gqa=6). GGML_KVARN_DEBUG_ROUTES=1 shows every layer on the fallback:

kvarn-route backend=HIP cc=16781568 wave=32 D=256 k=6 v=6 gqa=6 nq=2 nkv=256
route=portable-native entry=compact-tail fallback=generic_shape_rejected

Why: on HIP the specialized decode routes (split/vector) are disabled in fattn-kvarn-route-policy.h, and the WMMA eligibility check rejects head_dim > 128 (only CDNA MFMA allows 256). So all D=256 KVarN layers on RDNA3 run the 128-thread sequential portable kernel — correct but ~15x (pp) / ~22x (tg) slower than q8_0 at long context. ROCm docs confirm the HW can do f16 WMMA on gfx1100, so this is a tile-coverage gap, not hardware. Fix direction: extend WMMA to D=256, or port the CUDA split/vector decode to HIP.

2. Separate bug: prompt-generic-mma hangs the driver on D=128, nq ≥ ~256

With a D=128 model (Qwen3-4B-Thinking-2507, n_embd=2560, n_head=32, n_kv=8, gqa=4), KVarN correctly takes the fast route, but crashes at -p 256:

bench	result
-p 128 -n 32	pp128 2129 t/s, tg32 147 t/s — stable
-p 256 -n 64	pp256 printed with ±53% variance, then driver reset (WDDM TDR, no Event Viewer entry)
-p 256 -n 1	same crash

Crash route line: route=prompt-generic-mma ... nq=256 nkv=256. Identical commands with -ctk q8_0 -ctv q8_0 are stable (pp512: 4732 t/s), so the fault is in the KVarN prompt MMA kernel, unrelated to #122's portable fallback. Threshold is between pp128 (ok) and pp256 (crash).

Suggest splitting this issue: D=256 RDNA portable fallback (the perf report above) stays here as a roadmap item; the D=128 prompt-generic-mma crash should be a separate bug. Happy to bisect further (nq=192/224, kvarn4/5/6, GGML_KVARN_TEST_FORCE_PORTABLE_FATTN=1) if useful.

[johny-mnemonic]

I can confirm it's not just about ROCm. I am running Vulkan build on RDNA2 (RX6800) and here are my results:
beellama-v0.4.5 q8/q5.1

| model           |            test |           t/s |     peak t/s |        ttfr (ms) |     est_ppt (ms) |    e2e_ttft (ms) |
|:----------------|----------------:|--------------:|-------------:|-----------------:|-----------------:|-----------------:|
| qwen3.6-35b-a3b | pp2048 @ d65536 | 217.78 ± 0.00 |              | 281956.77 ± 0.00 | 281953.64 ± 0.00 | 281956.77 ± 0.00 |
| qwen3.6-35b-a3b |   tg32 @ d65536 |  32.22 ± 0.00 | 33.00 ± 0.00 |                  |                  |                  |

beellama-v0.4.5 kvarn5/kvarn4

| model           |            test |          t/s |     peak t/s |         ttfr (ms) |      est_ppt (ms) |     e2e_ttft (ms) |
|:----------------|----------------:|-------------:|-------------:|------------------:|------------------:|------------------:|
| qwen3.6-35b-a3b | pp2048 @ d65536 | 48.01 ± 0.00 |              | 1278308.22 ± 0.00 | 1278300.26 ± 0.00 | 1278308.22 ± 0.00 |
| qwen3.6-35b-a3b |   tg32 @ d65536 |  9.63 ± 0.00 | 12.00 ± 0.00 |                   |                   |                   |

Edit: I have interesting update.
I wanted to give up on kvarn for now and just use q8/q5.1 like normally, but with the added --kv-tail-tokens 1024. It tanked the performance to the same level as using kvarn cache. So seems like once I try to use these new kv cache changes it goes different code path and tanks the performance completely. Most probably it is using something specific for NV GPUs and hence doesn't work well with AMD.

[raufaser]

Status update (RX 7900 XTX, gfx1100):

1. The 15–22x gap — root cause confirmed, fix exists

As in my Sep 2 analysis, every D256 KVarN layer fell back to portable-native (fallback=generic_shape_rejected), because RDNA WMMA rejected head_dim > 128.

While enabling that path I found a real race on top: flash_attn_ext_f16_process_tile (ggml/src/ggml-cuda/fattn-mma-f16.cuh) is missing a __syncthreads() between k00 combine iterations whenever np == 1 with multiple iterations. Fast warps overwrite shared tile_Q while lagging warps still read it — so the newly enabled D256 path produced garbage, not just slow output.

The complete fix is 3 files:

File	Change
fattn-mma-f16.cuh	generalized barrier for all np==1 multi-iteration configs
fattn-kvarn-route-policy.h + device invariant	RDNA >128 → >256 eligibility
test-cuda-fattn-route-policy.cpp	test update (D256 now eligible, D512 still rejected)

Post-fix, short context (Qwen3.6-27B):

	pp	tg
q8_0	791	33.9
kvarn6	615 (−22%)	34.5 (flat)

The 15–22x is gone here — but I have not retested your exact setup (llama-server, -c 40000, 10k prompt). Please re-verify there before closing.

2. New finding: watch quality before shipping this on ROCm

KLD vs f16/bf16, Wikitext-2, same files/settings throughout:

ctx	kvarn6 KLD	PPL vs base
512	0.19	3.96 vs 3.42
32k	2.14	44.9 vs 6.3 ⚠️
64k	2.22	47.1 vs 6.3 ⚠️

But: forced-portable (GGML_KVARN_TEST_FORCE_PORTABLE_FATTN=1) gives KLD 0.004 on identical records, and 100+ isolated sweep configs (all bit pairs, tails, eager, 24/4 heads, n_kv up to 16k) all pass at RMSE ≤1e-3 with no bias.

So compression and kernel math are per-call exact — yet the WMMA prompt path carries ~4e-4/call noise (vs ~1e-5 portable) that accumulates over layers and context. Cross-check on RTX 4090 + same weights shows KLD 0.002, but note its prompt bypasses KVarN dispatch entirely, so that number flatters the kernels.

Bottom line: the fast path fixes the speed but is currently lossy on ROCm. It needs WMMA prompt-tile accuracy work before becoming the default.

3. Other reports in this thread

@gashap0n's Vulkan 3x and @johny-mnemonic's "--kv-tail-tokens 1024 tanks even q8_0 to kvarn levels" both point at AMD fallback/tail paths rather than KVarN math — same family as this issue. The D128 prompt-generic-mma driver hang from my Sep 2 comment is still a separate bug worth splitting out, IMO.

Happy to post the patch, GGML_KVARN_DEBUG_ROUTES=1 logs, and full KLD tables if useful.

[gashap0n]

@raufaser Could you create a pull request with the changes? That way I can test whether the regression still exists with my setup. Thanks!

[raufaser]

Status update (2026-09-07, RX 7900 XTX, v0.4.5 base)

Re-measured with Qwen3.6-27B-Q5_K_S, -fa 1 -ngl 99 -p 4096 -n 1028
(llama-bench):

cache	pp4096	tg
q8_0	804 t/s	33.4 t/s
kvarn4	107 t/s	30.3 t/s

Decode is fixed
The 22x tg collapse from the original report is gone.
30 vs 33 t/s is ordinary weight-bound parity.
That half was a D256 k00-combine race in the WMMA flash-attention kernel: shared staging memory is reused across loop iterations without a barrier when running single-warp-column configs (2 iterations at D256 with a 64-column tile), so fast warps overwrote data lagging warps still read. Fixed with the missing barrier, plus enabling the RDNA D256 route that was previously rejected before kernel launch.

Prefill turned out to be a second, separate bug (quality, not speed)
Measured GPU-native vs CPU-materialized attention error directly (D256, 6-bit K/V, production head/group layout, 256-query prompt tiles): the RDNA3 WMMA value-accumulator is fp16, giving ~3e-4 error per call vs ~1e-5 for an fp32 reference path, independent of sequence length.
It compounds through 64 layers: KLD 0.186 at 512 tokens, PPL 44.9 vs 6.3 at 32k context. (CUDA never showed this — its prompt path doesn't use the WMMA kernel.)

Fix being validated
route the HIP prompt-prefill through the exact fp32 direct-record path (same compressed records, exact math); decode stays on WMMA.
Perplexity confirmation (wikitext-2, -c 4096):
Q5_K_S f16 PPL 6.6375 vs kvarn6 6.6298; second weight file (UD-Q4_K_XL)
f16 6.6522, q8_0 6.6527, kvarn6 6.6269.
Quality at parity everywhere.

What remains:
prefill speed. The exact path has no tensor-core throughput, so prefill runs ~7.5x slower than q8_0 here (107 vs 804 t/s; decode unaffected).
The issue is now purely "prefill too slow", no longer "wrong + slow". Options from here: an fp32 WMMA accumulator (full recovery, needs RDNA3 fragment-layout work) or a reworked kernel that shares K tiles across queries.

[raufaser]

@gashap0n
I'm not proposing this for merge (prefill is still ~7x slower than q8_0, see my status update above; this is a correctness fix with a known perf cost, not shippable as-is).
But you're welcome to test the branch directly:

git remote add raufaser https://github.com/raufaser/beellama.cpp
git fetch raufaser kvarn-rocm-prompt-precision
git checkout raufaser/kvarn-rocm-prompt-precision

That's 4 commits on top of v0.4.5 (593e6d590).

What I'd like to know from your setup:
decode t/s (expect ~parity with q8_0)

and

perplexity or KLD vs f16 (expect parity), plus your GPU/VRAM.

Thanks for testing!

[raufaser]

Prefill update (same setup, branch kvarn-rocm-prompt-precision)

The exact portable prompt path got 2.7x faster with no accuracy change:

cache	pp4096 before	pp4096 now	tg
kvarn4	107 t/s	291 t/s	33.5 (unchanged)
q8_0	804 t/s	—	33.4

What changed (all in the portable direct-record kernel):

• 8 queries per block share one K/V token stream (was: one block per query, every block re-reading all records)
• per-token reduction via warp shuffles instead of a shared-memory tree
• token index math resolved once per block instead of once per thread

Perplexity still at parity (Q5_K_S f16 6.6375 vs kvarn6 6.6298).

Remaining gap to q8_0 (~2.8x) is raw per-element dequant work, that needs either a tensor-core (fp32 WMMA accumulator) path or a kernel tiling project.

Reporting for the record.