Repository navigation
[Config] [Perf] Add gfx950 a8w8 bpreshuffle and BF16 GEMM tunings for Qwen3.8 Flash Next - #6769
Conversation
…Qwen3.8-Flash-Next Tuned on MI355X (gfx950, 256 CUs) for amd/Qwen3.8-Flash-Next-Quark-MXFP4-AttnFP8-PLEFP8 at TP1, decode and prefill M values: - a8w8 bpreshuffle (PTPC-FP8 attention / GDN projections): +119 gfx950 rows. - BF16: +77 gfx950 rows (router incl. the 528-wide fused shared-expert router, lm_head 248320x2560, MTP and hyper-connection GEMMs). 5 existing gfx950 HC rows (10240x320 / 320x10240, M=32..4096) are replaced with faster solutions found by an exhaustive retune (e.g. 54.4 -> 41.3 us, 16.8 -> 13.6 us). - Untuned shape lists updated to match. No gfx942 rows are touched; no existing row is removed.
🏷️ CI GuideRuns automatically on every PR:
Extended tests (opt-in via labels):
One backend per PR: PR title tags & labels: |
#6769 tuned the same Qwen3.8-Flash-Next bf16 GEMMs on gfx950. Resolve qwen38_flash_next_bf16_{tuned,untuned}_gemm.csv as: - keep main's rows by default, in place; - add the 244 keys only this branch has; - of the 40 overlapping keys whose rows differ, take this branch's row where an A/B re-measurement on MI355X (same tensors, the tuned_gemm.solMap backend, run_perftest 10 warmup / 101 iters, median of 9 alternating rounds) shows it more than 3% faster: 12 keys (N=640 M=8..128, N=2560 K=2560 M=1/4/16/32/64, (128, 248320, 2560), (128, 13312, 2560)); - main's 33 gfx950 hipblaslt rows carry solution indices that do not exist in hipBLASLt 1.2.70204 (ROCm 7.2.4, the version aiter CI uses); hipb_mm exit()s on them, e.g. Qwen3.8 in ATOM dies at warmup on (16384, 10240, 2560) and the LM head (M<=64, N=248320) does the same. None of the kernel names map to a current solution. For the 18 that overlap, take this branch's row (tuned on ROCm 7.2.4, verified to run and match torch); re-tune the other 15 with gemm_a16w16_tune.py restricted to hipblaslt (the library #6769 chose for them) on ROCm 7.2.4; - untuned: main's list plus this branch's extra keys, plus the two tuned keys main's untuned list was missing, so every tuned key is listed. The 6 MiniMax-M3 rows in minimax_m3_eagle_bf16_tuned_gemm.csv (N=2560, K=6144) still share keys with this branch's rows on the new main, so they are dropped again as in the original commit; the merged bf16_tuned_gemm family has no duplicate shape. Co-authored-by: Cursor <cursoragent@cursor.com>
…le-tuning #6769 tuned the same Qwen3.8-Flash-Next a8w8 bpreshuffle GEMMs on gfx950. Resolve qwen38_flash_next_a8w8_bpreshuffle_{tuned,untuned}_gemm.csv as: - keep every row from main as is; - add the 10 keys only this branch has (N=2560, K=6144); - of the 31 overlapping keys whose rows differ, take this branch's row only where an A/B re-measurement on MI355X (same tensors, aiter's own launcher, run_perftest 10 warmup / 101 iters, median of 9 alternating rounds) shows it more than 3% faster: (128, 13312, 2560) and (256, 16384, 2560); - untuned: main's list plus this branch's extra keys, plus the two tuned keys main's untuned list was missing, so every tuned key is listed. The 9 MiniMax-M3 rows in a8w8_bpreshuffle_tuned_gemm_minimax_m3.csv that share keys with the N=2560, K=6144 rows are still duplicates on the new main, so they are dropped again as in the original commit; the merged a8w8_bpreshuffle_tuned_gemm family has no duplicate shape. Co-authored-by: Cursor <cursoragent@cursor.com>
Tuned on MI355X for amd/Qwen3.8-Flash-Next-Quark-MXFP4-AttnFP8-PLEFP8 at TP1, decode and prefill M values:
No gfx942 rows are touched; no existing row is removed.