Repository navigation
Conversation
|
Thanks for splitting it out and for updating Q8_KV.md instead of deleting it. A few things before I can take it:
I'll test the rebased version on an RTX 4070 laptop and an RTX 5090. |
…server (simplexai-labs#33) Builds its long texts from this repository's own files at run time, hides a code word at the chosen depths and asks through the API. Checked on the Coder: 1K, 32K (10/50/90%) and 128K found.
|
Hi @VecSzn — answering all five points; the branch is rebased and pushed, so the diff you see now 1. HIP (your either/or: new kernels build and pass on AMD, or keep the old readers there). The answer is "the new kernels build on HIP, minus the MMA tier". Concretely:
I did not take the "keep the old readers on HIP" option: with the gate the old readers buy 2. Measurable benefit — you asked for 32K/64K decode, int8 readers vs the q8_0 readers. RTX 4070 Laptop 8 GiB, Ornith 1.5 35B-A3B Q6_K,
Paired ratios (same round): The Incidental finding, in favour of replacing them: the old 3. The 15/15 table was the old readers' — agreed, and re-measured. It is now labelled as v0.1.43-poly.5's elementwise q8_0 readers in
The int8 readers quantize the query side, so scores move in the last bits and greedy can Caveats I won't paper over: Tiel and Occamy are no longer on this machine, so those rows have 4. The 3603/3603 dump — your suspicion was half right, and here is the precise claim.
It still means something, and stronger than I argued: And your query-quantization point does not apply to that A/B: it was arm A = local repo, Composition: 24 positions × 30 layers × 5 stages = 3600, plus 3 5. The stale "Gemma 4 is out of v1" line — it was my error, and it is fixed. I believed the rebase had removed it; it had not. It is gone now, and the AMD section's Rebased onto current main (PR 28/29 included): |
|
Thanks for the answers and the rebase. I tested 304c7b1 on an RTX 5090 (Linux, CUDA 12.8) with Qwen3.6-35B-A3B UD-Q4_K_M against main (26e4cb5):
With default settings (expert swaps, lookup guesses), 96K, 400 tokens, two runs each: 113.8 and 114.0 tok/s on main, 100.9 and 100.9 here. So on this card the new readers write 9-12% slower and read 1-4% faster. Against fp16, main's q8_0 kept the first 128 of 128 tokens at 32K and 45 at 96K; this PR 58 and 123.
The illegal memory access with main's q8_0 past 4,096 tokens: I couldn't reproduce it on Linux. With the 5090 held to 6.9 and 4.9 GiB free by another process, the default reserve, Smaller things:
On the speed: your laptop shows the new readers ahead at 64K, the 5090 shows them behind at 32K and 96K. Before they replace the current format for qwen35 they need to be at least level on a desktop card too. Gemma 4 and |
|
I ran 304c7b1 on an RX 9070 XT (16 GB, gfx1201; Fedora 44, ROCm 7 from AMD's Python wheels; Ryzen 7 9800X3D). The model is Tiel-Coder-35B-A3B-MTP UD-Q4_K_XL, which is qwen35moe like Qwen3.6, so the same attention code and readers run. It isn't your UD-Q4_K_M, so the numbers won't line up exactly with yours. The baseline is the PR's own base, fd2cb1b, not 26e4cb5. Both builds carry the same local patch, a ROCm VMM workaround that doesn't touch the attention code. On HIP, the PR takes the dp4a readers and the cuBLAS batch path. Blocking on AMD:
|
fd2cb1b, fixed experts (--expert-swaps 0 --cpu-share 0 --tail-skip 0), 8K prompt, 64 tokens |
sequences 1 / 2 / 3 / 4 |
|---|---|
| fp16, run 1 and run 2 | 64 / 40 / 16 / 10 (both runs) |
| q8_0, run 1 | 64 / 41 / 64 / 13 |
| q8_0, run 2 | 15 / 56 / 64 / 64 |
So on HIP, main's batch path already differs from the single-sequence path. With q8_0 it also changes from run to run. I can open a separate issue for that if it's useful. It does mean the check can't show whether this PR matches main on AMD. The crash doesn't need the check to show up, though.
Speed with your fixed settings
--expert-swaps 0 --cpu-share 0 --window 1 --tail-skip 0 --max-context 131072 --kv-whole, a real-text prompt, 128 tokens, greedy. The prompts came out at 38,263 and 122,583 tokens.
| Prompt | fp16, read / write | fd2cb1b's q8_0, read / write | this PR's int8, read / write |
|---|---|---|---|
| 38K | 2,752 / 48.2 tok/s | 2,752 / 50.7 tok/s | 2,753 / 53.2 tok/s |
| 123K | 1,682 / 32.3 tok/s | 1,657 / 33.4 tok/s | 1,691 / 36.9 tok/s |
Experts in VRAM at load: fp16 4,559, q8_0 5,502, int8 5,084.
Against fp16, the first 128 tokens kept were:
- 38K: q8_0 9, int8 9. Both part from fp16 at the same token.
- 123K: q8_0 93, int8 31.
Speed with default settings
--max-context 196608 --kv-whole, the 122,586-token prompt, 400 tokens, two runs each.
| read | write | |
|---|---|---|
| fd2cb1b's q8_0 | 1,899 / 1,911 tok/s | 48.5 / 49.7 tok/s |
| this PR's int8 | 1,940 / 1,682 tok/s | 54.5 / 60.3 tok/s |
So on this card, int8 writes 5–10% faster than the current q8_0 with fixed settings and about 12–20% faster with default settings. Prompt reading is level to 2% faster. That's the opposite of your 5090 result on writing, and in line with smashdog's laptop.
An earlier run with drafting on (--max-context 196608 --kv-whole) gave the same picture. Writing was level at 64K (62.4 tok/s both). Reading was 2–5% faster: 1,995 against 2,101 tok/s at 64K.
The expert cache at load varied between runs on this card. int8 loaded fewer experts than q8_0 in two of three pairings, by 210 and 418, and 71 more in the third. So it may take some VRAM for the fp16 twin. I haven't pinned that down.
poly_kv_q8_test doesn't build or pass on HIP as it is
- It doesn't compile.
src/poly/kv_q8_poly_test.cppis host C++, and__ushort_as_halfisn't available to host code on HIP (lines 56 and 230:'__ushort_as_half' was not declared in this scope). Going through the file's ownhalf_value()with amemcpyinto__halffixes it. - It stops at its first check.
check(probe_mma_layout(), ...)runs unconditionally, and on HIP the MMA body is the empty stub, so it prints[mma probe] 128 of 128 wrongand exits. The engine gates that tier ontensor_cores(), so the probe andprobe_mma_speed()should be skipped the same way under HIP.
With both changed locally, it gives ok: 34 checks (your 35 minus the MMA probe). Its own timings put the dp4a decode reads at 0.95–1.05× fp16 on the Qwen3.6 shape and 0.72–1.02× on the 64-wide shape.
--kv int8k is refused for this model ("the gemma4 backend's alone for now"), as the docs say.
|
Thanks @donschen, that's a lot of useful data.
@smashdog, with these numbers the speed is a split: the 9070 XT and your laptop write faster with int8, the 5090 slower. So I won't hold the PR on speed. The batch path is what blocks it: wrong tokens from the first steps on CUDA, a GPU fault on HIP. donschen's two fixes for |
|
I'll open the issue for main's batch path on HIP. @smashdog, here are the two --- a/src/poly/kv_q8_poly_test.cpp
+++ b/src/poly/kv_q8_poly_test.cpp
@@ -53,7 +53,11 @@ uint16_t half_bits(float f) {
return b;
}
-float half_value(uint16_t b) { return __half2float(__ushort_as_half((unsigned short) b)); }
+float half_value(uint16_t b) { // through memcpy: HIP has no host __ushort_as_half
+ __half h;
+ std::memcpy(&h, &b, sizeof h);
+ return __half2float(h);
+}
/// ggml's own 34-byte blocks for these rows, taken apart into the two arrays --kv int8 keeps.
struct Reference {
@@ -227,7 +231,7 @@ void case_error(const char* name, int n_kv, int head_dim, int rows) {
float worst = 0.0f;
for (size_t b = 0; b < blocks; ++b) {
- const float d = __half2float(__ushort_as_half(delta[b]));
+ const float d = half_value(delta[b]);
float amax = 0.0f;
for (int j = 0; j < kQ8_0Values; ++j)
amax = std::fmax(amax, std::fabs(src[b * kQ8_0Values + (size_t) j]));
@@ -774,7 +778,9 @@ int main() {
}
// n_kv = 2 and head_dim = 256 are Qwen3.6's; 256/32 tells how many blocks a head has
const int n_kv = 2, head_dim = 256;
+#ifndef STRATA_USE_HIP // no int8 mma tier on HIP: attend_q8 never takes it there (tensor_cores())
check(probe_mma_layout(), "the m16n8k32 fragment layout is what I think it is");
+#endif
// FIRST, while the card still has room for anything: attend() and attend_q8() against attend_batch, which is
// the attention the engine uses for every call of kStageRows (32) rows or more - a prompt chunk and, with
// --window, a commit of accepted guesses. attend() is asked for fewer rows than that, and this is the only
@@ -792,12 +798,14 @@ int main() {
case_attend_rows_batch("rows 8, window 512 of a 1024 ring", 256, 16, 8, 1024, 8, 512, 1024);
case_attend_rows_batch("rows 1, window 256 of a 1024 ring", 256, 16, 8, 1024, 1, 256, 1024);
case_attend_rows_batch("rows 8, window 256 of a 1024 ring", 256, 16, 8, 1024, 8, 256, 1024);
+#ifndef STRATA_USE_HIP
{
double tm = 0.0, td = 0.0;
probe_mma_speed(&tm, &td);
std::printf(" one 16x8x256 product x20000 x384 warps: mma %.1f ms, dp4a %.1f ms -> mma is %.1fx\n",
tm, td, td / tm);
}
+#endif
for (int kind = 0; kind < 4; ++kind) {
char name[96];
std::snprintf(name, sizeof name, "store_q8 rows=5 kind=%d", kind);On the ROCm VMM workaround: it's for a ROCm bug I could reproduce on its own. It's not the fix for the crash we actually see, which What the workaround fixesOn this card (RX 9070 XT, ROCm 7 from AMD's wheels, Linux 7.2), the card keeps its address translations for VRAM that has been unmapped. If you
The change is 11 lines in What it doesn't fix: crashes inside HIP when the attention caches give VRAM backWithout
The kernel log shows no GPU fault, and no reset other than the ones that come with a suspend. The workaround didn't change it. The backtrace is the same before and after, and the 11 crashes above were with the flush build. PR 31 didn't change it either. It happened with So I think it's another ROCm problem in the attention caches' unmap and re-reserve, and I haven't reproduced it outside the engine yet. Until then, the safe setting on AMD is #52I can't tell from here whether firetussin's 64K start fault is related. Ours happen when VRAM is given back mid-session, not at the start. If that start runs without |
|
A note on the batch check after #55: with long prompts |
…35moe and Gemma 4, replacing the elementwise readers Replaces --kv q8_0's format (KvFmt::kQ8_0, read back by dequantizing element by element into floats and accumulating in fp32) with an int8 that is a compute format, not just a storage one: - K and V keep codes and deltas in two arrays (codes back to back, so every four are one aligned word; a delta a 32-value block, delta = amax/127, roundf — ggml's quantize_row_q8_0_ref to the bit). - Decode reads them with __dp4a / int8-MMA kernels (plain, register- blocked and tensor-core tiers by head width and key count); the queries are quantized into the same form and a lane's run stays inside one 32-value block, so a delta multiplies once a run. - A prompt batch dequantizes the caches into an fp16 twin once, incrementally (O(rows) against an O(rows^2) attention), and hands the rows to cuBLAS's attend_batch as before. - --kv int8k keeps the values in fp16 (attend_q8_mixed) — a quarter of the fp16 cache saved with less of a loss; Gemma 4's window-plus-full layers use it. - Gemma 4 gains the whole set for the first time: the window layers' ring stores write in two pieces, the twin is vmm-mapped on demand (twin_cover), and card_row_bytes() now counts the int8k value side at fp16's width (it was short by 1.44x, which overran a parked conversation's D2H blob). - qwen35's new step() batch path (attn_rows()) stores and reads the quantized rows in the same shape. - q8_0 stays as an alias of int8. The HIP refusal is kept and covers int8 and int8k (the q8 kernels are CUDA-only). - poly_kv_q8_test (35 checks) pins the store against ggml's own quantizer and compares the three attention paths; the old kv_q8_test.cpp is replaced by it. - Q8_KV.md and GLM.md's --kv row are updated in a following commit with the speed and answer checks against the previous format. The default (--kv fp16) path is bitwise what it was.
… on HIP ROCm has no mma.sync.aligned.m16n8k32.s8, and attend_q8's use_mma only looked at head_dim and the key count, so a HIP build with up to 16384 keys really did launch that PTX and fail at load. Two changes: - use_mma now asks tensor_cores() as well (false under __HIPCC__, and on CUDA by ptxVersion), and STRATA_KV_MMA cannot override it — the same convention the fp16 MMA tier already followed; - the mma_s8 PTX body is compiled to an empty stub under HIP: the kernel is never selected there, but the template still has to instantiate. CMake drops the NOT STRATA_ENABLE_HIP gate on poly_kv_q8_test: the dp4a tiers only need __dp4a (hip_compat's intrinsics.hpp maps it) and the MMA tier is gated at runtime, so the test can cover the dp4a readers on AMD. poly_kv_q8_test: 35/35 on an RTX 4070 Laptop (arch 89); strata-poly on a 32K prompt gives 51.8 tok/s at decode and output identical to before the change.
…the int8-against-fp16 answers, make the HIP wording exact The 15/15 character-for-character table was measured on the previous elementwise q8_0 readers, which did not quantize the query side; keeping it unlabelled read as a claim about this change. It now says which readers produced it. Adds the re-measured table for the int8 readers (Ornith --kv int8, gemma-4-26B-A4B --kv int8 and --kv int8k, five prompts each, greedy, --no-draft): the differences are rewordings inside the reasoning chain, and the finished answers keep their result (KV OK, 42, Paris); one Gemma 4 prompt is bit-identical in all three arms. The HIP note said the MMA tier "a HIP build does not have"; with the tensor_cores() gate that is now a runtime choice and the PTX body is an empty stub under HIP, so the note says that instead.
The file is shared with the Strata project and its first line carries a BOM; my rebase dropped it. Put the BOM back and leave the line itself untouched.
__ushort_as_half is not available to host code on HIP, so half_value() goes through memcpy into a __half, and the two case_error() call sites use it too. The MMA layout probe and probe_mma_speed() run unconditionally, but under HIP the int8 MMA body is the empty stub (attend_q8 is gated on tensor_cores()), so the probe would report 128 of 128 wrong and stop the run; both are now behind #ifndef STRATA_USE_HIP, as the rest of the tree guards. 34 checks on AMD, 35 on CUDA.
With --batch every row belongs to its own sequence, so the codes a row writes and the codes it reads are that sequence's, not the bound one's -- the fp16 path reaches them through seq_buf(r.seq, ...) and the int8 path did not, so a batched sequence attended the wrong cache from the first token (on CUDA that read as answers parting at token 1-18; on HIP it faulted). The fp16 twin is one pair for the whole model and the incremental conversion in attn() assumes a single sequence growing row after row (f16_owner_/f16_rows_). A batch interleaves sequences, which leaves those two saying nothing useful, so attn_rows() converts the row's own sequence up to its own position instead and clears the incremental state afterwards. A batch is a few rows, so the cost is O(rows in the batch x that row's position) -- not the per-token decode cost the incremental conversion exists to avoid.
…e reader swap kQ8WorkRows sized the fp16 workspace the int8 prompt path dequantizes K and V blocks into; the old q8_0 readers had their own, and the swap took the code that used it without the constant, so it no longer compiled. fattn_q8() and fdec_q8() -- the flash-attention and decode kernels the prompt and decode paths do not take today -- still name the cache they read, which is the same two planes (k_code / k_delta, v_code / v_delta), so Kvq8 now lives in fattn.hpp where they are declared.
|
Rebased and re-run, and this time with On this card (RTX 4070 Laptop, 8 GiB, CUDA 13.4, Ornith 1.5-35B-A3B UD-Q6_K, I could not reproduce the parting you saw on the 5090: I ran it with and without What I did find and fix while chasing it, in With That is a real hole in the batch path, and the fix is cheap (a batch is a few rows), but I want to be straight about what I can and cannot claim: I did not manage to make it fail before the fix, so I have a code argument, not a reproduction. If your 5090 run still parts after this rebase, send me the exact command line -- Also landed:
|
|
Thanks, that fix fits what I saw: rows reading the bound sequence's codes would part at the first tokens, and on HIP read where they shouldn't, as donschen's fault did. The 5090 run that parted (Linux, CUDA 12.8, every expert in VRAM): No That pod is gone, so I'm checking the rebased branch on the RTX 4070 laptop with an MSVC build instead, with |
…le it The int8 arm runs its prompt stage 9-12% below fp16 on an RTX 5090 while its decode matches. The prompt path is fp16 either way -- the codes go into the twin and cuBLAS takes the same halves -- so the cost sits in the conversion or in the cache holding them, and neither is measured separately yet. Written down here instead of left in the review thread, with the two measurements that would explain it.
|
The laptop run of the rebased branch (2367529; 52b7f92 only adds docs). RTX 4070 Laptop 8 GB, Windows, MSVC 14.44 + CUDA 13.0, a 5,500-token prompt of real text,
So the int8 batch path still parts from the first token, on a card where main's q8_0 and the PR's own fp16 match. Your 2,048-token prompts passing would fit a bug that needs the rows past some length (a second workspace chunk of Two build notes from the same machine:
|
Three calls still passed attend_batch() and attend() the old q8_0 arguments (lines 101, 104 and 212), so poly_fattn_test -- built alongside strata-poly -- stopped compiling with this branch. The reference each case compares against is now the halves the int8 path would hand those kernels: the codes' whole cache dequantized, K then V. Both modes pass (`all passed`, q8 and f16 arms).
|
On the batch check, I still cannot make it part here. With your 5090 command line adapted (Ornith 1.5-35B-A3B UD-Q6_K, since I do not have Qwen3.6 UD-Q3_K_M or Tiel Coder locally):
and the same 5,500 with MiniCPM-V-4.7-35B-A3B Q6_K: 64 / 64 / 64 / 64. All with Two differences from your runs that may matter more than the prompt length:
If you still see it after |
|
Thanks for the test fix. One correction on the probe: the laptop run in my last comment used real text, not the random ids (those were only the 5090 run). The prompt is the first part of a long English document followed by a question, in Qwen3.6's chat template, cut to its first 5,500 tokens. The card is the same class as yours (RTX 4070 Laptop, 8 GB). So the difference left is the model: Qwen3.6 UD-Q3_K_M and Tiel Coder UD-Q3_K_XL part from the first tokens with int8, your Ornith UD-Q6_K and MiniCPM-V Q6_K don't. The quickest way to see it on your side is the file setup gives an 8 GB card by default: |
`--batch K --check-batch K` reported a divergence under --kv int8 that no
prompt of mine had shown: sequence 1 matched for all its tokens, sequence 4 for
none. The same run under --kv fp16 passed every sequence, so it was this
branch's own path and not the batch machinery.
Two things were wrong in `Qwen35::attn_rows` (step()'s rows, one sequence a row).
The fp16 twin, Block::k_f16, is one pair for the whole model -- that is what
the alloc comment buys (a pair a layer would cost as much VRAM as the
quantized cache saves). A prompt batch can use it because it is a single
sequence moving forward, which is the invariant f16_owner_/f16_rows_ keep.
`attn_rows` converted each row's own sequence into that same pair, so four
sequences at row 100 all wrote the pair's row 100: the last row converted won,
and every earlier one had already been read against another sequence's keys.
It also read through cuBLAS (`attend_batch`) while a decode reads the codes
directly with the dp4a kernels (`attend_q8`). Those two orders of products do
not agree to the bit, so even with the twin fixed, a row read on its own
through eval() and the same row inside a step() came out different -- and a
greedy answer built from either goes its own way.
The rows now read their own sequence's codes with `attend_q8`, the reader a
decode takes, and the prompt side of a --batch run does the same instead of
going to cuBLAS, so one request has one order of products whether or not it
shares a step() with another. Without --batch nothing changes: a prompt keeps
the cuBLAS path, which is 1.7x faster at 16k and 2.4x at 32k, and its rows are
not compared against anything.
Measured, Ornith-1.5-35B-A3B UD-Q6_K, --expert-swaps 0 --cpu-share 0 --window 1
--tail-skip 0 --threads 13 --batch 4 --check-batch 4 --max-new 4:
6,000 ids, Random(6).randrange(100, 3000), --kv int8
before: 4/4, 2/4, 1/4, 0/4 after: 4/4, 4/4, 4/4, 4/4
2,048 ids of the same noise, --batch 2 --max-new 1
before: want 565, got 197 after: want 197, got 197
real text, 256 and 512 ids, 5,500 ids, --kv int8 and --kv fp16: all 4/4
poly_kv_q8_test: 35 checks, ok.
The dump also stops forcing one row under --batch: step() writes several
sequences' rows at once and cannot do that in one, so a batch keeps its rows
and the dump prints the first of them.
|
Found it, and it was my bug rather than the batch machinery's — Your
Two separate faults in The fp16 twin ( Even with that fixed, a row read alone went through So both now take one reader: Also still passing after the change: real-text prompts at 256, 512 and 5,500 ids under both One more thing worth knowing for your runs: |
|
Confirmed on the laptop: 2c8ec1b with One cost of the fix I'd like to see measured before merging: with Could the prompt batches keep cuBLAS under |
The prompt reader differs between the two modes: without --batch a prompt row goes through attend_batch (dequantize into the fp16 twin once, then cuBLAS), with --batch it must take the same reader as a decode row because the two do not agree bit for bit and --check-batch compares them. Measured on an RTX 4070 Laptop, Ornith Q6_K, serial arms: 1.45x slower prompt at 16K and 1.86x at 32K against main's q8_0, level with --batch 0, decode unaffected. Also correct the 5090 note, which said the prompt path is fp16 either way -- true only without --batch.
The cost you asked to see:
|
| prompt | arm | --batch 4 prompt tok/s |
--batch 0 prompt tok/s |
decode tok/s (--batch 4) |
|---|---|---|---|---|
| 16K | this PR, --kv int8 |
811.1 | 1410.4 | 31.25 |
| 16K | main, --kv q8_0 |
1177.1 | 1409.2 | 31.25 |
| 32K | this PR, --kv int8 |
574.5 | 1379.3 | 28.65 |
| 32K | main, --kv q8_0 |
1068.6 | 1377.0 | 30.64 |
main / this PR, --batch 4: 1.45× at 16K, 1.86× at 32K. With --batch 0 the two are the
same number (0.999× and 0.998×) — so the whole gap is the prompt reader that --batch puts
the rows through, not the cache format or the decode path. Decode matches at 16K and sits 6.5%
below at 32K, which is the V-side being fp16 in int8k's half of the block; plain int8 keeps
both sides quantized.
Your estimate was 1.7× / 2.4×; measured here it is 1.45× / 1.86×. Same order, same direction.
Why option ① does not work
I tried exactly what you proposed before measuring — keep cuBLAS for the prompt stage under
--batch, and route only the decode-sized calls through attend_q8. It fails, and it fails at
one row, not four:
--batch 1 --check-batch 1on a 2048-id prompt:the first 0 of 1 tokens as on its own.- Same with
STRATA_FREEZE_EXPERTS=1and with--no-graphs: still fails, so it is neither the
expert split nor graph replay. --batch 0on the same prompt:197,197,62,220, twice.--batch 1under the proposal:
0,0,0,0— not a last-bit difference, the whole row is wrong.- The same proposal with
--kv fp16:197,197,62,220, fine. With a 256-id prompt: fine. With
2048 ids and--prompt-batch 2048(one call instead of two): fine.
So the break is int8 + a prompt read spread over more than one call + --batch, and the reason
is structural: Block::k_f16 / v_f16 is one pair of twins for the whole model, not one per
sequence or per row-batch (qwen35.cpp:496-500 says so; a pair per layer would be 737 MiB —
more than the int8 cache saves). attn_rows writes each row's prefix into that same pair, so
with more than one call in flight the later call converts over the row the earlier one has not
consumed yet. Under one call it happens to be ordered; across calls it is not.
The two readers therefore have to agree: eval() and attn_rows must both take the int8 path,
which is what the commit does, and what makes --batch pay the slower prompt reader. That is
the trade-off, and it is not one I can make disappear by ordering alone.
What this means for the server
"parallel": N with --batch N and --kv int8 will read long prompts at 574–811 tok/s where
main's q8_0 reads them at 1069–1177. If that is the shape of the deployment, --kv int8 is
the wrong arm on a card where the VRAM it saves is not buying experts back — the honest advice
in the docs is to keep q8_0 for a parallel server and take int8/int8k where the saved
cache actually converts into resident experts (and for Gemma 4, which q8_0 never covered).
Three ways forward, and I will take whichever you prefer:
- Leave it as is and document the
--batchcost (table above) indocs/Q8_KV.md. The PR is
correct today; this is a stated price. - Give
attn_rowsa per-sequence twin. It needs block-sized twins per concurrent sequence
(4 × 2 × per-row × layers), which on a 8 GB card is what we are trying to avoid — but on a
32 GB card it is affordable. A flag could pick between "one twin, one reader for everyone"
and "twins per sequence, cuBLAS for the prompt". More moving parts; I would rather not. - Take the
int8khalf first. Gemma 4 has no q8_0 at all today, soint8kthere costs
nothing against main — it is a pure gain. Narrowing this PR to Gemma 4 +int8kand leaving
qwen35moe'sint8for a later PR would let the useful half land now.
My preference is 1 + 3: state the cost, and if you would rather review a smaller diff, split the
qwen35moe int8 half out and land Gemma 4 first.
One correction to my earlier comment, since it matters for reading the table: the 803→? gain I
quoted was measured with --batch 0. Under --batch 4 the prompt stage is the regression above,
and the section in docs/Q8_KV.md that says "the prompt path is fp16 either way" is only true
when --batch is off — I am fixing that sentence in this PR.
|
Thanks for measuring it, and for trying the cuBLAS route before saying no; the shared twin explains why it can't work. Let's take 3. With So the PR I'd merge is Gemma 4's int8 / int8k: Gemma 4 has no quantized cache on main, so there it costs nothing against anyone. Keep the fixes you made along the way that apply to it (the per-sequence reads in the batch path, the |
|
Here is exactly how I would rebuild the branch if you confirm, so nothing is a surprise in the diff. Base. Rebased onto
Docs. Tests. Main's Verification before I push. Main's Two things I would rather have your call on than assume:
I will not touch the branch until you confirm this shape. |
|
That shape is right, thanks for laying it out.
Setup doesn't pass |
|
I took your "option 3" and rebuilt the branch. The rebuilt branch is #59 — please close #33 and work on #59 instead, since #33 still carries the old shape (qwen35moe on the int8 readers) and I would rather not have two PRs with the same code in different states. What changed in the rebuildThe scope is now exactly what you asked for: Gemma 4 only.
The symbols are renamed apart, so the two families read apart in the source
The int8 kernels are otherwise byte-for-byte the code that was in #33: the same
The measurements, on this machineRTX 4070 Laptop (8 GiB),
That is within run-to-run scatter, and I want to be plain about why: this file is Q4_K_M and this card cannot hold the experts either way (900 of 3840 routed experts in VRAM, the rest on the CPU), so the KV cache is a small part of what this configuration reads. The point of the format is the VRAM it gives back, not the clock — on a card that holds the whole model the KV cache is a bigger share, and your measurements would say more than mine. The answers: the four factual prompts keep their answer under both formats ( The kernel test ( Two things I would like you to decide
And one correction to my own earlier claim, since it is in this thread: the "3603 of 3603 lines bit-identical" I wrote earlier was an |
What this PR does
Rebased onto current main (PR 29 took the HIP rejection of
--kv q8_0out; this keeps that — theq8_0 alias rides the same HIP path it measured — and adds int8/int8k on top).
It replaces
--kv q8_0(PR #23, shipped in poly.5) with--kv int8—q8_0stays as an alias —and adds
--kv int8k(keys quantized, values fp16). Gemma 4, which #23 left out, gains the wholeset.
What was stored is the same numbers: the same Q8_0 rule (a fp16 delta per 32-value block, delta =
amax/127,
roundfcodes — pinned bitwise against ggml'squantize_row_q8_0_refby the test). Whatchanged is the layout and the readers. The old format kept ggml's 34-byte blocks and every attention
kernel dequantized them element by element into floats — two extra loads and four I2F/FMA per four
values, in every score and every output sum. The int8 here is a compute format: codes and deltas in
two arrays (codes back to back, so every four are one aligned word), the queries quantized into the
same form, a lane's run of a decode head inside one 32-value block, one
__dp4aper four productswith the delta multiplied once a run; wider heads take an int8 tensor-core MMA. A prompt batch
dequantizes the caches into an fp16 twin once — incrementally, O(rows) against the O(rows²)
attention — and hands the rows to cuBLAS's
attend_batchas before.Against the current q8_0: what the comparison shows
The reviewer asked for decode and prompt speed against the current q8_0 on the same card and
context, and the fp16-vs-q8 answer checks of Q8_KV.md. Both are in the updated
docs/Q8_KV.md:--dump-layerson one fixed prompt,one seed, three tokens (gemma-4-26B-A4B,
--kv int8k): 3603 of 3603 dump rows bitwise equal —the codes and deltas the two formats write are the same numbers, so the attention sees the same
cache.
engine-to-engine A/B over 36 requests gives decode ratio 1.0167, bootstrap 95% CI
[0.9813, 1.0577] (the CI contains 1.0 — statistically indistinguishable at decode), prefill
ratio 0.9844. The kernel-level test times the q8 decode path against the fp16 one at
0.88–0.96×: the decode reads half the bytes and pays the queries' quantize, which nets out level
with fp16 on this class of card — the saving is the VRAM (46.88% at full context, to experts), not
the tokens-per-second.
Ornith and Occamy, mean character agreement 1.0000.
The differences from #23, in one place
q8c_/q8d_), so the score kernel reads codes on both sides andnever touches a float key path.
attend_batch's prompt path dequantizes once into an fp16 twin (incrementally, on demand viatwin_coveron Gemma 4's vmm-backed cache) instead of re-dequantizing a key block per score row.--kv int8kmixes int8 keys with fp16 values (attend_q8_mixed) — a quarter of the fp16 cachesaved with less of a loss; Gemma 4's window-plus-full layers take it.
qwen35moe-only gate atEngine::openis gone (Gemma 4 joins). The int8 MMA tier isCUDA-only (
mma.syncPTX); the__dp4atiers compile throughhip_compat'sdp4a, so a HIPbuild takes the dp4a readers and the cuBLAS batch path (
STRATA_KV_MMA=0pins the same choice onCUDA). The AMD numbers in
docs/Q8_KV.mdare PR 29's elementwise readers, re-measured with thischange pending.
docs/Q8_KV.mdand the--kvrow ofdocs/GLM.mdare updated in place, not deleted.Verification
poly_kv_q8_test(replaceskv_q8_test.cpp): 35 checks pass — the store againstggml_quantize_chunkblock by block, byte-exact layout, and the three attention paths compared.--kv fp16path is bitwise what itwas.
card_row_byteswidth fix (it was short by 1.44× under int8k, which overran a parkedconversation's D2H blob — a real 0xC0000005): 36 requests, 0 failures.