perf(qwen36): grouped int8 tensor-core MoE prefill for Qwen3.6-35B-A3B

[9876543210-tc-0123456789]

Summary

Batched grouped int8 tensor-core MoE prefill for Qwen3.6-35B-A3B. On main this model's prefill falls back to the per-token decode loop — every routed expert's Q4_K/Q6_K weight is reloaded once per prompt token — so prompt processing runs at ~520–580 pp. This PR permutes the prompt's tokens into per-expert contiguous groups (block-aligned to the 128-row GEMM tile) so each expert weight is streamed once and reused across all of its tokens, and runs gate/up/down as mma.sync.m16n8k32 int8 tensor-core GEMMs. The batched FFN output tracks the decode MoE FFN (routing reuses launch_moe_router; output verified numerically faithful), so a following decode step is unchanged. The dense Qwythos/Qwen3.5 path is untouched (guarded by !dense_ffn). Scoped to kernels/ + runtime/.

Proof of speedup

⚠️ The on-device eval runs only when BOTH are true: (1) the box below is ticked, and
(2) at least one decode tok/s or prefill pp table shows a real end-to-end
improvement (after > before, filled from bench/scripts/bench.sh).

• Tested on RTX 5090 (sm_120)

Both main and this PR were built from source and benched on the same RTX 5090, one model-load per build (sweep mode), so the deltas are same-box.

Decode tok/s — unchanged (this PR does not touch the decode path):

	decode tok/s
before (main)	481.17
after (this PR)	481.78

Prefill pp tok/s (Qwen3.6-35B-A3B, from bench.sh / qwen3_gguf_bench <gguf> 128 <ctx>):

	prefill pp tok/s
before prefill (main)	562.75
after prefill (this PR)	7412.20

Full measured sweep (same box, fresh source builds of both):

ctx	before (main)	after (this PR)	speedup
512	578.14	1399.72	2.4×
4096	562.75	7412.20	13.2×
16384	543.17	12534.54	23.1×
32768	520.10	14026.89	27.0×

SPARKINFER_PREFILL_MOE=0 restores the token-loop baseline for A/B on the same binary.

# qwen3_gguf_bench 128, SPARKINFER_BENCH_SWEEP_CTXS=512,4096,16384,32768 (one load each)
BASELINE (main):
  {"512":{"decode_tps":500.71,"prefill_pp":578.14},"4096":{"decode_tps":481.17,"prefill_pp":562.75},
   "16384":{"decode_tps":462.36,"prefill_pp":543.17},"32768":{"decode_tps":431.96,"prefill_pp":520.10}}
CANDIDATE (this PR):
  {"512":{"decode_tps":500.90,"prefill_pp":1399.72},"4096":{"decode_tps":481.78,"prefill_pp":7412.20},
   "16384":{"decode_tps":463.03,"prefill_pp":12534.54},"32768":{"decode_tps":433.64,"prefill_pp":14026.89}}

Implementation

• New kernels/csrc/cuda/moe/moe_prefill_grouped_i8.cu (+ header): router-logits, block-aligned per-expert permute/gather, the grouped int8 m16n8k32 GEMM with a device-built (expert → row-tile) schedule, SwiGLU→int8, weighted scatter-back, and the gated shared-expert finalize.
• runtime/src/models/qwen35_prefill.cpp: MoE FFN branch; the projection int8/dequant scratch is sized on the true per-projection maxima (unchanged for the dense model, where moe_ffn is already the largest dim). UD mixed-quant down (Q5_K) uses a bf16 requant fallback.
• runtime/src/models/qwen35.cpp: batched_prefill_enabled widened to admit the Qwen3.6 MoE hybrid (env SPARKINFER_PREFILL_MOE).

[skyrocket2026]

⚪ sparkinfer auto-eval — ccf35db

