Optimize INT4 MoE prefill/decode (sym) XPU kernels with dedicated w4a16 tiling - #2111

Draft

a32543254 with Copilot wants to merge 49 commits into

mainfrom

copilot/optimize-int4-moe-performance

Commits on Jul 31, 2026

update the code
a32543254
committed
8be551d
opt 16 bit group gemm
a32543254
committed
ff76f39
perf: dedicated w4a16 tiles for int4 moe prefill + widen decode gemv
Copilot
authored
68bb067
fix: avoid illegal sycl::vec<uint16_t,32> in int4 MoE decode kernel
Copilot
authored
6cda6c8
perf: widen int8 MoE decode to 32-wide vectorized ladder
Copilot
authored
ed2c78f
perf: widen int2 MoE decode to 32-wide vectorized ladder
Copilot
authored
488fa34
perf: widen fp8 MoE decode to 32-wide vectorized ladder
Copilot
authored
2284e77
perf: hoist per-group/per-expert scale out of fp8 MoE decode inner loop
Copilot
authored
a039ad9
perf: defer per-group scale fold in native fp8 MoE prefill sym mainloop
Copilot
authored
f2c4739
merge: integrate opt_moe_kernel MoE kernel optimizations (WIP)
Copilot
authored
ec61621
fix: resolve S4 MoE dispatch conflict (opt_moe env-policy + our w4a16 tiles)
Copilot
authored
90ba433
fix: restore INT4 MoE perf by defaulting S4 direct DPAS path ON
Copilot
authored
f4efc89
revert: restore int4 MoE to original base implementation
Copilot
authored
5466928
Changes before error encountered
Copilot
authored
515f839

Commits on Aug 4, 2026

Merge branch 'main' into copilot/optimize-int4-moe-performance
a32543254
authored
ff20e3a
feat: route int4-sym MoE decode to shared S4 DPAS grouped GEMM

Copilot
and
a32543254
authored
6c3d26f
test: add int4-sym decode DPAS vs scalar perf comparison; docs

Copilot
and
a32543254
authored
0c94659
fix: split orphaned int8 block out of int4-sym dpas perf test

Copilot
and
a32543254
authored
ffd2926
fix: lower default MoE auto-decode threshold 256 -> 32

Copilot
and
a32543254
authored
87b617e

Commits on Aug 5, 2026

test: add int4-sym DPAS vs scalar token-count sweep for MoE decode threshold tuning

Copilot
and
a32543254
authored
23ba0db
perf: hoist per-group scale and split accumulators in int4 MoE decode GEMV

Copilot
and
a32543254
authored
a1480af
perf: coalesce int4 MoE decode weight loads via on-device N-tiled repack

Copilot
and
a32543254
authored
c20c0b2
perf: block tokens-per-expert in coalesced int4 MoE decode GEMV

Copilot
and
a32543254
authored
8c732ed
fix: parameterize MoEDecodeRepackInt4 kernel name to avoid duplicate mangled names

Copilot
and
a32543254
authored
cc939f8
perf: gate coalesced int4 decode on repack amortization to fix tiny-batch regression

Copilot
and
a32543254
authored
590ce59
fix: qualify TOKEN_BLOCK with moe_decode_detail namespace

Copilot
and
a32543254
authored
d80a9a1
perf: hoist routing filter out of int4 coalesced decode hot loop

Copilot
and
a32543254
authored
7db3d5f
feat: add dedicated w4a16_policy_m_8 decode DPAS dispatch for S4 sym

Copilot
and
a32543254
authored
c8ef1b8
docs+test: document decode S4 DPAS path and add A/B parity assertion

Copilot
and
a32543254
authored
a3c82ce

Commits on Aug 6, 2026

perf: reuse persistent atomic buffer in S4 DPAS dispatch

Copilot
and
a32543254
authored
3396082
perf: trim device-timeline overhead on S4 DPAS decode hot path

Copilot
and
a32543254
authored
ba2c18d
perf(ark): use A_avg_M tile ladder for int4 moe sym decode (match vllm-xpu-kernels)

Copilot
and
a32543254
authored
5b88c2e
perf(ark): route decode-sized int4-sym moe through the asym scalar GEMV

Copilot
and
a32543254
authored
1f7b631
perf(ark): decode int4-sym nibbles through the asym unsigned path (^0x88 sign-flip)

Copilot
and
a32543254
authored
211e95f
perf(ark): vectorize int4 decode GEMV loads, hoist activation sums, pool the repack buffer

Copilot
and
a32543254
authored
e4d25c5
test,docs: cover the new int4 decode paths and document the tuning knobs

Copilot
and
a32543254
authored
4ec9d7f
perf: speed up int4-sym MoE decode with word-native nibble decode

Copilot
and
a32543254
authored
610a005
docs: update int4 decode perf notes for word-native nibble decode

Copilot
and
a32543254
authored
a282794
perf: word-native FP8 decode + decode-tuned FP8 DPAS dispatch

Copilot
and
a32543254
authored
6cd8192
test: FP8 decode-mode + DPAS parity/perf coverage; docs: FP8 decode section (EN+CN)

Copilot
and
a32543254
authored
3423da8
fix: restore test_perf_fp8_per_tensor definition; make fp8_decode_mode allocation-free

Copilot
and
a32543254
authored
c6f67a1
perf: add K-split lane-parallel FP8 MoE decode GEMV

Copilot
and
a32543254
authored
7102f0e
docs: describe the K-split FP8 decode GEMV (EN + CN)

Copilot
and
a32543254
authored
e43b98f

Commits on Aug 7, 2026

perf: N-block the FP8 K-split decode GEMV and drop the per-call routing sync

Copilot
and
a32543254
authored
51328f3
perf: raise the MoE auto-dispatch decode cutoff from 32 to 128 tokens

Copilot
and
a32543254
authored
03ca877