perf(qwen35): fused Q4K/Q6K->int8 dequant + token-parallel GDN conv + lane-parallel fallback attention (+78% pp @32k on top of #465)#464
Conversation
706b1fc to
8e7e027
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 scored decode (128 ctx) | ? tok/s |
| Qwen3.5 correctness | top-1 0.0% · KL ? |
| Qwen3.6 vs same-box main | 481.85 tok/s → +0.0% (+0.2) |
| Qwen3.6 scored decode (512 ctx · 512-context) | 482.08 tok/s |
| Qwen3.6 correctness | top-1 93.0% · KL 0.0559 |
| Qwen3.6 128-token no-regression gate | 488.5 tok/s vs main 488.74 tok/s · pass |
| Qwen3.6 512-context no-regression gate | 482.08 tok/s vs main 481.85 tok/s · pass |
| Qwen3.6 4k-context no-regression gate | 462.55 tok/s vs main 462.59 tok/s · pass |
| Qwen3.6 16k-context no-regression gate | 450.26 tok/s vs main 450.44 tok/s · pass |
| Qwen3.6 32k-context no-regression gate | 425.39 tok/s vs main 425.37 tok/s · pass |
| Qwen3.5 optimize | eval:REJECT · ? tok/s · fail |
| Qwen3.6 optimize | eval:none · 482.08 tok/s · pass |
| Qwen3.6 optimize — Qwythos-9B (Q4_K_M) guard accuracy | top-1 94.0% · KL 0.0327 · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 128 | 488.5 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 512 | 482.08 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 4k | 462.55 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 16k | 450.26 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 32k | 425.39 tok/s · pass |
Rejected — primary (Qwythos-9B (Q4_K_M)) produced no verdict (infra error).
RTX 5090 (sm_120) · 128-token decode scored vs same-box main · built from source · correctness vs llama.cpp. Automated — not merged; merge manually after review.
|
The
The primary Qwythos scoring run never executed (
For reference, measured on RTX 5090 (sm_120) against post-#463 main tooling, this PR's claim is prefill-only: Qwythos pp Could the bot re-run the Qwythos primary eval? Happy to rebase/push again if that's what re-triggers it. |
8e7e027 to
e2f8dce
Compare
❌ sparkinfer auto-eval —
|
| metric | value |
|---|---|
| label | eval:REJECT |
| Qwen3.5 score | eval-qwen35:XL (fail) |
| Qwen3.6 score | eval-qwen36:REJECT (fail) |
| Qwen3.5 vs same-box main | 7662.72 tok/s → +83.6% (+6409.4) |
| Qwen3.5 scored decode (4096 ctx · 4k-context) | 286.84 tok/s |
| Qwen3.5 scored prefill (32768 ctx · 32k-context) | 14072.17 pp tok/s · eval-prefill:XL |
| Qwen3.5 correctness | top-1 93.3% · KL 0.038 |
| Qwen3.5 128-token no-regression gate | 296.29 tok/s vs main 296.7 tok/s · pass |
| Qwen3.5 4k-context no-regression gate | 286.84 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.5 tok/s vs main 285.44 tok/s · pass |
| Qwen3.5 4k prefill no-regression gate | 12821.44 pp tok/s vs main 7792.88 pp tok/s · pass |
| Qwen3.5 32k prefill no-regression gate | 14072.17 pp tok/s vs main 7662.72 pp tok/s · pass |
| Qwen3.5 64k prefill no-regression gate | 14182.92 pp tok/s vs main 7825.4 pp tok/s · pass |
| Qwen3.5 128k prefill no-regression gate | 0.0 pp tok/s · pass |
| Qwen3.6 vs same-box main | 462.2 tok/s → +0.0% (+0.0) |
| Qwen3.6 scored decode (4096 ctx · 4k-context) | 462.76 tok/s |
| Qwen3.6 correctness | top-1 87.8% · KL 0.0629 |
| Qwen3.6 128-token no-regression gate | 488.49 tok/s vs main 488.75 tok/s · pass |
| Qwen3.6 512-context no-regression gate | 481.81 tok/s vs main 481.93 tok/s · pass |
| Qwen3.6 4k-context no-regression gate | 462.76 tok/s vs main 462.2 tok/s · pass |
| Qwen3.6 16k-context no-regression gate | 450.12 tok/s vs main 450.27 tok/s · pass |
| Qwen3.6 32k-context no-regression gate | 425.15 tok/s vs main 425.36 tok/s · pass |
| Qwen3.5 optimize | eval:XL · 14072.17 tok/s · fail |
| Qwen3.5 optimize — Qwen3.6-35B-A3B guard accuracy | top-1 87.8% · KL 0.0629 · FAIL |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 128 | 296.29 tok/s · pass |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 4k | 286.84 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.5 tok/s · pass |
| Qwen3.6 optimize | eval:REJECT · 462.76 tok/s · fail |
| Qwen3.6 optimize — Qwythos-9B (Q4_K_M) guard accuracy | top-1 93.3% · KL 0.038 · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 128 | 488.49 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 512 | 481.81 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 4k | 462.76 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 16k | 450.12 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 32k | 425.15 tok/s · pass |
Rejected — correctness gate: top1=0.877551 (need >= 0.9), kl=0.062888 (need <= 0.2).
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.
|
The reject is a near-threshold flake on the Qwen3.6 guard, not a regression from this diff — three observations:
Could the Qwen3.6 guard be re-run? Happy to push a rebase to re-trigger if that's the preferred mechanism. |
…nv + fused Q4K/Q6K->int8 dequant Three batched-prefill bottlenecks on the Qwythos (Qwen3.5) path, found by profiling against llama.cpp (which led Qwythos prefill 3.7x at pp4096): 1. Prefill attention (prefill_attn_window.cu): the windowed/tiled kernels walked keys one per query-warp — a 5-shuffle reduction plus two dependent expf per key, latency-bound at 18.3 ms/layer. New lane-parallel kernel: one key per lane (zero shuffles to score 32 keys), tile-level online softmax, 4 queries per warp so each K/V smem read feeds 4 accumulators. 5.7 ms/layer (3.2x), same sink+window semantics, LANEPAR=0 restores. 2. GDN conv (batched_prefill.cu): pf_gdn_conv_kernel looped all N tokens sequentially with a __syncthreads per token on 64 blocks. The causal 4-tap conv over the raw in-buffer inputs is token-parallel; new kernel is one block per (token, head). 2.7 ms -> ~0.1 ms per layer. 3. Weight dequant (dequant_gguf.cu): every projection dequanted Q4_K/Q6_K to a bf16 scratch then re-read + row-quantized it to int8 for the int8 tensor-core GEMM. Fused deq_rows_i8 kernel decodes superblocks in registers and writes int8 + per-row scale directly (1B/value written, no bf16 bounce), and the N<=8192 int8 gate is lifted (own arena, falls back to bf16 GEMMs on alloc failure instead of the token loop). RTX 5090 (sm_120), Qwythos-9B Q4_K_M, default window path: prefill pp4096 6724 -> 12848 tok/s (+91%) prefill pp32768 2663 -> 12791 tok/s (+380%) decode unchanged (294.3 vs 294.6 tok/s, run noise) prefill_check 4096/16: top-1 13/16, KL 0.0164 (main kernels: 13/16, 0.0191) Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
e2f8dce to
b699eb5
Compare
|
To be precise about which gate blocked what, since the dual-model table is easy to misread:
So the only failing measurement in the whole eval is taken on Qwen3.6-35B — the model this diff never executes on (batched prefill early-returns on |
❌ sparkinfer auto-eval —
|
| metric | value |
|---|---|
| label | eval:REJECT |
| Qwen3.5 score | eval-qwen35:XL (fail) |
| Qwen3.6 score | eval-qwen36:REJECT (fail) |
| Qwen3.5 vs same-box main | 7662.72 tok/s → +83.7% (+6416.8) |
| Qwen3.5 scored decode (32768 ctx · 32k-context) | 285.42 tok/s |
| Qwen3.5 scored prefill (32768 ctx · 32k-context) | 14079.51 pp tok/s · eval-prefill:XL |
| Qwen3.5 correctness | top-1 93.8% · KL 0.0244 |
| Qwen3.5 128-token no-regression gate | 296.25 tok/s vs main 296.7 tok/s · pass |
| Qwen3.5 4k-context no-regression gate | 286.52 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.49 tok/s vs main 285.44 tok/s · pass |
| Qwen3.5 4k prefill no-regression gate | 12818.75 pp tok/s vs main 7792.88 pp tok/s · pass |
| Qwen3.5 32k prefill no-regression gate | 14079.51 pp tok/s vs main 7662.72 pp tok/s · pass |
| Qwen3.5 64k prefill no-regression gate | 14186.56 pp tok/s vs main 7825.4 pp tok/s · pass |
| Qwen3.5 128k prefill no-regression gate | 0.0 pp tok/s · pass |
| Qwen3.6 vs same-box main | 425.36 tok/s → -0.1% (-0.3) |
| Qwen3.6 scored decode (32768 ctx · 32k-context) | 425.05 tok/s |
| Qwen3.6 correctness | top-1 0.0% · KL 99.0 |
| Qwen3.6 128-token no-regression gate | 488.31 tok/s vs main 488.75 tok/s · pass |
| Qwen3.6 512-context no-regression gate | 455.95 tok/s vs main 481.93 tok/s · fail |
| Qwen3.6 4k-context no-regression gate | 460.63 tok/s vs main 462.2 tok/s · pass |
| Qwen3.6 16k-context no-regression gate | 449.68 tok/s vs main 450.27 tok/s · pass |
| Qwen3.6 32k-context no-regression gate | 425.05 tok/s vs main 425.36 tok/s · pass |
| Qwen3.5 optimize | eval:XL · 14079.51 tok/s · fail |
| Qwen3.5 optimize — Qwen3.6-35B-A3B guard accuracy | top-1 0.0% · KL 99.0 · FAIL |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 128 | 296.25 tok/s · pass |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 4k | 286.52 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.49 tok/s · pass |
| Qwen3.6 optimize | eval:REJECT · 425.05 tok/s · fail |
| Qwen3.6 optimize — Qwythos-9B (Q4_K_M) guard accuracy | top-1 93.8% · KL 0.0245 · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 128 | 488.31 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 512 | 455.95 tok/s · fail |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 4k | 460.63 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 16k | 449.68 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 32k | 425.05 tok/s · pass |
| regressions | regression-qwen35-32k, regression-qwen35-32k |
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.
|
Ran the accuracy gate locally to characterize the failing Qwen3.6 guard — full seed matrix, both trees built clean from source on an RTX 5090 (main =
Two facts fall out:
So the ≥0.90 top-1 bar sits inside the prompt-sampling noise band of the Qwen3.6 UD-Q4_K_M weights themselves — any PR (including an empty diff) has roughly a coin-flip guard on an unlucky seed. Re-run requested at |
|
✅ sparkinfer auto-eval —
|
| metric | value |
|---|---|
| label | eval:XL |
| Qwen3.5 score | eval-qwen35:XL (pass) |
| Qwen3.6 score | eval-qwen36:none (pass) |
| Qwen3.5 vs same-box main | 7662.72 tok/s → +83.5% (+6395.1) |
| Qwen3.5 scored decode (65536 ctx · 64k-context) | 285.21 tok/s |
| Qwen3.5 scored prefill (32768 ctx · 32k-context) | 14057.87 pp tok/s · eval-prefill:XL |
| Qwen3.5 correctness | top-1 93.4% · KL 0.0442 |
| Qwen3.5 128-token no-regression gate | 295.97 tok/s vs main 296.7 tok/s · pass |
| Qwen3.5 4k-context no-regression gate | 286.22 tok/s vs main 286.72 tok/s · pass |
| Qwen3.5 32k-context no-regression gate | 284.84 tok/s vs main 285.38 tok/s · pass |
| Qwen3.5 64k-context no-regression gate | 285.21 tok/s vs main 285.44 tok/s · pass |
| Qwen3.5 4k prefill no-regression gate | 12808.24 pp tok/s vs main 7792.88 pp tok/s · pass |
| Qwen3.5 32k prefill no-regression gate | 14057.87 pp tok/s vs main 7662.72 pp tok/s · pass |
| Qwen3.5 64k prefill no-regression gate | 14149.94 pp tok/s vs main 7825.4 pp tok/s · pass |
| Qwen3.5 128k prefill no-regression gate | 0.0 pp tok/s · pass |
| Qwen3.6 vs same-box main | 462.2 tok/s → +0.2% (+0.7) |
| Qwen3.6 scored decode (4096 ctx · 4k-context) | 462.9 tok/s |
| Qwen3.6 correctness | top-1 92.4% · KL 0.0597 |
| Qwen3.6 128-token no-regression gate | 488.89 tok/s vs main 488.75 tok/s · pass |
| Qwen3.6 512-context no-regression gate | 482.16 tok/s vs main 481.93 tok/s · pass |
| Qwen3.6 4k-context no-regression gate | 462.9 tok/s vs main 462.2 tok/s · pass |
| Qwen3.6 16k-context no-regression gate | 449.92 tok/s vs main 450.27 tok/s · pass |
| Qwen3.6 32k-context no-regression gate | 425.5 tok/s vs main 425.36 tok/s · pass |
| Qwen3.5 optimize | eval:XL · 14057.87 tok/s · pass |
| Qwen3.5 optimize — Qwen3.6-35B-A3B guard accuracy | top-1 92.4% · KL 0.0597 · pass |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 128 | 295.97 tok/s · pass |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 4k | 286.22 tok/s · pass |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 32k | 284.84 tok/s · pass |
| Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 64k | 285.21 tok/s · pass |
| Qwen3.6 optimize | eval:none · 462.9 tok/s · pass |
| Qwen3.6 optimize — Qwythos-9B (Q4_K_M) guard accuracy | top-1 93.4% · KL 0.0442 · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 128 | 488.89 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 512 | 482.16 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 4k | 462.9 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 16k | 449.92 tok/s · pass |
| Qwen3.6 optimize — Qwen3.6-35B-A3B 32k | 425.5 tok/s · pass |
Verified speedup over same-box origin/main — 14057.87 tok/s (main was 7662.72 tok/s).
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.
|
✅ Auto-merged as the round's |
…tensor-ai-lab#464 gittensor-ai-lab#465 from latest evals Record merged Qwen3.5 prefill wins (gittensor-ai-lab#455 XL @64k, gittensor-ai-lab#465 XL, gittensor-ai-lab#463 M, gittensor-ai-lab#464 XL @32k 14057.87 pp tok/s) and repair optimization journey.
Summary
Three profiling-guided fixes to the Qwythos (Qwen3.5) batched-prefill path. Rebased on and composing with the merged #463 (half-precision smem tiles) and #465 (int8-MMA prefill attention): the two big wins here are in stages #465 doesn't touch, and they stack on top of it.
dequant_gguf.cu). Every projection dequantized the weight to a bf16 scratch (2B/value written) then re-read and row-quantized it to int8 for the perf(qwen35): int8 tensor-core prefill GEMM for Qwythos [DRAFT — awaits #398] #422 int8 tensor-core GEMM.deq_rows_i8_kerneldecodes superblocks in registers and writes int8 + per-row scale directly (1B/value, no bf16 bounce), one block per row instead of one thread per superblock. TheN <= 8192int8-GEMM gate is also lifted — int8 scratch moves to its own arena, so an alloc failure at huge N degrades to the bf16 GEMMs instead of the token loop (this is why main's pp32k ran bf16 GEMMs until now).batched_prefill.cu).pf_gdn_conv_kernellooped all N tokens sequentially with a__syncthreads()per token on only 64 blocks (2.7 ms/layer, ~21 ms/layer at 32k). The causal 4-tap conv reads only in-buffer raw inputs (the prompt starts at position 0), so it is token-parallel: one block per (token, head), ~0.1 ms/layer. Same taps/SiLU/L2-norm math and conv-state layout.SPARKINFER_PREFILL_GDN_CONV_SEQ=1restores.prefill_attn_window.cu) — written before perf(qwen35): int8 tensor-core prefill attention for Qwythos #465 landed; with perf(qwen35): int8 tensor-core prefill attention for Qwythos #465's MMA kernel default-on this now only upgrades theSPARKINFER_PREFILL_ATTN_MMA=0fallback path (the scalar kernels walk keys one per query-warp with a 5-shuffle reduce + two dependentexpfper key; this one scores 32 keys per warp in parallel, 18.3 → 5.7 ms/layer, exact fp32 math). Kept env-gated (SPARKINFER_PREFILL_ATTN_LANEPAR=0restores the scalar kernels); no effect on the default MMA path.Decode paths, CUDA graphs, and the Qwen3.6 model 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):Correctness.
qwen3_gguf_prefill_check(batched prefill vs token-by-token fill, prefix 4096,cont 16), same box, same build flags, both on top of the merged int8-MMA attention:
this PR top-1 13/16 · KL 0.0164 · seed 290 vs main @ 9fdf68a top-1 13/16 · KL 0.0152 ·
seed 290 — identical top-1 and seed, KL at the same level. (The fused dequant quantizes from
the exact fp32 dequant instead of the bf16-rounded scratch — pre-#465 it measured tighter than
the then-main: KL 0.0164 vs 0.0191.)
ctest7/7 pass.