metric	value
label	eval:none
Qwen3.5 score	eval-qwen35:REJECT (fail)
Qwen3.6 score	eval-qwen36:none (pass)
Qwen3.5 vs same-box main	295.16 tok/s → +1.0% (+2.9)
Qwen3.5 scored decode (128 ctx · 128-context)	298.1 tok/s
Qwen3.5 scored prefill	not measured (0 pp tok/s on all contexts)
Qwen3.5 correctness	top-1 93.3% · KL 0.0406
Qwen3.5 128-token no-regression gate	298.1 tok/s vs main 295.16 tok/s · pass
Qwen3.5 4k-context no-regression gate	284.24 tok/s vs main 285.0 tok/s · pass
Qwen3.5 32k-context no-regression gate	284.58 tok/s vs main 283.18 tok/s · pass
Qwen3.5 64k-context no-regression gate	285.1 tok/s vs main 283.23 tok/s · pass
Qwen3.5 4k prefill no-regression gate	287.71 pp tok/s vs main 320.41 pp tok/s · fail
Qwen3.5 32k prefill no-regression gate	284.07 pp tok/s vs main 300.15 pp tok/s · fail
Qwen3.5 64k prefill no-regression gate	284.11 pp tok/s vs main 279.06 pp tok/s · pass
Qwen3.5 128k prefill no-regression gate	0.0 pp tok/s · pass
Qwen3.6 vs same-box main	433.38 tok/s → +2.0% (+8.6)
Qwen3.6 scored decode (16384 ctx · 16k-context)	441.98 tok/s
Qwen3.6 correctness	top-1 93.9% · KL 0.0373
Qwen3.6 128-token no-regression gate	471.48 tok/s vs main 465.03 tok/s · pass
Qwen3.6 512-context no-regression gate	464.34 tok/s vs main 459.25 tok/s · pass
Qwen3.6 4k-context no-regression gate	444.87 tok/s vs main 440.42 tok/s · pass
Qwen3.6 16k-context no-regression gate	441.98 tok/s vs main 433.38 tok/s · pass
Qwen3.6 32k-context no-regression gate	420.01 tok/s vs main 412.82 tok/s · pass
Qwen3.5 optimize	eval:REJECT · 298.1 tok/s · fail
Qwen3.5 optimize — Qwen3.6-35B-A3B guard accuracy	top-1 93.9% · KL 0.0373 · pass
Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 128	298.1 tok/s · pass
Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 4k	284.24 tok/s · pass
Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 32k	284.58 tok/s · pass
Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 64k	285.1 tok/s · pass
Qwen3.6 optimize	eval:none · 441.98 tok/s · pass
Qwen3.6 optimize — Qwythos-9B (Q4_K_M) guard accuracy	top-1 93.3% · KL 0.0406 · pass
Qwen3.6 optimize — Qwen3.6-35B-A3B 128	471.48 tok/s · pass
Qwen3.6 optimize — Qwen3.6-35B-A3B 512	464.34 tok/s · pass
Qwen3.6 optimize — Qwen3.6-35B-A3B 4k	444.87 tok/s · pass
Qwen3.6 optimize — Qwen3.6-35B-A3B 16k	441.98 tok/s · pass
Qwen3.6 optimize — Qwen3.6-35B-A3B 32k	420.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.

[skyrocket2026]

❌ sparkinfer auto-eval — ccf35db

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	295.16 tok/s → +1.0% (+2.9)
Qwen3.5 scored decode (128 ctx · 128-context)	298.1 tok/s
Qwen3.5 scored prefill (4096 ctx · 4k-context)	287.71 pp tok/s · eval-prefill:none
Qwen3.5 correctness	top-1 93.3% · KL 0.0406
Qwen3.5 128-token no-regression gate	298.1 tok/s vs main 295.16 tok/s · pass
Qwen3.5 4k-context no-regression gate	284.24 tok/s vs main 285.0 tok/s · pass
Qwen3.5 32k-context no-regression gate	284.58 tok/s vs main 283.18 tok/s · pass
Qwen3.5 64k-context no-regression gate	285.1 tok/s vs main 283.23 tok/s · pass
Qwen3.5 4k prefill no-regression gate	287.71 pp tok/s vs main 320.41 pp tok/s · fail
Qwen3.5 32k prefill no-regression gate	284.07 pp tok/s vs main 300.15 pp tok/s · fail
Qwen3.5 64k prefill no-regression gate	284.11 pp tok/s vs main 279.06 pp tok/s · pass
Qwen3.5 128k prefill no-regression gate	0.0 pp tok/s · pass
Qwen3.6 vs same-box main	433.38 tok/s → +2.0% (+8.6)
Qwen3.6 scored decode (16384 ctx · 16k-context)	441.98 tok/s
Qwen3.6 correctness	top-1 93.9% · KL 0.0373
Qwen3.6 128-token no-regression gate	471.48 tok/s vs main 465.03 tok/s · pass
Qwen3.6 512-context no-regression gate	464.34 tok/s vs main 459.25 tok/s · pass
Qwen3.6 4k-context no-regression gate	444.87 tok/s vs main 440.42 tok/s · pass
Qwen3.6 16k-context no-regression gate	441.98 tok/s vs main 433.38 tok/s · pass
Qwen3.6 32k-context no-regression gate	420.01 tok/s vs main 412.82 tok/s · pass
Qwen3.5 optimize	eval:REJECT · 298.1 tok/s · fail
Qwen3.5 optimize — Qwen3.6-35B-A3B guard accuracy	top-1 93.9% · KL 0.0373 · pass
Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 128	298.1 tok/s · pass
Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 4k	284.24 tok/s · pass
Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 32k	284.58 tok/s · pass
Qwen3.5 optimize — Qwythos-9B (Q4_K_M) 64k	285.1 tok/s · pass
Qwen3.6 optimize	eval:none · 441.98 tok/s · pass
Qwen3.6 optimize — Qwythos-9B (Q4_K_M) guard accuracy	top-1 93.3% · KL 0.0406 · pass
Qwen3.6 optimize — Qwen3.6-35B-A3B 128	471.48 tok/s · pass
Qwen3.6 optimize — Qwen3.6-35B-A3B 512	464.34 tok/s · pass
Qwen3.6 optimize — Qwen3.6-35B-A3B 4k	444.87 tok/s · pass
Qwen3.6 optimize — Qwen3.6-35B-A3B 16k	441.98 tok/s · pass
Qwen3.6 optimize — Qwen3.6-35B-A3B 32k	420.01 tok/s · pass
regressions	regression-4k-pp, regression-32k-pp

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.