auto-round
feat: W4A8 ARK XPU MoE kernel (int4 weight / int8 compute) with prefill + decode
#2143
Open
Open
feat: W4A8 ARK XPU MoE kernel (int4 weight / int8 compute) with prefill + decode #2143
copilot-swe-agent
wants to merge
83 commits into
main
from
copilot/copilotoptimize-int4-moe-performance
update the code
8be551d3
opt 16 bit group gemm
ff76f390
perf: dedicated w4a16 tiles for int4 moe prefill + widen decode gemv
68bb067d
fix: avoid illegal sycl::vec<uint16_t,32> in int4 MoE decode kernel
6cda6c8c
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 'copilot/optimize-int4-moe-performance' into W4A8 MoE work
05ad2e16
feat: add W4A8 (int4 weight / int8 compute) ARK XPU MoE kernel
ef97754e
feat: add W4A8 MoE Python API and perf/accuracy benchmark script
f7898197
docs: add EN + CN README for the W4A8 MoE benchmark script
dc5ac502
fix: pin source tensors in the W4A8 prepack cache to prevent address-…
3e969ce1
copilot-swe-agent
changed the title ![pre-commit-ci[bot]](https://avatars.githubusercontent.com/in/68672?v=4&s=60)
[pre-commit.ci] auto fixes from pre-commit.com hooks
8cb7034c
perf(ark): W4A8 MoE decode K-split mapping + roofline-aware benchmark
1ff56798
docs(ark): document the W4A8 roofline, decode K-split mapping and new…
22a7a65c
fix(ark): compare the two decode mappings by SNR, not bit-equality
bb7bdce4
perf(ark): widen the prefill tile to 256x256 and make the decode chun…
ed2e3ca9
docs(ark): document the prefill tile / decode chunk knobs, MiniMax sh…
725b2c82
fix(ark): carry the 32-byte W4A8 decode chunk as 16-byte sycl::vec su…
35e89aee
perf(ark): top the W4A8 prefill tile ladder at 128x128 from the measu…
d112a37e
perf(ark): widen W4A8 act-quant messages and drop the float C shadow …
210c9ad4
test(ark): sweep the W4A8 act-quant message width and assert bit-iden…
1197fe21
perf(ark): add the measured 256x128 prefill rung and an interior-tile…
95e78533
perf(ark): take the 256-row prefill tile everywhere and batch the act…
fba16566
docs(ark): record the new sweep and the act-quant unroll knob (EN + CN)
a6b07c5f
docs(ark): pin the W4A8 roofline to the Arc Pro B60 it was measured o…
6d16b009
docs(ark): note the 24 GB B60 ceiling on the W4A8 prepack cache (EN +…
d8f3e298
perf(ark): 2D block store epilogue + single-pass activation quantizer…
b56f301b
docs: document the single-pass quantizer, the 2D store and the correc…
3660ad43
docs: correct the D-vs-weight byte figures for the 384 rows/expert batch
d7b09b11
perf: re-tune the W4A8 prefill tile ladder on the 384 rows/expert sweep
038f02ed
docs: refresh the W4A8 prefill tuning tables from the 384 rows/expert…
a561aa0e
test: add 8K-prompt prefill case and tile sweep to W4A8 perf harness
a1d2ab55
perf: retire the W4A8 prefill 256-row tile rung from the auto ladder
4318efaa
perf: cut W4A8 prefill traffic with int8-in and fused top-k reduce
fb2fca6a
docs: document the W4A8 prefill call contracts and correct the A-tile…
dce6713d
fix: pass prefetch_dist and reduce through the W4A8 prefill launch macro
a3b1241e
fix: bound the W4A8 pre-quantized contract test to one ULP instead of…
49307cd9
fix: build W4A8 pre-quantized contract cases on the int8 grid
064b3cf8
Login to write a write a comment.
Login via GitHub