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

update the code

8be551d3

opt 16 bit group gemm

ff76f390

perf: dedicated w4a16 tiles for int4 moe prefill + widen decode gemv

68bb067d

copilot-swe-agent assigned

copilot-swe-agent 95 days ago

copilot-swe-agent assigned

a32543254 95 days ago

fix: avoid illegal sycl::vec<uint16_t,32> in int4 MoE decode kernel

6cda6c8c

copilot-swe-agent requested a review from

a32543254 95 days ago

perf: widen int8 MoE decode to 32-wide vectorized ladder

ed2c78fa

perf: widen int2 MoE decode to 32-wide vectorized ladder

488fa34b

perf: widen fp8 MoE decode to 32-wide vectorized ladder

2284e77c

perf: hoist per-group/per-expert scale out of fp8 MoE decode inner loop

a039ad95

perf: defer per-group scale fold in native fp8 MoE prefill sym mainloop

f2c47391

merge: integrate opt_moe_kernel MoE kernel optimizations (WIP)

ec616218

fix: resolve S4 MoE dispatch conflict (opt_moe env-policy + our w4a16…

90ba4332

fix: restore INT4 MoE perf by defaulting S4 direct DPAS path ON

f4efc895

revert: restore int4 MoE to original base implementation

54669280

Changes before error encountered

515f839a

fix: restore fp8 per-expert MoE prefill large-M tile policy threshold

93cde8c0

fix: restore fp8 per-group MoE prefill large-M tile policy threshold

f8877638

revert: undo fp8 per-group MoE prefill threshold change

c8374a86

fix: switch ark.moe auto dispatch to total-token threshold

3b3d7148

Merge branch 'main' into copilot/optimize-int4-moe-performance

ff20e3a4

feat: route int4-sym MoE decode to shared S4 DPAS grouped GEMM

6c3d26f2

test: add int4-sym decode DPAS vs scalar perf comparison; docs

0c946597

fix: split orphaned int8 block out of int4-sym dpas perf test

ffd2926c

fix: lower default MoE auto-decode threshold 256 -> 32

87b617ea

test: add int4-sym DPAS vs scalar token-count sweep for MoE decode th…

23ba0dbf

perf: hoist per-group scale and split accumulators in int4 MoE decode…

a1480af3

perf: coalesce int4 MoE decode weight loads via on-device N-tiled repack

c20c0b2c

perf: block tokens-per-expert in coalesced int4 MoE decode GEMV

8c732ed8

fix: parameterize MoEDecodeRepackInt4 kernel name to avoid duplicate …

cc939f8d

perf: gate coalesced int4 decode on repack amortization to fix tiny-b…

590ce59a

fix: qualify TOKEN_BLOCK with moe_decode_detail namespace

d80a9a13

perf: hoist routing filter out of int4 coalesced decode hot loop

7db3d5f2

feat: add dedicated w4a16_policy_m_8 decode DPAS dispatch for S4 sym

c8ef1b8e

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

a3c82ced

perf: reuse persistent atomic buffer in S4 DPAS dispatch

33960828

perf: trim device-timeline overhead on S4 DPAS decode hot path

ba2c18d7

perf(ark): use A_avg_M tile ladder for int4 moe sym decode (match vll…

5b88c2ec

perf(ark): route decode-sized int4-sym moe through the asym scalar GEMV

1f7b6313

perf(ark): decode int4-sym nibbles through the asym unsigned path (^0…

211e95f8

perf(ark): vectorize int4 decode GEMV loads, hoist activation sums, p…

e4d25c5d

test,docs: cover the new int4 decode paths and document the tuning knobs

4ec9d7fe

perf: speed up int4-sym MoE decode with word-native nibble decode

610a005d

docs: update int4 decode perf notes for word-native nibble decode

a2827941

perf: word-native FP8 decode + decode-tuned FP8 DPAS dispatch

6cd81921

test: FP8 decode-mode + DPAS parity/perf coverage; docs: FP8 decode s…

3423da8f

fix: restore test_perf_fp8_per_tensor definition; make fp8_decode_mod…

c6f67a13

perf: add K-split lane-parallel FP8 MoE decode GEMV

7102f0ec

docs: describe the K-split FP8 decode GEMV (EN + CN)

e43b98f7

perf: N-block the FP8 K-split decode GEMV and drop the per-call routi…

51328f36

perf: raise the MoE auto-dispatch decode cutoff from 32 to 128 tokens

03ca8774

test: add qwen3-moe int4-sym prefill and decode perf benchmarks

353543da

docs: document Qwen3-MoE int4-sym prefill/decode perf tests

42949b74

Merge branch 'main' into copilot/optimize-int4-moe-performance

006d7f05

a32543254 marked this pull request as ready for review 75 days ago

a32543254 requested a review from

Zhenzhong1 75 days ago

opt the performance

21d58d66

Merge branch 'main' into copilot/optimize-int4-moe-performance

c01f7206

docs: point S4 decode dispatch references at the generated TU after m…

b3080671

style: wrap ARK MoE decode lines to the 120-char limit

86b1a105

Zhenzhong/moe opt (#2216)

67a23536

a32543254 requested a review from

luoyu-intel 71 days ago

Merge branch 'main' into copilot/optimize-int4-moe-performance

4553a863

fix: include fp8 dpas header in generated moe decode fp8 TU

8b129103

Merge branch 'main' into copilot/optimize-int4-moe-performance

41ae7db2

Merge branch 'main' into copilot/optimize-int4-moe-performance

24a3cc99

Zhenzhong1 approved these changes on 2026-08-31

luoyu-intel commented on 2026-09-01

Merge branch 'main' into copilot/optimize-int4-moe-performance

5995a2e7

Merge branch 'main' into copilot/optimize-int4-moe-performance

716fd475

Merge branch 'main' into copilot/optimize-int4-moe-performance

17d1e3c4

refactor: serve MoE DPAS atomic counter from DeviceMemoryPool scratch

35db0328

docs: clarify serialization assumption for the pooled atomic counter

f5ff0841

Merge branch 'main' into copilot/optimize-int4-moe-performance

20b24a95

refactor: serve MoE DPAS prefill atomic counter from the scratch pool

026a363f

docs: record that the pooled counter relies on the in-kernel reset

1a9f704c

Merge branch 'main' into copilot/optimize-int4-moe-performance

5c4a7ece

Merge branch 'main' into copilot/optimize-int4-moe-performance

5965dee5

perf: split FP8 decode DPAS instantiations into per-policy TUs

4bd927ea

docs: document FP8 decode per-policy TU split; note ladder location

532de07e

perf: split FP8 decode scalar-GEMV kernels into per-mode TUs

8e8f231f

docs: document FP8 decode per-mode TU split (EN + CN)

da9ea044

Merge branch 'main' into copilot/optimize-int4-moe-performance

fcf8587a

perf: split bf16/fp16 MoE GEMM tile policies into per-policy TUs

daa5aafe

Merge branch 'main' into copilot/optimize-int4-moe-performance

aae50b4f

refactor(ark): serve MoE decode scratch from DeviceMemoryPool

6269ffcc

docs(ark): document device-keyed decode scratch; harden OOM path

d4932b5f

Merge branch 'main' into copilot/optimize-int4-moe-performance

35b2a74c

Merge branch 'main' into copilot/optimize-int4-moe-performance

09547b06

luoyu-intel approved these changes on 2026-09-04

add the code

22d4ce21

a32543254 merged 21bab556 into main 59 days ago

a32543254 deleted the copilot/optimize-int4-moe-performance branch 59 days ago

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

Merged

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

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

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

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

Merged