KVARN KV Performance optimizations

[valujin]

I've made lot of work to significantly improve Kvarn KV Cache performance on my RTX 3090 for Qwen dense and MoE models, especially for concurrent requests without losing quality. Right now, the performance is on par with Q8 KV and Kvarn KV is faster on large context than Q8 KV. I hope, you will find my contribution useful. Since I've done it for myself, currently all docs and comments on my changes are in Russian.

https://github.com/valujin/beellama-kvarn

[Anbeeld]

Looks interesting, I'll check it out. Thanks!

[kirstenjg]

Please look into this, I just tested and my results are even better than OP's.
Just confirm that his implementation does not affect KLD in any significant way.

[Arien02]

I second this too, tested with rtx3090 and Qwen3.8-27B-UD-Q4_K_XL.gguf and I can confirm it's the fastest llama implementation so far.
And another discovering, latest v0.4.4 it's slower than v0.4.3 for my setup, don't know why but something it's loosing TPS...

[Anbeeld]

Ported most of the changes, additionally adapting them to v0.4.4 changes that landed closer to release.

Performance is indeed better, especially with multiple slots.

Please test it thoroughly on the v0.4.5 branch, especially in terms of stability, correctness, and edge cases.

AI-written implementation details and benchmarks

Completed a port of the KVarN correctness and CUDA decode work from valujin/beellama-kvarn onto the current v0.4.5 architecture.

The main port landed in 72bace072 (Optimize KVarN CUDA decode paths). The final long-context compact-plan, windowing, and stage-slot follow-up described below is currently in the local candidate and has not yet been committed.

The implementation was adapted rather than cherry-picked because the optimization fork was based on an older v0.4.4 tree and its allocator, stage ownership, cache state, and CUDA routing differ materially from current v0.4.5.

---

What was implemented

1. Compact-plan lifecycle correctness

The cached indirect KVarN read plan is now invalidated whenever its source cache metadata can change:

• llama_kv_cache_kvarn_context::next()
• llama_kv_cache_kvarn_context::apply()

Previously, a plan could survive across ubatches even though the occupied and pending physical cells had changed. That could silently produce incorrect reads depending on slot count and scheduling.

2. Safe GQA-6 CUDA query loading

The CUDA MMA path logically supports six GQA rows, but ldmatrix physically reads an eight-row extent.

The query tile now:

• reserves storage for all eight physical rows;
• initializes the padded rows;
• retains six logical GQA heads.

This prevents out-of-bounds or uninitialized shared-memory reads without changing logical attention dimensions.

3. Compact-plan duplicate bitmap

Compact plan construction now uses a capacity-sized bitmap instead of std::set for duplicate filtering.

It preserves:

• occupied/pending deduplication;
• first-occurrence semantics before physical layout;
• capacity checks;
• sentinel padding.

4. Physical record-group-safe compact plans

Unified-cache ownership can interleave logical sequences while their KVarN records remain physically disjoint. A logically ordered plan could therefore look contiguous to a CUDA cooperative record loader even though its adjacent rows came from different physical records.

The final plan implementation:

• groups cells by their physical 128-cell KVarN record;
• sorts selected records by physical group;
• gives every selected record its own monotonic segment;
• retains -1 holes within incomplete records;
• never falls back to interleaved logical order for sparse groups.

This is the change responsible for recovering the large two-slot performance improvement.

5. Descriptor setup optimizations

The descriptor path now includes:

• 1024-thread descriptor reduction;
• low-32-bit physical-cell extraction;
• preservation of high-bit encoded stage-slot metadata;
• division-free scanning for the common single-stream case.

The old fork treated indices more like plain physical cells. The port explicitly preserves the current v0.4.5 encoded stage-slot representation.

6. CUDA decode-kernel optimizations

The following mechanisms were ported and adapted:

• packed fragment permutation;
• paired low-bit unpack;
• fully masked split skipping;
• half-width shared quantization axes;
• split-128 geometry;
• launch bounds and occupancy-aware selection;
• query-row tiling with Q_TILE=2/3;
• reuse of decoded K/V data across query rows.

