perf(qwen35): weight-amortized Q4_K tile-GEMM prefill for Qwythos above 64k (5.5x pp @128k)#409
Conversation
⏸ Merge conflict — rebase before evalThis branch conflicts with Please rebase onto |
|
The round's |
|
Drafting this one — it only touches 128k prefill, and #419 took 128k out of the Qwythos eval, so there's no scored context left for it to move. It's rebased on current main and builds clean, and the Main question: is 128k coming back? "for now" sounds temporary, so I've parked this rather than closed it — if it's staying off, I'll just close it. Worth noting it's the one context no other open prefill PR reaches, since #398 caps at 65536 (its prefill attention is dense O(N²)). And if the reason for dropping it was eval wall-clock, main's 128k prefill run is ~7 minutes on its own — which is the thing this removes. |
|
The round's |
1be78fd to
44dcad4
Compare
⚪ sparkinfer auto-eval —
|
| metric | value |
|---|---|
| label | eval:none |
| Qwen3.5 score | eval-qwen35:none (pass) |
| Qwen3.6 score | eval-qwen36:none (pass) |
| Qwen3.5 vs same-box main | 283.16 tok/s → -0.0% (-0.1) |
| Qwen3.5 scored decode (65536 ctx · 64k-context) | 283.1 tok/s |
| Qwen3.5 scored prefill | not measured (0 pp tok/s on all contexts) |
| Qwen3.5 correctness | top-1 92.9% · KL 0.0375 |
| Qwen3.5 128-token no-regression gate | 294.67 tok/s vs main 295.19 tok/s · pass |
| Qwen3.5 4k-context no-regression gate | 284.58 tok/s vs main 284.95 tok/s · pass |
| Qwen3.5 32k-context no-regression gate | 283.1 tok/s vs main 283.2 tok/s · pass |
| Qwen3.5 64k-context no-regression gate | 283.1 tok/s vs main 283.16 tok/s · pass |
| Qwen3.5 4k prefill no-regression gate | 319.98 pp tok/s vs main 320.45 pp tok/s · pass |
| Qwen3.5 32k prefill no-regression gate | 299.89 pp tok/s vs main 300.16 pp tok/s · pass |
| Qwen3.5 64k prefill no-regression gate | 278.77 pp tok/s vs main 279.01 pp tok/s · pass |
| Qwen3.5 128k prefill no-regression gate | 0.0 pp tok/s · pass |
| Qwen3.6 vs same-box main | 480.19 tok/s → -0.0% (-0.2) |
| Qwen3.6 scored decode (128 ctx · 128-context) | 480.01 tok/s |
| Qwen3.6 correctness | top-1 90.2% · KL 0.0436 |
| Qwen3.6 128-token no-regression gate | 480.01 tok/s vs main 480.19 tok/s · pass |
| Qwen3.6 512-context no-regression gate | 472.69 tok/s vs main 472.96 tok/s · pass |
| Qwen3.6 4k-context no-regression gate | 453.2 tok/s vs main 453.43 tok/s · pass |
| Qwen3.6 16k-context no-regression gate | 446.29 tok/s vs main 446.65 tok/s · pass |
| Qwen3.6 32k-context no-regression gate | 425.01 tok/s vs main 425.33 tok/s · pass |
| Qwen3.5 optimize | eval:none · 283.1 tok/s · pass |
| Qwen3.5 optimize — Qwen3.6-35B-A3B guard accuracy | top-1 90.2% · KL 0.0436 · pass |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 128 | 294.67 tok/s · pass |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 4k | 284.58 tok/s · pass |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 32k | 283.1 tok/s · pass |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 64k | 283.1 tok/s · pass |
| Qwen3.6 optimize | eval:none · 480.01 tok/s · pass |
| Qwen3.6 optimize — Qwythos-9B (Q4_K_M) guard accuracy | top-1 92.9% · KL 0.0375 · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 128 | 480.01 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 512 | 472.69 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 4k | 453.2 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 16k | 446.29 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 32k | 425.01 tok/s · pass |
Within the significance gate — no verified speedup over same-box main.
RTX 5090 (sm_120) · 128/512/4k/16k/32k guarded · Qwen3.5 prefill at 4k/32k/64k · scored vs same-box main · built from source · correctness vs llama.cpp. Automated — not merged; merge manually after review.
44dcad4 to
bf0a845
Compare
❌ sparkinfer auto-eval —
|
| metric | value |
|---|---|
| label | eval:REJECT |
| Qwen3.5 score | eval-qwen35:REJECT (fail) |
| Qwen3.6 score | eval-qwen36:REJECT (fail) |
| Qwen3.5 vs same-box main | 285.35 tok/s → +0.1% (+0.2) |
| Qwen3.5 scored decode (65536 ctx · 64k-context) | 285.56 tok/s |
| Qwen3.5 scored prefill (32768 ctx · 32k-context) | 5129.66 pp tok/s |
| Qwen3.5 correctness | top-1 91.5% · KL 0.0409 |
| Qwen3.5 128-token no-regression gate | 296.65 tok/s vs main 296.54 tok/s · pass |
| Qwen3.5 4k-context no-regression gate | 286.71 tok/s vs main 286.56 tok/s · pass |
| Qwen3.5 32k-context no-regression gate | 285.47 tok/s vs main 285.32 tok/s · pass |
| Qwen3.5 64k-context no-regression gate | 285.56 tok/s vs main 285.35 tok/s · pass |
| Qwen3.5 4k prefill no-regression gate | 4879.35 pp tok/s vs main 6311.38 pp tok/s · fail |
| Qwen3.5 32k prefill no-regression gate | 5114.5 pp tok/s vs main 5475.75 pp tok/s · fail |
| Qwen3.5 64k prefill no-regression gate | 5129.66 pp tok/s vs main 5509.77 pp tok/s · fail |
| Qwen3.5 128k prefill no-regression gate | 0.0 pp tok/s · pass |
| Qwen3.6 vs same-box main | 425.37 tok/s → +0.0% (+0.0) |
| Qwen3.6 scored decode (32768 ctx · 32k-context) | 425.69 tok/s |
| Qwen3.6 correctness | top-1 88.0% · KL 0.0836 |
| Qwen3.6 128-token no-regression gate | 488.95 tok/s vs main 488.74 tok/s · pass |
| Qwen3.6 512-context no-regression gate | 482.21 tok/s vs main 481.85 tok/s · pass |
| Qwen3.6 4k-context no-regression gate | 462.81 tok/s vs main 462.59 tok/s · pass |
| Qwen3.6 16k-context no-regression gate | 450.48 tok/s vs main 450.44 tok/s · pass |
| Qwen3.6 32k-context no-regression gate | 425.69 tok/s vs main 425.37 tok/s · pass |
| Qwen3.5 optimize | eval:REJECT · 285.56 tok/s · fail |
| Qwen3.5 optimize — Qwen3.6-35B-A3B guard accuracy | top-1 88.0% · KL 0.0836 · FAIL |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 128 | 296.65 tok/s · pass |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 4k | 286.71 tok/s · pass |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 32k | 285.47 tok/s · pass |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 64k | 285.56 tok/s · pass |
| Qwen3.6 optimize | eval:REJECT · 425.69 tok/s · fail |
| Qwen3.6 optimize — Qwythos-9B (Q4_K_M) guard accuracy | top-1 91.5% · KL 0.0409 · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 128 | 488.95 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 512 | 482.21 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 4k | 462.81 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 16k | 450.48 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 32k | 425.69 tok/s · pass |
No context cleared the 2% significance gate while at least one context regressed. Auto-closing this PR.
RTX 5090 (sm_120) · 128/512/4k/16k/32k guarded · Qwen3.5 prefill at 4k/32k/64k · scored vs same-box main · built from source · correctness vs llama.cpp. Automated — not merged; merge manually after review.
|
This one only targets 128k — it's gated to That's also why the prefill gates show regressions: the branch is behind #422, so the bot built my stale copy (4879 pp @4k) against a main that already has the int8 GEMM (6311 pp). A rebase makes those numbers identical — it isn't something this PR's code does. The So the real question is 128k coming back? It's still the only prefill context that falls back to the token path, since batched prefill caps at 65536. If it's staying off, I'll close this rather than leave it half-evaluated. |
bf0a845 to
2a5705c
Compare
❌ sparkinfer auto-eval —
|
| metric | value |
|---|---|
| label | eval:REJECT |
| Qwen3.5 score | eval-qwen35:REJECT (fail) |
| Qwen3.6 score | eval-qwen36:none (pass) |
| Qwen3.5 vs same-box main | 285.38 tok/s → +0.0% (+0.0) |
| Qwen3.5 scored decode (32768 ctx · 32k-context) | 285.42 tok/s |
| Qwen3.5 scored prefill (4096 ctx · 4k-context) | 6311.1 pp tok/s |
| Qwen3.5 correctness | top-1 90.4% · KL 0.0378 |
| Qwen3.5 128-token no-regression gate | 296.34 tok/s vs main 296.7 tok/s · pass |
| Qwen3.5 4k-context no-regression gate | 286.73 tok/s vs main 286.72 tok/s · pass |
| Qwen3.5 32k-context no-regression gate | 285.42 tok/s vs main 285.38 tok/s · pass |
| Qwen3.5 64k-context no-regression gate | 285.45 tok/s vs main 285.44 tok/s · pass |
| Qwen3.5 4k prefill no-regression gate | 6311.1 pp tok/s vs main 7792.88 pp tok/s · fail |
| Qwen3.5 32k prefill no-regression gate | 5475.44 pp tok/s vs main 7662.72 pp tok/s · fail |
| Qwen3.5 64k prefill no-regression gate | 5510.17 pp tok/s vs main 7825.4 pp tok/s · fail |
| Qwen3.5 128k prefill no-regression gate | 0.0 pp tok/s · pass |
| Qwen3.6 vs same-box main | 462.2 tok/s → +0.1% (+0.4) |
| Qwen3.6 scored decode (4096 ctx · 4k-context) | 462.63 tok/s |
| Qwen3.6 correctness | top-1 95.4% · KL 0.0363 |
| Qwen3.6 128-token no-regression gate | 488.72 tok/s vs main 488.75 tok/s · pass |
| Qwen3.6 512-context no-regression gate | 482.03 tok/s vs main 481.93 tok/s · pass |
| Qwen3.6 4k-context no-regression gate | 462.63 tok/s vs main 462.2 tok/s · pass |
| Qwen3.6 16k-context no-regression gate | 450.3 tok/s vs main 450.27 tok/s · pass |
| Qwen3.6 32k-context no-regression gate | 425.35 tok/s vs main 425.36 tok/s · pass |
| Qwen3.5 optimize | eval:REJECT · 285.42 tok/s · fail |
| Qwen3.5 optimize — Qwen3.6-35B-A3B guard accuracy | top-1 95.4% · KL 0.0363 · pass |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 128 | 296.34 tok/s · pass |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 4k | 286.73 tok/s · pass |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 32k | 285.42 tok/s · pass |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 64k | 285.45 tok/s · pass |
| Qwen3.6 optimize | eval:none · 462.63 tok/s · pass |
| Qwen3.6 optimize — Qwythos-9B (Q4_K_M) guard accuracy | top-1 90.4% · KL 0.0378 · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 128 | 488.72 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 512 | 482.03 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 4k | 462.63 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 16k | 450.3 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 32k | 425.35 tok/s · pass |
No context cleared the 2% significance gate while at least one context regressed. Auto-closing this PR.
RTX 5090 (sm_120) · 128/512/4k/16k/32k guarded · Qwen3.5 prefill at 4k/32k/64k · scored vs same-box main · built from source · correctness vs llama.cpp. Automated — not merged; merge manually after review.
Closed: 3 evaluations with no verified speedup or rejectionThis PR received 3 completed sparkinfer auto-evaluations labeled Open a fresh PR if you have a new optimization to try. |
Summary
forward_tokeningests a prompt one token at a time, so a 128k prompt re-reads the whole Q4_K weight set 131072 times and prefill pp sits flat at ~285 tok/s, pinned to DRAM bandwidth. This PR ingests prompts longer than 64k in token chunks, issuing every projection as one Q4_K x Q8_1 dp4a tile GEMM so each weight superblock is read once per tile of token rows instead of once per token — weights are never dequantized, and 131072 goes 284.80 -> 1575.10 pp (5.5x). Gated to ctx > 65536, so 4k/32k/64k and the whole decode path are untouched.Proof of speedup
sm_120)Decode tok/s (end-to-end, from
bench/scripts/bench.sh— fill if this PR targets decode):Prefill pp tok/s (Qwythos / Qwen3.5 — fill if this PR targets prefill; use
--ctx 4096,32768,65536, or131072and copy theprefill ppline — report your best context):Context sweep — every context at or below the 65536 gate runs main's token loop byte-for-byte, so before/after there are literally the same code path:
How. The prompt is walked in 1024-token chunks; every projection (GDN
qkv/gate/alpha/beta/out, full-attentionq/k/v/o, dense SwiGLUgate/up/down) is issued as a single Q4_K x Q8_1 dp4a tile GEMM, so a weight superblock is read once per tile of token rows rather than once per token.Weights stay 4-bit from DRAM into registers — nothing is dequantized — and the arithmetic is the same
vec_dot_q4_K_q8_1the decode GEMV runs. The LM head (248320x4096 Q4_K, 572 MB = 10% of the weight set) runs for the final prompt token only. Scratch is O(chunk), not O(prompt), so 128k allocates what 4k does (VRAM 16.2 GB, unchanged above).Gated to ctx > 65536 (
SPARKINFER_PREFILL_MMQ_MINCTX, default 65536;SPARKINFER_PREFILL_MMQ=0disables). Everything that is not a weight read stays on the kernels decode already uses: the Gated-DeltaNet conv + recurrence run per token (a serial recurrence over state, not weights, so batching it buys nothing), and the full-attention layers call the same sink + sliding-window sparse-KV kernels (fa_kv_window_select/flash_decode_split_sparse) once per query, under the same depth-adaptive KV-split policy — so prefill attends to exactly the keys decode will later read. That is also why 128k is reachable: pastsparse_min_ctxthe baseline attends to sink + recent window, so a prefill that respects that window has no O(N^2) wall.Correctness
The chunked fill must leave the KV cache + Gated-DeltaNet/conv state such that decode continues exactly as the token-by-token fill would. Ingesting the same prompt both ways and comparing the next-token distribution at the prompt boundary (gate: top1 >= 0.90, KL <= 0.20):
Those runs pin
SPARKINFER_FAMMA4=0 SPARKINFER_ATTN_GQ8=0. Without them the two paths diverge at ctx >~ 5k for two reasons that are independent of this PR and already have open fixes — #389 (hd256 int8-MMA correctness) and #393 (sparse-KV combine drops the attention output gate). This path computes the gated, correctly-strided result; the token path currently does not, so the two disagree until those land. Both reproduce on the token path alone, without this PR — main's own argmax moves when either is switched off:The accuracy gate is unaffected either way: it teacher-forces
forward_tokenat 2048 (qwen3_gguf_score), which this path never touches.Relation to open PRs
#398 (
batched prompt prefill for Qwythos) is the only other PR that amortizes prefill weight reads, and it caps at 65536 (SPARKINFER_PREFILL_BATCHED_MAXCTX— its prefill attention is dense causal O(N^2)), falling back to the token path at 131072. That is the one context this PR touches, so the two do not overlap at any measured context, and if #398 merges first this PR still applies as-is.The implementations differ throughout: #398 dequantizes Q4_K to bf16 and runs cp.async + WMMA tensor-core GEMMs with a one-launch sequential GDN scan and its own dense int8 paged attention; this PR keeps weights 4-bit and runs a dp4a Q4_K x Q8_1 tile GEMM, reuses the in-tree per-token GDN kernels unchanged, and reuses the in-tree sparse-KV attention.
#387 (
skip LM head on prefill) keeps prefill token-by-token; skipping the LM head for non-final prompt tokens is inherent to a one-pass ingest here rather than a separable change.#389 / #393 are the two correctness fixes the A/B above depends on; no code overlap with this PR.
#345 (
grid-target KV-split cap for int8-MMA long-context) tunes the KV-split policy this PR mirrors — a trivial rebase if it lands first.