[JartX]

Overview

Adds ROCM_OCTAVE, an attention backend that keeps the KV cache of
full-attention layers at 3-4 bits per value, with native HIP kernels for
RDNA3 (gfx1100) and MI300X (gfx942). Select it with
--kv-cache-dtype octave_k3v3, octave_k3v4 or octave_k3v3_compact
(216, 248 and 196 bytes per slot at head size 256, against 1024 for fp16).

RFC: #59773. docs/design/rocm_octave.md describes the format and the
per-architecture contract.

Claims

• Capacity:
  • On 4× RX 7900 XTX with Qwen3.8-27B, 4.6× the KV tokens of fp16.
  • On one MI300X, 1.95× the tokens of AITER's fp8 cache (7.0M vs 3.6M, i.e.
    53 vs 27 concurrent 128k requests).
• Quality: GSM8K and a 126k needle-in-a-haystack test match fp16 and fp8
  on both architectures. Next-token NLL at 380k is +0.0009 to +0.0015 nats
  over fp16.
• Speed:
  • On MI300X, decode for one request matches AITER fp8. Per layer it is
    faster from 32k on (65 vs 103 µs at 128k).
  • On RDNA3 at 380k, per layer, decode is 4.6× faster than TurboQuant and
    MTP verification 11× faster.
• Works with MTP and CUDA graphs. No host syncs in the metadata builder.
  We have served it in production on RDNA3 since October 2, 2026.
• Guarded to the validated architectures, gfx1100 and gfx942:
  • CMake builds the sources and registers the ops only for them.
  • supports_compute_capability rejects other devices.
  • The tests skip elsewhere.

Validation

Models. Qwen3.8-27B W4A16 (compressed-tensors, int4, group 32). To
reproduce, cyankiwi/Qwen3.8-27B-AWQ-INT4
is a public checkpoint with the same scheme and architecture; KV token
counts will differ slightly.

Tests (pytest tests/kernels/attention/test_rocm_octave.py), on this
branch at 28aa54e (merged with main 3897f95):

	gfx1100	gfx942
test_rocm_octave.py	105 passed	87 passed, 18 skipped (sparse decode is gfx1100-only)
Octave kernels with scratch in the built .so	0 of 108	0 of 108

• The tests compare store, decode and prefill with the PyTorch reference for
  all three cache dtypes. That covers:
  • the native and the portable decode;
  • 1 and 4 query tokens (MTP);
  • padded queries and 1-256 KV splits;
  • fp16 and bf16;
  • prefill over a cached prefix with several requests;
  • the sparse (QSA) decode.
• To check that the tests can fail, I broke the kernels on purpose several
  times: a wrong codebook, a least-squares scale, a wrong log2(e), a missing
  query debias. Each time the relevant tests failed.
• VLLM_GPU_SYNC_CHECK=error on gfx942: served cyankiwi/Qwen3.8-27B-AWQ-INT4
  with octave_k3v3 and MTP k=3 through 40 GSM8K questions with no sync
  error (97.5%).
• pre-commit passes on the changed files, except update-dockerfile-graph,
  which needs /bin/bash and cannot run on my machine. mypy (CI stage)
  passes.
• CI: the AMD kernels/attention step runs test_rocm_octave.py on its
  MI300 runners (gfx942). gfx1100 is not in CI; it was validated locally as
  above.

MI300X end to end. Setup:

• 1× MI300X, TP1, fp16, MTP k=3, --max-model-len 131072; only the cache
  dtype changes between rows.
• GSM8K: tests/evals/gsm8k/gsm8k_eval.py, chat, thinking off, 200
  questions, temperature 0.6.
• NIAH: details below.
• Speed: vllm bench serve.
• The fp8 row is forced with --attention-backend ROCM_AITER_FA. With only
  VLLM_ROCM_USE_AITER=1, vLLM picks ROCM_ATTN.