Unsupported shapes continue to use the descriptor-native MMA fallback rather than failing or silently changing cache interpretation.

7. Small verification-batch routing

Split decode is now eligible through n_q <= 8.

This covers common speculative verification shapes while preserving generic MMA fallback for larger batches. Multi-query split routing requires an attention mask.

A temporary n_q <= 3 experiment was rejected because it did not provide the same stability/performance balance.

8. Evidence-based split geometry selection

The final selector uses:

• split-64 for ordinary single-query decode below 32K;
• split-128 for single-query decode at 32K and above;
• split-128 earlier when multi-query verification benefits from query tiling.

This replaced a broader initial split-128 policy that regressed ordinary-depth workloads.

9. Bounded large-window materialization

The default KVarN materialization window was reduced from 65,536 to 32,768 tokens.

This bounds transient K/V scratch when two long prompts share a 128K cache. The partial windows retain the online-softmax maximum, numerator, and denominator and are merged afterward.

Chunked-window CUDA parity coverage was added. The accepted RMSE is 3e-4, reflecting the changed floating-point reduction order.

10. Explicit replacement-stage slots in delayed sealers

The delayed high-shared, low-shared, head-wide, direct, and workspace sealing paths now use the incoming replacement group’s encoded host-assigned stage slot.

They no longer assume that a live group’s slot can always be reconstructed from its group number modulo the tail depth.

This did not explain every observed long-prefill failure because the production target uses eager record commitment, but it is a real correctness defect in the delayed paths and is retained with focused CPU/CUDA parity tests.

---

How this differs from the initial implementation

The first adaptation was deliberately conservative:

• preserve first-seen logical group order;
• align only small n_q=1..8 decode/verification plans;
• cap alignment growth at 25%;
• leave large prefill and sparse plans dense.

That version improved ordinary two-slot decode, but it did not fully enforce the CUDA record-loader invariant.

Initial approach

The initial policy assumed that dense fallback was safe when alignment expansion was large. This proved false: a sparse logical plan could still place rows from separate physical records next to each other and make them look like a contiguous physical tile.

It also retained logical group order to minimize observable ordering changes.

Final approach

After reproducing failures with larger two-slot contexts, the plan was changed to:

1. sort by physical record group;
2. reserve one complete monotonic segment per selected record;
3. preserve holes with -1;
4. remove the small-batch-only alignment helper;
5. remove the 25% sparse-growth fallback;
6. apply the physical record invariant to every compact unified-cache plan.

The API was simplified again after the growth policy was removed.

Other changes made during diagnosis

• A startup crash caused by consulting base()->get_ubatch() while constructing synthetic reserve graphs was fixed by using valid current slot metadata. The policy that needed that query was later removed entirely.
• The original 65,536-token materialization window was reduced to 32,768 after full-window tests showed excessive transient scratch at 128K.
• Alignment, exact tail, CUDA graph use, low-shared-memory stores, ubatch size, and portable attention were tested independently to separate compact-plan failures from other long-run failures.
• The final two-slot benchmark was redesigned so both slots were actually decoding simultaneously. The original cold-prompt harness measured one slot decoding while the other was still prefilling, which substantially understated two-slot decode throughput.

---

What was not implemented

1. The old fork allocator/group-claim implementation

The implementations from the old group-claim and allocator commits were not copied.

Current v0.4.5 already has newer:

• allocator ownership;
• planning reservations;
• stage-slot assignment;
• cache state and checkpoint handling.

Copying the old allocator would have replaced current ownership semantics and risked regressions in rollback, state restore, MTP verification, and slot reuse.

Only the physical read-plan mechanism was adapted.

2. The old fork stage-owner implementation

The old modulo/owner stage system was not ported.

Current v0.4.5 uses explicit host-assigned stage slots encoded in cache indices. The CUDA changes were adapted to consume those assignments rather than restoring the old ownership model.

3. Wholesale cherry-picking of 54bc3d29, c3ca5bcc, 432024c3b, or 1b2f628ca

These commits bundled mechanisms with obsolete allocator and state assumptions. They were used as architectural evidence, not treated as drop-in patches.

---

Additional changes not present in the optimization fork

