perf(qwen35): chunk-parallel GDN scan + vector-store int8 dequant + single-pass row quantize + skinny GEMM (+23% prefill pp)

[fansilas]

Summary

Five kernels in the Qwythos batched prefill, each a new file wired as a one-line guard with its own
env toggle. ctx=4096 prefill 16704 → 20549 pp (+23.0%) on top of #474; 32k +18.6%, 64k +18.3%;
decode unchanged.

lever (new file)	what was wrong	before → after
Chunk-parallel GDN scan prefill_gdn_chunk.cu	pf_gdn_scan_kernel (my #398) is the last sequential stage: one rank-1 update per token, and all 128 column-warps of a v-head re-load the same k_t/q_t every step — ~5.3 TFLOP/s (~5% of fp32 peak), and not DRAM-bound. Rewritten in WY/UT chunk form (S~ = S/Γ cancels the gate, the rank-1 chain becomes a triangular solve): serial chain N → N/C, inner work becomes dense matmuls every state column reuses.	59.3 → 34.1 ms
Vector-store GGUF→int8 dequant dequant_rows_i8_fast.cu	deq_rows_i8_kernel (#464) maps thread t to value t, so each thread stores one byte and a warp emits 32 B where the memory system moves 128. 4 consecutive values/thread → 4-byte store, 128 B/warp. Bit-identical.	19.5 → 11.8 ms
Single-pass row quantize prefill_quant_rows.cu	pf_quantize_rows_i8 is <<<rows, 32>>> — one warp per row — and streams each row twice (amax, then scale+store). One block per row, row held in registers.	17.1 → 4.09 ms
Skinny bf16 GEMM prefill_gemm_skinny.cu	GDN gate projections (n_out=32) fall through to the 128×128-tiled pf_gemm → grid 1 × 32 = 32 blocks on a 170-SM part, 75% of each tile padding. Compute-bound (A is L2-resident), so: wmma + full grid, no padding.	7.5 → 3.54 ms
Coalesced Q4_K dequant dequant_gguf_fast.cu	deq_q4k_kernel gives one thread a 256-value super-block and scalar-stores all 256 outputs → lanes write 512 B apart, ~16x write amplification (~7% of DRAM peak). Warp per super-block, one 16-byte store per lane. Bit-identical.	4.42 → 0.35 ms (+19.05 ms off GGUF load)

Every existing line is untouched (guards: batched_prefill.cu +9, prefill_gemm_i8.cu +5,
dequant_gguf.cu +8). SPARKINFER_PREFILL_GDN_CHUNK, SPARKINFER_DEQUANT_ROWS_I8_FAST,
SPARKINFER_PREFILL_QUANT_ROWS, SPARKINFER_PREFILL_GEMM_SKINNY, SPARKINFER_DEQUANT_COALESCED —
each =0 restores the previous kernel.

Correctness: top-1 0.9707 / KL 0.0034 at ctx=4096 over 512 positions (main 0.9609 / 0.0047;
gate 0.90 / 0.20). Measured at 512 rather than 128, where the metric quantises to 0.78pp per token.

Proof of speedup

• Tested on RTX 5090 (sm_120)

Decode tok/s (no-regression — decode does not touch this path):

	decode tok/s
before (main)	291.69
after (this PR)	292.08

Prefill pp tok/s (Qwythos-9B Q4_K_M, best context = --ctx 4096):

	prefill pp tok/s
before prefill (main)	16704.4
after prefill (this PR)	20549.4

ctx	before	after	delta		ctx	decode before	decode after
4096	16704.4	20549.4	+23.0%		128	301.13	301.52
32768	17534.0	20844.6	+18.6%		4096	291.69	292.08
65536	17580.5	20791.1	+18.3%		32768	290.97	290.98
					65536	291.05	291.33

Measured exactly as bench_sweep_run does — one process, all scored contexts, so max_seq is sized
for 64k as in the eval. Five interleaved A/B pairs, same binary toggled. Each scoring context clears
on its own (4k 5/5, 32k 5/5, 64k 5/5), so the result does not hinge on which context is selected.

# RTX 5090, idle box, main @ d72b98a, eval's own sweep shape:
#   SPARKINFER_EVAL_PREFILL=1 SPARKINFER_BENCH_SWEEP_CTXS=128,4096,32768,65536 \
#   SPARKINFER_BENCH_SWEEP_REPS=1 ./build/runtime/qwen3_gguf_bench <gguf> 128 sweep
# 5 interleaved pairs, gain on each run's best context (build_pp_score_select's rule):
#   pair1 ctx=4096  17037.4 -> 20885.3  +22.6%      pair4 ctx=4096  16685.8 -> 20522.4  +23.0%
#   pair2 ctx=4096  17001.2 -> 20597.4  +21.2%      pair5 ctx=4096  16219.2 -> 20425.4  +25.9%
#   pair3 ctx=4096  16704.4 -> 20549.5  +23.0%

# Correctness — qwen3_gguf_prefill_check <gguf> 4096 512:
  all toggles=0 (main):  TOP1 492/512 0.9609   KL 0.00473
  all toggles=1 (this):  TOP1 497/512 0.9707   KL 0.00338

# Dequant bit-exactness — 51.2M pseudo-random Q4_K values, FNV-1a over the raw bf16 bits:
  SPARKINFER_DEQUANT_COALESCED=0 -> 14116085882651877130
  SPARKINFER_DEQUANT_COALESCED=1 -> 14116085882651877130

# nsys (--cuda-graph-trace=node, ctx=4096), per-prefill (reps=1-vs-2 diffed):
  pf_gdn_scan  59.3 -> 34.1 | deq_rows_i8 19.5 -> 11.8 | pf_quantize_rows_i8 17.1 -> 4.09
  pf_gemm(bf16) 7.5 -> 3.54 | deq_q4k 4.42 -> 0.35     | pf_gemm_i8 (#474) 109.5 -> 109.5 (untouched)

Relation to open PRs

• No open non-draft PR shares any file with this one (perf(qwen36): grouped int8 tensor-core MoE prefill for Qwen3.6-35B-A3B #403, perf(qwen36): Q6_K->Q4_K requant of LM head #445, feat(runtime): MTP speculative decode for Qwythos NextN GGUF #509 → qwen36 / runtime; feat(prefill): Qwen3.6 MoE batched prefill + parity fixes #482
  is a draft), so the copycat guard never computes containment for it.
• perf(qwen35): int8-native prefill GEMM (mma.sync m16n8k32) for Qwythos #474 (merged) is untouched: its mma.sync m16n8k32 GEMM is now 44% of the prefill wall. I did
  try replacing its XOR swizzle (whose comment notes rows 4 apart collide 2-way) with an 80 B padded
  stride — bank-conflict-free but 20% slower (109.5 → 131.7 ms), because the conflict was already
  hidden behind the mma pipe. Its shape is the right one.
• Lever 2 improves perf(qwen35): fused Q4K/Q6K->int8 dequant + token-parallel GDN conv + lane-parallel fallback attention (+78% pp @32k on top of #465) #464's merged deq_rows_i8_kernel (store width only — its fused decode is kept
  verbatim, output bit-identical). Also builds on my merged perf(qwen35): batched prompt prefill for Qwythos (one weight-amortized GEMM pass) #398 and perf(qwen35): int8 tensor-core prefill attention for Qwythos #465.

[skyrocket2026]

❌ sparkinfer auto-eval — c14c929

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	16083.12 tok/s → +23.3% (+3745.4)
Qwen3.5 scored decode (4096 ctx · 4k-context)	286.64 tok/s
Qwen3.5 scored prefill (4096 ctx · 4k-context)	19828.57 pp tok/s · eval-prefill:XL
Qwen3.5 correctness	top-1 91.9% · KL 0.036
Qwen3.5 128-token no-regression gate	296.41 tok/s vs main 296.56 tok/s · pass
Qwen3.5 4k-context no-regression gate	286.64 tok/s vs main 286.65 tok/s · pass
Qwen3.5 32k-context no-regression gate	285.01 tok/s vs main 285.08 tok/s · pass
Qwen3.5 64k-context no-regression gate	285.07 tok/s vs main 285.09 tok/s · pass
Qwen3.5 4k prefill no-regression gate	19828.57 pp tok/s vs main 16083.12 pp tok/s · pass
Qwen3.5 32k prefill no-regression gate	20551.89 pp tok/s vs main 17616.92 pp tok/s · pass
Qwen3.5 64k prefill no-regression gate	20591.97 pp tok/s vs main 17754.56 pp tok/s · pass
Qwen3.5 128k prefill no-regression gate	0.0 pp tok/s · pass
Qwen3.6 vs same-box main	462.01 tok/s → +0.0% (+0.0)
Qwen3.6 scored decode (4096 ctx · 4k-context)	462.32 tok/s
Qwen3.6 correctness	top-1 0.0% · KL 1.8742
Qwen3.6 128-token no-regression gate	488.56 tok/s vs main 488.84 tok/s · pass
Qwen3.6 512-context no-regression gate	482.06 tok/s vs main 481.77 tok/s · pass
Qwen3.6 4k-context no-regression gate	462.32 tok/s vs main 462.01 tok/s · pass
Qwen3.6 16k-context no-regression gate	444.68 tok/s vs main 444.67 tok/s · pass
Qwen3.6 32k-context no-regression gate	416.7 tok/s vs main 417.12 tok/s · pass
Qwen3.5 optimize	eval:XL · 19828.57 tok/s · fail
Qwen3.5 optimize — Qwen3.6-35B-A3B guard accuracy	top-1 0.0% · KL 1.8742 · FAIL
Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 128	296.41 tok/s · pass
Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 4k	286.64 tok/s · pass
Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 32k	285.01 tok/s · pass
Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 64k	285.07 tok/s · pass
Qwen3.6 optimize	eval:REJECT · 462.32 tok/s · fail
Qwen3.6 optimize — Qwythos-9B (Q4_K_M) guard accuracy	top-1 91.9% · KL 0.036 · pass
Qwen3.6 optimize — Qwen3.6-35B-A3B 128	488.56 tok/s · pass
Qwen3.6 optimize — Qwen3.6-35B-A3B 512	482.06 tok/s · pass
Qwen3.6 optimize — Qwen3.6-35B-A3B 4k	462.32 tok/s · pass
Qwen3.6 optimize — Qwen3.6-35B-A3B 16k	444.68 tok/s · pass
Qwen3.6 optimize — Qwen3.6-35B-A3B 32k	416.7 tok/s · pass

Rejected — correctness gate: top1=0.0 (need >= 0.9), kl=1.874206 (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.

[fansilas]

Hi @skyrocket2026
Just rebased onto latest main. The earlier eval-qwen36:REJECT (Qwen3.6 top-1 0.0%) turned out to be the stale-Q int8 QK-norm bug fixed in #518 — it was in my old base, not this PR. Confirmed locally: Qwen3.6 output is byte-identical with my changes on vs off, and identical to a clean main build (all my kernels are on the qwen35 batched-prefill path, which the MoE model never hits). With the fix in the base, Qwythos holds at +23% / XL and correctness is unchanged. It should be green on the re-run.

[skyrocket2026]

❌ sparkinfer auto-eval — 95d846c

metric	value
label	eval:REJECT
Qwen3.5 score	eval-qwen35:REJECT (fail)
Qwen3.6 score	eval-qwen36:REJECT (fail)
Qwen3.5 scored decode (128 ctx)	0 tok/s
Qwen3.5 correctness	top-1 0.0% · KL 99
Qwen3.6 scored decode (128 ctx)	0 tok/s
Qwen3.6 correctness	top-1 0.0% · KL 99
Qwen3.5 optimize	eval:REJECT · 0 tok/s · fail
Qwen3.6 optimize	eval:REJECT · 0 tok/s · fail

Infra failure during eval (not a verified PR regression) — re-run recommended.

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.

[github-actions]

Closed: 3 evaluations with no verified speedup or rejection

This PR received 3 completed sparkinfer auto-evaluations labeled eval:none or eval:REJECT (limit: more than 2).

Open a fresh PR if you have a new optimization to try.