KV cache	KV tokens	GSM8K	NIAH	decode at 126k	TPOT 1 user, 8k / 32k / 100k	TTFT 100k	16 users, 8k
octave_k3v3	7,014,973	98.0	28/28	52 tok/s	15 / 18 / 16 ms	20.2 s	221 tok/s
octave_k3v4	6,391,547	97.5	28/28	52 tok/s	15 / 15 / 19 ms	20.2 s	226 tok/s
octave_k3v3_compact	7,446,005	97.0	28/28	50 tok/s	16 / 18 / 16 ms	20.1 s	211 tok/s
fp8, ROCM_AITER_FA	3,597,048	97.0	28/28	50 tok/s	15 / 17 / 16 ms	18.6 s	235 tok/s
fp16, ROCM_ATTN	1,952,899	97.0	28/28	6 tok/s	22 / 45 / 108 ms	23.1 s	120 tok/s

AITER is ahead with 16 users and in time to first token at 100k, which points
at our prefill over a cached prefix; that is the next thing to work on.

MI300X: NIAH in detail

The haystack is vLLM's own Python source, cut to 126k tokens. 16 keys are
inserted every ~7.7k tokens, each a line such as
# PROJECT KEY Aldebaran: WK3DEGZD, with star names as look-alike labels and
random 8-character values. The same document is cut at 32k, 64k and 126k,
which hold 4, 8 and 16 keys. For each key present the model is asked, thinking
off, for only that key's value; an answer counts if it contains the exact
value. Keys sit between 3% and 95% of the context.

KV cache	32k (4 keys)	64k (8 keys)	126k (16 keys)	another key's value returned
octave_k3v3	4/4	8/8	16/16	0
octave_k3v4	4/4	8/8	16/16	0
octave_k3v3_compact	4/4	8/8	16/16	0
fp8, ROCM_AITER_FA	4/4	8/8	16/16	0
fp16, ROCM_ATTN	4/4	8/8	16/16	0

The model saturates this task at 128k. So it shows that the 3-bit cache loses
no retrieval there; it does not rank the formats. The harder test is the 800k
key curve in the RFC, where the keys start to blur only past 640k, because of
the model and not the cache.

MI300X: vllm bench serve in detail

Setup:

• Random prompts (--dataset-name random), 512 output tokens with
  --ignore-eos, --seed 0, no prefix reuse.
• One user at a time: 6 prompts at 8k, 4 at 32k and 3 at 100k.
• 16 concurrent users: 64 prompts at 8k.

No request failed. Each cell is median TTFT / median TPOT / output tok/s.
Random tokens give MTP a low acceptance, the same for every row.

KV cache	8k, 1 user	32k, 1 user	100k, 1 user	8k, 16 users
octave_k3v3	1.55 s / 15.3 ms / 51.3	5.99 s / 18.1 ms / 33.4	20.2 s / 15.6 ms / 17.9	3.34 s / 62.1 ms / 221
octave_k3v4	1.55 s / 15.4 ms / 52.4	6.06 s / 15.3 ms / 35.8	20.2 s / 19.2 ms / 17.1	3.95 s / 57.0 ms / 226
octave_k3v3_compact	1.57 s / 16.2 ms / 50.2	6.73 s / 17.6 ms / 32.4	20.1 s / 15.6 ms / 17.9	3.27 s / 57.2 ms / 211
fp8, ROCM_AITER_FA	1.57 s / 15.2 ms / 49.3	5.53 s / 17.1 ms / 34.1	18.6 s / 16.2 ms / 19.0	3.29 s / 49.9 ms / 235
fp16, ROCM_ATTN	1.64 s / 22.5 ms / 36.8	5.78 s / 44.6 ms / 16.9	23.1 s / 108.4 ms / 6.2	7.26 s / 112.4 ms / 120