The following work was added because it was required by current v0.4.5 or found during validation:

1. GQA-6 eight-row physical query padding.
2. Low-32-bit physical-cell decoding while preserving encoded stage slots.
3. 32K bounded window materialization for 128K unified caches.
4. Explicit replacement-slot handling in every delayed CUDA sealer.
5. The measured 32K split-64/split-128 selector boundary.
6. Current-architecture stage-slot parity tests.
7. Chunked-window online-softmax parity tests.
8. Static source-safety tests for encoded indices, route limits, and removed experimental switches.
9. A synchronized hot-cache two-slot benchmark harness. It preloads both long prompts, then starts both measured requests with sub-0.2 ms launch skew and verifies that each request reuses the full cached prompt and reprocesses only one token.

---

Verification

The final candidate was checked with:

• a complete Windows CUDA sm_86 build;
• fresh serialized CTest: 82/82 passed;
• earlier expanded/all-test coverage: all 104 tests passed across the full run plus one isolated retry of a transient unrelated CUDA test;
• test-kvarn CPU/CUDA parity;
• direct CUDA FlashAttention parity;
• KVarN metadata, exact-tail, stage-slot, rollback, state, and route tests;
• n_q=2..16 route coverage;
• all 36 ordered KVarN CUDA bit pairs in the expanded build;
• 169 standard CUDA quantization pairs;
• static safety tests;
• git diff --check.

---

Benchmark environment

Unless noted otherwise:

• GPU: RTX 3090, CUDA0, sm_86;
• CUDA: 13.3;
• driver: 616.56;
• model: Qwen3.6-27B IQ3_XXS;
• model SHA-256: a6cd1e829a8ba24a78edeb8f078b724ea7a9873856f1916e356eea5c537c0b40;
• v0.4.4 package: commit f8cd4e6dd;
• v0.4.5 package: commit 476aa6c4b;
• final two-slot candidate: 476aa6c4b-dirty, including the follow-up physical-plan/window/stage work;
• batch/ubatch for the standard decode and prefill matrix: 2048/512;
• three samples unless noted.

The older published fork baseline was commit a749684, not the provided v0.4.4 package. The tables below use the provided f8cd4e6dd package.

---

v0.4.4 versus v0.4.5 results

16K single-slot decode

64 generated tokens:

Cache	v0.4.4	v0.4.5	Difference
BF16, tail 0	37.420	37.035	−1.03%
KVarN4, tail 0	33.042	34.451	+4.26%
KVarN4, tail 1024	32.042	33.470	+4.46%
Q4_0, tail 0	39.222	39.183	−0.10%
Q4_0, tail 1024	38.175	38.399	+0.59%

The controls show that the ordinary-depth gain is specific to KVarN rather than a general build-wide speedup.

16K prefill

Cache	v0.4.4	v0.4.5	Difference
BF16, tail 0	1244.96	1238.29 tok/s	−0.54%
KVarN4, tail 0	1176.11	1177.13 tok/s	+0.09%
KVarN4, tail 1024	1166.07	1166.44 tok/s	+0.03%
Q4_0, tail 0	1234.33	1234.69 tok/s	+0.03%
Q4_0, tail 1024	1205.29	1205.23 tok/s	−0.005%

There is no material prefill change at 16K.

64K single-slot decode

Tail 1024, 64 generated tokens:

Cache	v0.4.4	v0.4.5	Difference
KVarN4	18.347	20.036 tok/s	+9.21%
KVarN5	18.368	19.822 tok/s	+7.92%

The long-context gain is larger than the 16K gain, consistent with the descriptor and split-geometry changes.

Single-slot MTP

KVarN4, tail 1024, 16K context, three trials:

Build	Decode
v0.4.4	65.766 tok/s
v0.4.5	66.965 tok/s
Difference	+1.82%

---

Proper two-slot long-context comparison

The earlier cold-prompt server benchmark did not isolate simultaneous two-slot decode: one slot often decoded while the other was still prefilling.

The final benchmark therefore used:

• --parallel 2;
• --kv-unified;
• 131,072-token cache;
• 63,692 cached prompt tokens per slot;
• F16 precision tail 1024;
• 256 measured tokens per slot;
• both prompts preloaded before measurement;
• one token reprocessed per measured request;
• launch skew below 0.19 ms;
• three valid measured trials for every build/cache combination.

Aggregate server-reported decode

Cache	v0.4.4	Final local v0.4.5	Difference
KVarN4	6.497 ± 0.114	29.346 ± 0.013 tok/s	4.52× / +351.7%
KVarN5	6.052 ± 0.035	29.298 ± 0.548 tok/s	4.84× / +384.1%

Average per slot

Cache	v0.4.4	Final local v0.4.5
KVarN4	3.249	14.673 tok/s
KVarN5	3.026	14.649 tok/s

Wall-clock aggregate throughput

This includes request and synchronization overhead:

Cache	v0.4.4	Final local v0.4.5	Difference
KVarN4	6.368	26.615 tok/s	4.18× / +317.9%
KVarN5	5.916	26.103 tok/s	4.41× / +341.2%

Long-prompt prefill in the same trials

Cache	v0.4.4	Final local v0.4.5	Difference
KVarN4	680.14	672.98 tok/s	−1.05%
KVarN5	673.42	672.40 tok/s	−0.15%

The large two-slot improvement is therefore a decode-path result, not a general prefill change.

The result does not exactly reproduce the fork’s published 5.9× Qwen3.8 result, but it confirms the same mechanism on a matched comparison against the provided v0.4.4 package: the final local build is approximately 4.5–4.8× faster with both long-context slots actively decoding.

[valujin]

I and my AI assistant have made farther mixed optimizations and fixes that improve QoL significantly for 24GB card users). Some of them are KVarN related (like MTP KV cache support and optimizations), some of them not (like VRAM free VISION with almost the same performance).

Update: what's new since the ported batch

All measured on RTX 3090 / Qwen3.8-27B-UD-Q4_K_M, ctx 212992, 4 slots, kvarn4 target cache.

Bug fixes — the first two are not KVarN-specific and likely affect the base as well

1. Speculation checkpoint rollback was quadratic. get_seq_rm_capability() declared a 1-token reserve that can_seq_rm() does not actually enforce for a structural cache. The server saw the capability as insufficient and re-serialized the entire draft cache on every speculation cycle, through a tail serializer that walks cells from zero per entry. perf showed state_tail_payload_slots at 17.95% of cycles with GPU at 63%. After the fix: 42.7 → 59.4 t/s at 20k depth, GPU 63% → 99%.

2. Draft rollback was computed on the wrong axis. The MTP draft numbers its tokens by count (n_past), while the server required rollback to ckpt.pos_max + 1 — the target's position axis. These are identical on text, so it never showed. On vision models they diverge: one image is thousands of tokens but only a few positions under M-RoPE. Result: GGML_ABORT on the first image whenever the draft cache is quantized. Fixed by using ckpt.n_tokens. Verified on 15 image requests across 5 configurations, byte-identical output, zero seq_rm refusals.

3. Shared-memory OOB in the KVarN decode kernel: p_sh is declared with 6 rows at MAX_GQA == 6, but load_ldmatrix reads 8. Harmless in practice, but real — fixed the same way q_sh already was, via a P_ROWS bound.

New capability

4. MTP draft cache on KVarN: 832 MiB f16 → 131 MiB kvarn2. This is safe by construction rather than by measurement — the draft only proposes tokens and the target verifies them with its own cache, so draft-cache error can lower acceptance but cannot corrupt output. Measured acceptance actually went up (0.875 vs 0.769). Total MTP cost 2522 → 1814 MiB.

5. Projector weights in pinned host memory, opt-in via --mmproj-host-weights: frees 886 MiB of VRAM, image encode +5–16%, output byte-identical, LLM speed untouched. Off by default.

One analysis that may interest ggml-org

The devices[].integrated = false gate (their PR ggml-org#16308) is masking a scheduler race, not a hardware limitation. llama.cpp deliberately places model input tensors in the CUDA host buffer; while CUDA reports it as unsupported, the scheduler copies those inputs to VRAM and calls ggml_backend_synchronize first — and that copy is the only barrier between the CPU writing inputs and kernels reading them. Naively enabling host buffers removes the barrier along with the copy. The fix is a separate weights-only buffer type. Isolated by enabling the gate without moving any weights: encode time unchanged, output corrupted with the identical hash. This appears to explain the open ggml-org#25992 (server returning other sessions' answers verbatim at -np 4).

On quality: none of these changes touch the target cache numerics. Verification was byte-exact output comparison at temperature 0, single slot — stricter than KLD for this purpose. Note that at --parallel > 1 response hashes are not reproducible run-to-run on this build (batching order), so multi-slot hash comparison is not a valid method.

Not resolved — please treat as known

• One anomaly is reproducible but not root-caused: with a kvarn2 draft cache at --parallel 1, depth ≈2123 and one specific prompt, the first request after start differs from all subsequent ones (the draft itself diverges). Does not reproduce at --parallel 4 or at depth 20115. GGML_KVARN_SPLIT_MAX_Q=1 masks it. compute-sanitizer --tool initcheck is clean, though the pool uses VMM so that is not conclusive.
• Unexplained: a coarser separate draft (Q4_K_S sidecar) reaches higher acceptance at 45k depth than the built-in layer — 0.804 vs 0.627. Measured stably over three runs, not understood.

[Anbeeld]

I'll take a look.

Does the mmproj change result in weights being in RAM but the model still being processed on GPU?

[valujin]

Does the mmproj change result in weights being in RAM but the model still being processed on GPU?

Yeah, you got it right. mmproj is not that big and weights used once per request - PCIe link is sufficient for that. In my testing I got several percent speed reduction compared to full on GPU mmproj well known classic.

[Anbeeld]

There's the --no-mmproj-offload option that offloads both weights and exectution to CPU, which is safe fallback, but GPU execution is indeed better if there won't be any issues.

[valujin]

I know and I hate it:)

1. Qwen mmproj require at least 1024 image tokens for reasonable output - this is 6.3s per image on my Ryzen 9950x. 2048 token images take 23s per image. That's insane.
2. If flash attention is on and --no-mmproj-offload is set - time per image just doubles (13s per 1024 token image). Using ZenDNN help, since it's used outside flash attention path.

My next thought was - why no use system RAM for on GPU CUDA computation. There are two ways:

• GGML_CUDA_ENABLE_UNIFIED_MEMORY - it's already implemented in original llama.cpp, correct, but only frees 456 MiB, and 328 under load. Not worth it.
• ggml_backend_cuda_host_buffer_type — it's already in ggml, pinned host memory that CUDA kernels read over PCIe. But it's gated behind devices[].integrated, which is hardcoded to false with a comment blaming corrupted output (Eval bug: Broken/no Gemma 3n output on CUDA (Nvidia Jetson Orin Nano) ggml-org/llama.cpp#15034).

I flipped the gate and got the corruption too. Turns out it has nothing to do with weights or hardware. llama.cpp deliberately puts the model's input tensors in that same host buffer, and while CUDA reports the buft as unsupported, the scheduler is forced to copy them to VRAM with a ggml_backend_synchronize first — that sync is the only barrier before the kernels read them. Open the gate and the barrier disappears together with the copy, so the next ubatch's set_input() overwrites memory the previous one is still reading.

So instead of opening the general path I added a separate weights-only buffer type, reachable only via ggml_backend_reg_get_proc_address — model inputs can't land there, barrier stays.

Two gotchas: plain zero-copy is 2.9x slower because cuBLAS re-reads each weight tile over PCIe, so the weight gets staged into a device buffer with one cudaMemcpyAsync first (then it's a flat +56ms, one pass of 888 MiB). And clip.cpp sets threads only on backend_cpu, so a projector on any other backend silently runs on 4 threads.

886 MiB freed, encode +5-16%, output byte-identical, LLM speed untouched - this is a beast for me) Already running OCR tests for the past 6hrs with --spec-draft-type-k kvarn2 --spec-draft-type-v kvarn2 --mmproj-host-weights on my production dataset and everything is 100% correct.