prefill: the fused int8 prompt experts on Turing (sm_75), opt-in (+9 % prefill on RTX 2080 Ti)

[rafatxf]

On Turing the prompt path's experts always go through llama.cpp's MMQ: the fused int8 kernels (moe_fused.cu for the Strata Q2_0 pack, moe_fused_iq.cu's native_kernel for the native packs, #136) need sm_80. On an RTX 2080 Ti MMQ's Q2_0 kernel runs at ~20-30 TOPS (255 registers with spills, one CTA per SM) and is the largest single cost of a long prompt read.

This ports both to sm_75, opt-in on top of STRATA_PF_FUSED=1 with STRATA_PF_FUSED_SM75=1:

• cp.async → plain 16-byte copies (the barriers around them stay), prefetch.global.L2 kept;
• mma m16n8k32 / m16n8k16 with s8 operands → four / two m8n8k16 (rows g / g + 8 by k 0-15 / 16-31). The int32 sums are exact in any order, so d is what the sm_80 instruction gives; MAGIC is passed as the C operand of each pair's first m8n8k16;
• two activation stages below sm_80 (64 KB of shared memory per SM), smem_bytes / the launch / the occupancy query take the device's depth (host_astages());
• native_kernel takes Q2_0 gate/up (64-value blocks: a block per stage, as the down rows already read theirs), so the canonical Q2_0 pack (Q2_0 / Q2_0 in every layer) runs fused. gu_covered() lists Q2_0 on CUDA;
• on Turing only that pair is set up (DevInfo::t75: native_supported answers false for the others, which keep MMQ);
• Turing, Q2_0 only (T75_DEV): a stage's two 32-value halves share the block's d, so a row's dot is d (dx_0 dot_0 + dx_1 dot_1) - 11 % fewer instructions (the kernel is bound by its FP32 epilogue on Turing, not the tensor pipe). sm_80+ keeps its code and bits.

prefill_fused_iq_test gets the Q2_0 / Q2_0 pair, and skips the pairs a Turing card does not set up instead of failing.

Setup

RTX 2080 Ti 22 GB (sm_75), i7-6700, CUDA 13.0, SM 1650 MHz, 240 W; v0.1.43. Canonical Q2_0 pack, --prefill auto:16384, STRATA_BF16_TC=1, STRATA_SELECT_SIMT=2.

Kernels (prefill_fused_iq_test --only="Q2_0 / Q2_0", one layer, 512 experts, top 10)

tokens	MMQ path	fused
2,048	12.99 ms	6.61 ms	1.96×
8,192	30.75 ms	21.18 ms*	1.45×
16,384	55.02 ms	37.15 ms	1.48×

* before the Q2_0 epilogue change. Fused vs MMQ rel RMS 1.6e-5 - 2.7e-5; both 1.09e-2 against the FP32 reference (the int8 activations' own error). prefill_fused_moe_test (Strata pack) passes too: 5.29 ms fused vs MMQ 3.67×, rel RMS 6.7e-5.

Engine

98K prompt (16K chunks, two runs each, deterministic run to run): 58.3 → 53.3 s (+9.4 %). 4K: 3.63 → 3.54 s.

Accuracy in the engine: layer 0 of a 32K prompt (identical input on both paths), every one of the 163,840 (token, expert) rows of the expert output against MMQ's: median rel 1.6e-7, p99.9 2.6e-3, max 7.6e-3 (the rows where H's int8 rounding falls the other way); against the FP16 expert path (STRATA_PREFILL_MMQ=0) MMQ and fused have the same error (median 1.08e-2).

Teacher-forced, as in #655 / #2071 (the first P tokens through the prompt path, the next 2,048 through the verify windows, STRATA_LOGPOS_TOPK=20):

P	KL (MMQ ‖ fused)	KL (FP16 ‖ MMQ)	perplexity FP16 / MMQ / fused
16,384	0.0101	0.0090	7.308 / 7.255 / 7.501
32,768	0.0226	0.0110	4.561 / 4.557 / 4.354
65,536	0.0186	0.0262	4.110 / 4.096 / 4.122

The perplexity moves either way (+3.4 %, −4.5 %, +0.6 %) - trajectory noise: the QSA selection is discrete, so any change of the prompt path's rounding gives KLs of this size (enabling only one layer's fused experts already moves the first logit by KL ~1e-3, as much as the FP16 path in every layer). The verify phase's time over the same 2,048 tokens: MMQ 42.36 s, fused 42.28 s.

The server here will run with STRATA_PF_FUSED=1 STRATA_PF_FUSED_SM75=1.

🤖 Generated with Claude Code