Skip to content

Qwen3.6: opt-in int8 K/V cache and single staging buffer - #2742

Closed
ggeorgovassilis wants to merge 2 commits into
Niko1221:mainfrom
ggeorgovassilis:q36-int8-kv
Closed

ggeorgovassilis wants to merge 2 commits into
Niko1221:mainfrom
ggeorgovassilis:q36-int8-kv

Conversation

@ggeorgovassilis

Copy link
Copy Markdown

Title

Qwen3.6: opt-in int8 K/V cache and single staging buffer

Summary

Two opt-in improvements to the Qwen3.6-35B-A3B own-engine path, both off by default so the default path stays byte-identical to the last release:

  1. int8 (q8_0) K/V cache behind --kv int8 — roughly 4x smaller K/V cache (~44 KiB/token fp32 -> ~11 KiB/token int8), so an 8 GB card doubles its context (16K -> 32K).
  2. Single staging buffer behind STRATA_Q36_SINGLE_STAGE=1 — halves the prompt-streaming staging VRAM (one buffer instead of two double-buffered).

What changed

  • CMakeLists.txt — build wiring for the new q8_0 K/V kernels.
  • include/strata/kernels/q36_kv_q8.hpp — new int8 K/V kernel interface.
  • src/kernels/cuda/q36_kv_q8.cu — CUDA q8_0 K/V cache kernel.
  • src/kernels/q36_kv_q8_parity.cpp — CPU reference kernel for parity testing.
  • src/qwen35/q36_engine.cu — the --kv int8 engine path; STRATA_Q36_SINGLE_STAGE=1 (single staging buffer), numeric STRATA_Q36_NO_MMQ, STRATA_Q36_PLACE_DEBUG, and an attention-scores-buffer shrink (256 -> 64 MB) so staging placement still succeeds on tight VRAM.
  • docs/Q36_INT8_KV.md — design note with the measured numbers.
  • docs/DETAILS.md — links the new section into the table of contents.

Extra Notes

  • --kv int8 defaults to --kv fp32, which is byte-identical to the previous release.
  • Benchmarks were measured on an RTX 4060 (8 GB); see docs/Q36_INT8_KV.md.

ggeorgovassilis and others added 2 commits October 15, 2026 21:32
Quantize the resident K/V cache of the Qwen3.6 path (src/qwen35/q36_engine.cu)
from fp32 to q8_0 (32 int8 codes + one fp16 scale per block) when --kv int8 is
passed; the default --kv fp32 path is byte-identical. Adds the quantize/
dequantize primitive (include/strata/kernels/q36_kv_q8.hpp,
src/kernels/cuda/q36_kv_q8.cu) and a GPU parity test
(src/kernels/q36_kv_q8_parity.cpp). Implements docs/Q36_INT8_KV.md.

Measured on an RTX 4060 Max-Q Laptop 8GB / Ryzen 7 7735HS / 62GB with
Qwen3.6-35B-A3B-UD-Q4_K_XL.gguf (see docs/Q36_INT8_KV.md):
- greedy decode is bit-for-bit identical to fp32 on short and long prompts;
  the 248,320-logit dump on the long prompt differs by at most 0.79.
- the context window doubles: fp32 16,384 vs int8 32,768 (the 131,072 target
  is not reachable; the ceiling is dense weights + expert-streaming staging +
  the attention-score buffer, not K/V).
- decode is ~4% slower at 4K (16.0 -> 15.4 tok/s); prefill unchanged.

CUDA verified; the HIP build of the shared source is still to be run.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
- STRATA_Q36_SINGLE_STAGE=1 halves the streamed-prompt staging VRAM (one
  buffer instead of two double-buffered)
- STRATA_Q36_NO_MMQ=1 is now numeric (MMQ stays on by default)
- STRATA_Q36_PLACE_DEBUG logs the staging allocation
- shrink the attention-scores buffer (256 -> 64 MB) when the staging
  buffers only just miss, so placement still succeeds

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
@ggeorgovassilis

Copy link
Copy Markdown
Author

Thanks — but this overlaps almost entirely with an older PR, #2698 (TorBere, opened the evening before this one): the same --kv int8 K/V cache for the Qwen3.6-35B-A3B own-engine path, with --kv fp32 byte-identical by default. Neither PR mentioned the other, so I checked how far the overlap goes.

Both heads are based on 296b59d (current main) and are individually MERGEABLE/CLEAN. Merged together locally (git merge-tree --write-tree 3a9e6c8 d85b465):

Auto-merging src/qwen35/q36_engine.cu
CONFLICT (content): Merge conflict in src/qwen35/q36_engine.cu

10 conflict hunks, all in src/qwen35/q36_engine.cu: the Engine fields, the init() allocations for the attention / GDN / MTP layers, the k_attn_prep / k_kv_f16 / k_attn_part launch sites, and the --kv parser and usage text. Essentially every place either PR touches the K/V cache, the other touches too.

The overlap is semantic, not only textual — the two use different formats and different strategies:

#2698 this PR
group kKVG = 64 Q36_KV_Q8_GROUP = 32 (llama.cpp q8_0)
scale __half* uint16_t* holding fp16 bits
kernels template <bool KI8, bool VI8> on the existing kernels, one kv4() dispatcher separate k_attn_prep_q8 / k_attn_part_q8 / k_kv_f16_q8
stores independent: k8v32 / k32v8, plus a STRATA_Q36_KV env fallback both stores together: int8 / fp32 only
decode vectorized uint2 K reads (8 codes/load), sc[k] * group scale folded into a shared V table per-element dequantize in the inner loop
test / docs none q36_kv_q8_parity, docs/Q36_INT8_KV.md, CMake wiring

The different group size means a different cache layout, so at most one of the two can land as-is.

What in this PR does not overlap, and is independent of the K/V format: STRATA_Q36_SINGLE_STAGE=1, the numeric STRATA_Q36_NO_MMQ parsing, STRATA_Q36_PLACE_DEBUG, and the attention-scores buffer shrink (256 -> 64 MB) in place_experts().

I'm happy to close this one and let #2698 carry the K/V cache, re-opening the staging/placement work as a PR of its own — or to keep this one and let #2698 close instead. @TorBere, which do you prefer?

Whichever lands should probably keep the parity test and the design note from here: #2698 has neither, and its measurement is on a different card (RTX 3060 12 GB) from the 8 GB ceiling this one was written around.

@Niko1221

Copy link
Copy Markdown
Owner

Thanks for the pull request.

We're resetting the tracker for 0.2.0-beta. There are more than 1,300 open issues and pull requests, too many for us to review properly, so we're closing everything that isn't being worked on right now. Closed here does not mean rejected.

To reopen it takes about two minutes:

  1. Rebase on main if it no longer applies, and re-run your checks on 0.2.0-beta.
  2. Start the title with its area: [Engine], [GUI], [TUI], [Benchmark], [Docs] or [Other].
  3. Fill in the pull request template: the problem and result, and the validation (what you ran, on which revision, what you saw).
  4. Comment /reopen. A check runs and either reopens the pull request or lists what is still missing.

The rules are in CONTRIBUTING.md.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants