[yangyuhuiling]

Adds a second gfx950 FlyDSL route for 1×32 ue8m0 (mxfp8) operands, beside the 4-wave mxscale_preshuffle kernel. Ported from mori, which vendored this pipeline from aiter's own gemm_a8w8_8wave.py and added the scaled path.

Stacked on #6368. The first three commits here are that PR; review it first. The mainloop copied into the new kernel carries its corrected wait_barrier count.

Why it is faster

v_mfma_scale_f32_16x16x128_f8f6f4 takes e8m0 scales as instruction operands, so a 32-wide ue8m0 GEMM needs no dequantisation arithmetic at all — no second accumulator, no running rescale, no division. One MFMA consumes a whole K=128 step with its four 32-blocks already dequantised in hardware. That is what separates it from the 128-wide blockscale kernels, whose promote/rescale chain exists precisely because a 128-wide block cannot be expressed that way.

The mainloop is gemm_a8w8_8wave's, with the scale operands threaded into the four MFMA calls. Two things around it are not, and both are about addresses rather than arithmetic:

• A and B are swapped. fx.gemm(atom, c, b, a, c) computes B^T A^T = (A B)^T, moving a lane's four accumulator values from four consecutive rows to four consecutive columns — 8 contiguous bytes in a row-major C, where the unswapped layout can only ever emit buffer_store_short. Two cross-lane stages (permlane16_swap, then a ds_bpermute lane transpose) take a row to 64 contiguous bytes. This is gcnasm's mfma_adaptor_swap_ab and the store stages from its quad-subtile template.

• The scales are K-block major and packed four to a dword. A block group's sixteen lanes want sixteen consecutive rows of one K block, so K-block major makes that one coalesced access where the quantiser's own [M, K/32] row major makes it sixteen — same instruction count, 3.2× the cache accesses, +50 to +67% on the whole GEMM. And a lane's four A tiles differ only by sixteen rows and share the K block, so packing their bytes into one dword lets opsel_b pick between them at no VALU cost (−5 to −7.5%). CK does the same in preShuffleScaleBuffer_gfx950.

Measured

Against gemm_a8w8_blockscale_bpreshuffle with 2-D e8m0 scales — the route aiter already takes for these operands, which dispatches to the batched kernel at B=1. Identical logical operands; the scale shuffles are weight-prep work and are not timed. MI355X, min of 3 medians per arm, one process per point.

N	K	M	tile	new	existing	Δ
5120	2048	1024	128	35.5 µs	42.7 µs	−16.9%
5120	2048	1280	128	37.1	48.4	−23.2%
5120	2048	1536	128	39.2	51.9	−24.4%
5120	2048	1792	256	42.0	55.5	−24.3%
5120	2048	2048	256	42.7	60.3	−29.2%
5120	2048	4096	256	65.2	91.8	−29.0%
5120	2048	16384	256	157.0	304.9	−48.5%
8192	1280	1024	128	30.6	43.6	−30.0%
8192	1280	1280	256	34.6	48.6	−28.7%
8192	1280	1536	256	34.5	50.5	−31.7%
8192	1280	1792	256	36.3	57.6	−37.0%
8192	1280	2048	256	37.9	61.4	−38.3%
8192	1280	4096	256	56.4	103.4	−45.4%
8192	1280	16384	256	175.8	334.0	−47.3%

Correctness

• Bit-identical to mori on four shapes covering both N tiles, including M=64 where the tail block masks its stores. That is the check that matters here: a torch reference says the port is close, this says it is the same kernel. It is also what would catch a silently dropped kwarg — fx.gemm takes **kwargs, so a misspelled scale_a compiles and computes an unscaled product.
• Against a torch reference, 30 configurations (3 shapes × 5 M × 2 tiles): zero elements off by more than 5%, rel_l2 ≈ 2e-4 throughout. ue8m0 scales are exact powers of two, so applying one is lossless and that residual is fp8 quantisation plus bf16 output rounding — nothing else hides in it.
• The existing ptpc kernel is unaffected. Mfma16x16x128 gains an optional scaled path; passing neither scale leaves the unscaled call byte-for-byte as it was, and test_gemm_a8w8_8wave.py still passes 6/6.

Integration