Without MTP, to compare with UltraQuant (#57057), which needs bf16 here:

KV cache	KV tokens	8k, 1 user	32k, 1 user	100k, 1 user	8k, 16 users
octave_k3v3 (fp16)	9,437,184	1.57 s / 51.9 ms / 18.2	5.28 s / 52.2 ms / 16.0	19.1 s / 52.6 ms / 11.1	10.9 s / 82.0 ms / 157
ultraquant_4bit (bf16)	7,730,626	1.84 s / 58.5 ms / 16.1	6.24 s / 66.0 ms / 12.8	20.3 s / 87.5 ms / 7.9	14.0 s / 94.5 ms / 135

On gfx942 UltraQuant's FlyDSL decode is not used, because it targets gfx950
and does not cover 6 query heads per KV head, so this is its Triton fallback.
With MTP k=3 its output was not usable on this setup (GSM8K 1%, NIAH 0/28);
without MTP it is correct (GSM8K 98.0, NIAH 28/28). This says nothing about
UltraQuant on gfx950.

MI300X: attention per layer

benchmarks/attention_benchmarks, the whole model on one GPU: 24 query heads, 4 KV heads.

	octave_k3v3	fp8, AITER
decode, 1 token at 8k	25 µs	16 µs
decode, 1 token at 32k	32 µs	33 µs
decode, 1 token at 128k	65 µs	103 µs
decode, 16 requests at 32k	227 µs	323 µs
decode, 64 requests at 8k	220 µs	352 µs
MTP verify, 4 tokens at 32k	61 µs	—

Prefill, against ROCM_AITER_FA with an fp16 cache:

• 8k with no prefix: 2.76 vs 2.50 ms.
• Two 2k chunks over an 8k prefix: 3.28 vs 6.76 ms.
• A 2k chunk over a 32k prefix: 5.44 vs 7.15 ms.

RDNA3 (4× RX 7900 XTX, TP4) evaluation

Setup:

• 8 GB of KV cache per GPU, same image and flags for every row.
• GSM8K: chat, thinking off, 500 questions, temperature 0.6, no MTP so that
  TurboQuant is correct.
• NLL: one 380k-token document.

KV cache	Bytes	GSM8K	NLL over fp16
fp16	1024	91.8	—
turboquant_k3v4_nc	230	89.8	+0.00155
octave_k3v4	248	91.4	+0.00086
octave_k3v3	216	92.4	+0.00125
octave_k3v3_compact	196	91.8	+0.00150

Attention per layer at 380k (one TP4 shard):

	fp16	turboquant_k3v4_nc	octave_k3v3_compact	octave_k3v3
decode, 1 token	2370 µs	2565 µs	563 µs	614 µs
MTP verify, 4 tokens	48.8 ms	6.65 ms	0.58 ms	0.62 ms
prefill, 2048-token chunk	125.6 ms	86.6 ms	60.8 ms	63.9 ms

Time to first token at 380k, end to end with MTP and a cold prefix:

• octave_k3v3_compact: 297 s;
• fp16: 371 s;
• TurboQuant: 621 s.

At 100k all three are within 8%, because the GEMMs and the all-reduce
dominate.

When I replaced the metadata builder the backend first inherited with its
own, I checked that nothing moved (TP4, MTP, GSM8K 500 at concurrency 6):

• GSM8K: 92.2 → 91.8;
• MTP acceptance: 83.33% → 83.39%;
• output: 216.1 → 216.8 tok/s.

Sparse attention (qwen4_exp QSA) evaluation

Setup:

• 4x RX 7900 XTX, Qwen3.8-Flash-Next W4A16 (asymmetric, g32), TP4, fp16,
  --max-num-batched-tokens 512, no prefix caching.
• NLL from prompt_logprobs over 12 documents of ~8k tokens, split at
  position 2048, and over 24 short documents.
• Values in nats over fp16; mean NLL repeats to about ±0.003 nats.

KV cache	KV tokens	NLL pos >= 2048	NLL pos < 2048	NLL short
fp16	47,409	1.0856	1.6707	1.1788
octave_k3v3, no window	155,930	+0.0395	+0.0851	+0.0708
octave_k3v3	155,930	+0.0163	-0.0221	+0.0085

That run used our fork for the fp16 port of QSA and two RDNA3 MoE fixes,
which go in separate PRs. On this branch alone the QSA path is bf16, like
upstream.

Details

What changes. Most of the PR is new files:

• vllm/v1/attention/ops/rocm_octave.py: the format and its PyTorch
  reference, which defines the byte layout.
• vllm/v1/attention/backends/rocm_octave_attn.py: the backend and its
  metadata builder.
• csrc/attention/octave_*, the kernels:
  • store, reduce, rotation and a portable decode, shared by both
    architectures;
  • WMMA decode and prefill for gfx1100;
  • MFMA decode and prefill for gfx942;
  • the op registration in octave_torch_bindings.cpp.
• Tests, docs, and --model plus the new cache dtypes in
  benchmarks/attention_benchmarks.

Changes to existing files are small:

• the three cache dtypes in vllm/config/cache.py,
  vllm/utils/torch_utils.py and vllm/v1/kv_cache_interface.py;
• the backend in the registry and in the ROCm platform's selection;
• the sources in CMakeLists.txt;
• customize_spec now runs under the current vLLM config, because the
  backend needs the model's rotary dims to size the slot;
• in qwen4_exp (QSA), Octave as the main cache through a sparse decode,
  with the new fused QSA prepare turned off for an Octave cache, since it
  writes raw BF16 K/V.

Nothing changes for other backends or cache dtypes.

Design choices a reviewer may want to know about:

• Native kernels only. There is no Triton fallback; an architecture is
  added once it has its own kernels and passes the tests and an end-to-end
  run. The RFC lists what each one needs.
• The metadata builder is the backend's own. It builds the per-query
  maps on the device without host syncs, supports UNIFORM_BATCH CUDA
  graphs, and refreshes the maps between fused MTP draft steps.
• On gfx942:
  • QK runs on v_mfma_i32_16x16x32_i8, with the query split into high and
    low int8 halves.
  • V is widened to fp16 by one v_perm: each codebook byte becomes an exact
    fp16 denormal, which the MI300 MFMA keeps.
  • The decode file is built with -amdgpu-mfma-vgpr-form, so the
    accumulators do not halve the occupancy.
• Supported models. Head size 256 with 64 RoPE dims, i.e. Qwen3-Next and
  Qwen3.5 / 3.6 / 3.8.
  • The decode handles up to 8 query heads per KV head on each GPU.
  • The sparse decode and the exact sink/recent window are QSA-only and
    gfx1100-only for now.
  • Head size 128 is the next step.
• Size. This is a large PR: about 1.4k lines of gfx942 kernels on top of
  the RDNA3 ones. Each architecture lives in its own file behind the same
  octave_prefill / octave_decode entry points, so it can be reviewed and
  later removed on its own.

Not a duplicate. I searched open PRs and issues for octave,
turboquant rocm, kv cache 3-bit, kv cache quantization hip and
lloyd-max kv. The closest are:

• UltraQuant ([Attention] Add UltraQuant 4-bit KV cache backend (FlyDSL D=256) #57057): FP4 with FlyDSL for gfx950, built on TurboQuant.
• KVarN ([RFC]: KVarN: a calibration-free, variance-normalized sub-8-bit KV-cache quantization backend #46613 / [Attention] Add KVarN: calibration-free variance-normalized KV-cache quantization backend #46812): Triton, CUDA-first.
• [Kernel][ROCm] CDNA HIP attention kernels for INT8/INT4 per-token-head KV cache (MI300) #43835: HIP int8/int4 for CDNA.
• Feat/tq rocm decode v2 #40396: a Triton TurboQuant decode for MI300X.
• [ROCm] TurboQuant KV cache support for Gemma 4 #53231: TurboQuant for Gemma 4 on ROCm.
• [ROCM][Perf] Add ROCM_SEGMENTED_ATTN attention-backend works across RDNA3/3.5/4 #59132: Triton attention for RDNA without KV quantization.

None targets RDNA with native kernels, and none runs below 8 bits with MTP on
ROCm.

AI assistance. I used AI assistance to write this PR: Claude, and
Qwen3.8-27B running locally on the boxes it was measured on. I have reviewed
every changed line, ran the tests and evaluations above, and can answer for
the change.

---

Pull Request Checklist

• I used vLLM's /pr-checklist skill. (Mandatory for agents, optional for humans).

• AI assistance was used during the creation of this PR.

• Design Fit: Minimizes impact on core components, reuses existing functionality, and justifies added complexity.

• Testing and Validation: Validates the change and ensures any added tests are meaningful and reliable, with CI coverage or documented CI resource constraints and validation performed outside CI.

• Code Quality and Style: Keeps code and comments clear and concise, and updates relevant documentation and examples.

• Pull Request Contents: Includes a brief summary and relevant links, supports claims with evidence, explains root causes and implementation trade-offs, and follows the contributing guide.

[mergify]

Documentation preview: https://vllm--59774.org.readthedocs.build/en/59774/

[mergify]

This pull request has merge conflicts that must be resolved before it can be
merged. Please rebase the PR, @JartX.

https://docs.github.com/en/pull-requests/collaborating-with-pull-requests/working-with-forks/syncing-a-fork

[mergify]

This pull request has merge conflicts that must be resolved before it can be
merged. Please rebase the PR, @JartX.

https://docs.github.com/en/pull-requests/collaborating-with-pull-requests/working-with-forks/syncing-a-fork

[vidriagon]

Qwen3.8-27B on Dual RX 7900 XTX: Benchmarking KV Cache Quantization in vLLM

I wanted to see how well Qwen3.8-27B would run on two AMD Radeon RX 7900 XTX GPUs (24 GB each), especially during long conversations. The goal was straightforward: reduce GPU memory usage by quantizing the KV cache, while keeping generation speed and answer quality as close as possible to the unquantized setup.

I compared five KV-cache configurations in vLLM, all using float16 and tensor parallelism across both GPUs. Three packing variants—k3v3, k3v4, and k3v3_compact—use our RDNA3 backend, Octave (vLLM PR #59774). int8_per_token_head uses a separate quantized kernel implementation, while auto leaves the KV cache unquantized. The model was cyankiwi/Qwen3.8-27B-AWQ-INT4.

Main results

KV-cache mode	Generation speed at ~180K context tokens	KV-cache capacity (tokens)	GSM8K accuracy
k3v3_compact	44.8 tokens/s	1,385,454	75%
k3v3	44.4 tokens/s	1,143,333	74%
k3v4	43.8 tokens/s	1,005,882	75%
int8_per_token_head	31.8 tokens/s	505,185	72%
auto (unquantized)	Test not completed	252,671	73%

The three Octave variants sustained roughly 44–45 tokens/s even with ~180K tokens of context. k3v3_compact also provided the highest KV-cache capacity. The separate INT8 implementation slowed down more at long context lengths, while unquantized auto performed substantially worse.

The capacity figures above describe the aggregate KV-cache token capacity reported by vLLM, not the maximum length of a single conversation. For these tests, --max-model-len was set to 200,000 tokens.

One important correction came out of the testing. Our first measurements made long-context generation look extremely slow, but the issue was in our benchmark script rather than the quantized cache. We estimated how long prompt processing (prefill) would take and accidentally counted much of that time as generation. We fixed the measurement by using streaming and timing the interval between the first and last received output tokens.

That revealed the real bottleneck: processing a very long prompt before generation begins. At around 180K context tokens, the quantized modes took roughly two minutes to produce the first token. Once generation started, however, they still ran at around 44 tokens/s. This distinction matters when judging how responsive the model feels in everyday use.

Did quantization affect quality?

We did not observe a clear quality difference in these tests. All five configurations answered our 10 reasoning questions correctly and passed the needle-in-a-haystack retrieval checks at the tested context lengths. We then used vLLM's GSM8K evaluation (gsm8k_eval.py --num-shots 5 --max-tokens 256, same command for all five variants) and obtained accuracy scores between 72% and 75%. That sample is too small to establish a statistically significant difference between the modes, but the results suggest that KV-cache quantization did not cause a noticeable loss of accuracy in these tests.

We also measured server throughput using the official vllm bench serve tool (vllm bench serve --dataset-name random --random-input-len 2000 --random-output-len 500 --num-prompts 60 --max-concurrency 4, same command for all five variants, just pointed at each server):

KV-cache mode	Aggregate generation throughput
k3v3	100.4 tokens/s
k3v4	99.9 tokens/s
k3v3_compact	99.3 tokens/s
int8_per_token_head	99.6 tokens/s
auto	68.1 tokens/s

Which one would I use?

For this setup, I would choose k3v3_compact. It offers the largest KV-cache capacity, maintains strong generation speed with long contexts, and showed no noticeable accuracy loss in our tests. k3v3 is a very close alternative; based on these results, it is difficult to justify using unquantized auto.

These findings apply to two RX 7900 XTX GPUs and this particular vLLM build. Other hardware or software configurations may produce a different ranking.

Tests carried out October 10–13, 2026. Backend: Octave, vLLM PR #59774.

[JartX]

Thanks a lot for taking the time to run this on 2× 7900 XTX with a third-party AWQ checkpoint. An independent setup like this is exactly the kind of data that's hard to get on our own boxes.

Two small suggestions, if you have time for another pass:

1. GSM8K. gsm8k_eval.py with 5 shots and --max-tokens 256 tends to underreport chat models, because a lot of answers get cut before the final number. Running it through the chat endpoint with thinking disabled and ~2048 max tokens usually gives a much more representative score (this model family is in the mid/high 90s that way). Your relative comparison between the five modes still holds; this would just make the absolute numbers comparable with other reports.
2. auto at long context. If it fits, a run of the unquantized auto at the longest context it can hold (even shorter than 180K) would give a same-length baseline for the 44 tok/s figure.

On the ~2 min TTFT at 180K: that's mostly prefill and expected at that length on two cards. If you happen to have the auto TTFT at a shorter matching length, it would also help confirm Octave adds no prefill overhead.

[tjtanaa]

@JartX @vidriagon Given that we can reach a longer context length now, can you please evaluate on long context dataset e.g. NIAH ?

[JartX]

@tjtanaa Of course
Tell me what tests or evaluations you need. I have already run full-scale development sessions with a 512K context window.

Whatever you ask for, I will do.

[vidriagon]

Thanks for the detailed feedback, really helpful. Reran both points:

GSM8K via chat endpoint. You were right — with --use-chat-completions, enable_thinking: false, and --max-tokens 2048 instead of the completions endpoint at 256, every mode jumped into the mid/high 90s, confirming answers were getting cut before the final number in the original run:

KV-cache mode	GSM8K accuracy (chat, thinking off, 2048 tokens)
k3v3_compact	98.0%
int8_per_token_head	97.5%
auto	97.5%
k3v4	97.0%
k3v3	97.0%

Relative ranking is basically unchanged (all within noise), but now the absolute numbers should be comparable to other reports. Thanks for catching that.

auto at a matching shorter context. Ran the same streaming methodology at ~8K input tokens, same length for both:

	TTFT	Decode
auto	6.3s	14.1 tok/s
k3v3_compact	5.9s	53.2 tok/s

TTFT is basically identical, so that confirms Octave isn't adding any prefill overhead — the ~2 min figure at 180K is just prefill scaling with context, same as you'd expect. But the decode gap is already there at 8K, not something that only shows up at long context — auto is ~3.8x slower than k3v3_compact even at this length. So the 44 tok/s figure at 180K isn't really a "long-context" effect on auto, it's the same gap you'd see anywhere, just compounding with the usual bandwidth cost of raw fp16 KV cache. @tjtanaa

[tjtanaa]

@JartX Given that this is only validated on gfx1100, please make sure that the tests are gated to only run on this hardware. We will relax the condition once validation has been done on other platform.

[tjtanaa]

we should also override this supports_compute_capability to only validated architecture.

[JartX]

@tjtanaa done, pushed in 884e0fb: I added get_rdna_version() next to get_cdna_version() in vllm/platforms/rocm.py, and Octave now gates on get_rdna_version() == 3 / get_cdna_version() == 3, with no on_gfx* left in the backend or the tests.

The kernels are still only built and dispatched for gfx1100 and gfx942, the chips I validated, so another RDNA3 chip passes selection but gets a clear "not validated" error from the kernel dispatch instead of running untested code. Happy to open it up to the whole gfx11 family if you prefer.

Tested on gfx1100: test_rocm_octave.py 105 passed.

[tjtanaa]

@JartX This is not tested in the unit test. Please add a unit test for this.

[JartX]

Added test_kernels_follow_declared_cache_layouts in test_rocm_octave.py (5da887e): for every layout in supported_kv_cache_layouts() it builds the cache view vLLM allocates (customize_spec + dense_kv_cache_views, two KV heads so LBHNC and LBNHC really differ) and checks that store, decode (native and portable tier) and prefill match a contiguous cache bit for bit. I also checked it catches a kernel that ignores the strides (LBNHC then fails). gfx1100: 107 passed.

[JartX]

/ci run

[github-actions]

✅ Triggered Buildkite CI #97126 for commit 95dc71b2ff7a.

ℹ️ This PR is 358 commits behind upstream main. Merged upstream main into this branch as 95dc71b2ff7a before starting CI. Pull the branch before pushing more changes.

[JartX]

/ci run

[github-actions]

✅ Triggered Buildkite CI #97134 for commit 678cc6e19d26.

ℹ️ This PR is 4 commits behind upstream main. Merged upstream main into this branch as 678cc6e19d26 before starting CI. Pull the branch before pushing more changes.

[tjtanaa]

LGTM