Deliberately not reachable from the tuned-CSV dispatch. The scale layout is incompatible with the mxscale_preshuffle kernel's, and the two buffers are indistinguishable by shape and dtype — converting on the critical path would eat the margin, and getting them crossed computes garbage rather than raising. So callers select this explicitly via gemm_a8w8_mxfp8_8wave_flydsl and shuffle with the two new helpers:

b_sh = shuffle_weight(w_fp8, (16, 16))              # shared with the ptpc kernel
a_sc = shuffle_mxfp8_a_scale(x_exps)                # [M, K/32] e8m0 bytes
b_sc = shuffle_mxfp8_b_scale(w_exps)                # [N/32, K/32] e8m0 bytes
gemm_a8w8_mxfp8_8wave_flydsl(x_fp8, b_sh, a_sc, b_sc, out)

The weight layout is shared: shuffle_weight(w, (16, 16)) serves both this and the ptpc 8-wave kernel, verified byte-identical to mori's preshuffle_b, so a server shuffles once.

block_n defaults to pick_block_n(M, N), which gates on the wide grid size rather than on M — the narrow tile exists to widen a launch-starved grid, and the grid depends on N as much as on M. Timing both tiles at 13 points confirms it picks the faster one everywhere; the table behind the constant is in the source.

Constraints

• M must be a multiple of 64, the packed A scale's group. It need not be a multiple of BLOCK_M: the grid is ceildiv(M, BLOCK_M) and the tail block masks its stores. The op rejects rather than pads, because the A scale is built over the padded M and repairing its packed layout afterwards is not a reshape — the caller has to pad before quantising.
• BLOCK_M is fixed at 256: the packed scale puts a lane's four 16-row M tiles in one dword, which is exactly that.
• K a multiple of 128 and at least 256; N a multiple of block_n and of 32.

Not included

mori's Mxfp8GemvOp (decode, M ≤ 32, its own kernel and tuned table) is a separate follow-up. So is pushing the swap-ab store stages back into the ptpc 8-wave kernel, where they are worth −6.7% bit-exactly; that wants its own bit-exactness check rather than riding along here.

🤖 Generated with Claude Code

[github-actions]

🏷️ CI Guide

Runs automatically on every PR:

• ✅ Pre-checks (submodule verification, code formatting)
• ✅ Aiter op tests (gfx942 + gfx950)
• ✅ Triton tests on MI35X (only when aiter/ops/triton/** or related paths are changed)

Extended tests (opt-in via labels):

Label	Tests
ci:gfx1250-ffm-triton	Run the five-shard gfx1250 FFM Triton test suite
ci:triton-300x	Run an additional Triton test job on MI300X in PRs (added automatically when gfx942 configs change); main branch always runs both MI35X and MI300X
ci:triton-355	Run the full Triton test suite on MI35X, not only the tests the change affects
ci:multi-gpu	Aiter multi-GPU tests on the 8-GPU runner
ci:sglang	SGLang integration tests: DeepSeek-R1-MXFP4 accuracy, Qwen 3.5 accuracy
ci:atom	ATOM benchmark: DeepSeek-R1-0528, GPT-OSS-120B
ci:atom_full	ATOM accuracy suite for PR and main models from ATOM models_accuracy.json
ci:vllm	vLLM benchmark: DeepSeek-V4-Pro, MiniMax-M3-MXFP4
ci:all	All standard extended tests (excludes ci:atom_full)

Only add ci:atom_full for FlyDSL or Triton upgrades.
Add labels via the sidebar or gh pr edit 6370 --add-label <label>

One backend per PR:
A PR changes one kernel backend: [Triton/Gluon] (Triton and Gluon count as one), [HIP], [ASM], [CK], [OPUS] or [FlyDSL]. If the title ends up with two backend tags, split the PR -- as stacked pull requests when one part cannot merge without the other.

PR title tags & labels:
Component tags ([Triton/Gluon], [HIP], [CK], [ASM], ...) are added to the PR title and as PR labels automatically from the changed files and re-synced on every push — change-type tags like [fix]/[Perf], op tags like [MLA], and human labels (ci:*) are left untouched. Add the no-auto-title label to stop the title rewrites; labels stay in sync either way.

[solinzby1]

This is duplicate functionality — it's already on main. Also, the 950 interface currently supports non-shuffle scale in the current architecture.The shuffle variant also exists, it's just not enabled.