[HIP] gfx950 A8W4 fused_moe configs for DSv4.1 TP-sharded experts

[valarLip]

Motivation

DeepSeek-V4.1-Flash's routed experts are MXFP4. Its reference implementation runs them with MXFP8 activations (A8W4). On gfx950 with separated gate/up rows, aiter picks MXFP4 activations (A4W4) for these shapes instead. That leaves the routed output ~12% off the reference and costs ~1.4pp of GSM8K at TP4.

The fix on the ATOM side is to run these layers with interleaved gate/up rows and quant_dtype_a=fp8. However, aiter had no tuned A8W4 rows for the TP-sharded V4.1 shapes, so every call fell back to the heuristic.

Changes

dsv41_fp8_{tuned,untuned}_fmoe.csv: 64 rows.

• model_dim 5120; per-rank inter 640 (TP4) and 1152 (TP2).
• Target experts 384×6 and DSpark draft experts 128×3.
• 16 token counts, 1–32768.

How the rows were chosen:

• The default sweep misses a kernel family. TUNE_ONLY=flydslv2 only generates fused-fp8 stage1 + moe2_layout stage2. The heuristic's family (bf16-out _gui stage1 + non-layout stage2) never becomes a candidate.
• That family wins at decode sizes. A v2-only table was up to 1.25× slower than the heuristic at 32–64 tokens, and raised V4.1 TP4 decode TPOT from 8.2 ms to 12.8 ms.
• Per-shape merge. Each row is the faster of a v2 sweep and a TUNE_ONLY=flydsl (v1-family) sweep.
• Checked on the production op. Every row was then compared against the heuristic on the production fused_moe call.
• Heuristic fallback rows. In 7 shapes the tuned pick measured >3% slower than the heuristic. Those rows carry the heuristic's own kernel pair (from the no tuned FlyDSL config ... kn1/kn2 log), with us set to the measured op time. us1/us2/err are left empty, since the tuner never timed that pair.

csrc/kernels/topk_softmax_kernels.cu: one-line build fix for #5652.

• ROCm 7.1.1's clang rejects __shared__ alignas(64) DTYPE s_gate[...] ("'alignas' attribute cannot be applied to types").
• That breaks module_moe_asm and, with it, the MoE tuner.
• Changed to __attribute__((aligned(64))), as used elsewhere in csrc.

Test Plan

1. Per-shape production-op timing. Ran test_fmoe with interleaved gate/up, FP8 activations and AITER_BF16_FP8_MOE_BOUND=0 on all 64 shapes, once with this table and once with the pre-change table. Compared shape by shape.
2. No fallback warnings. Started the 7 heuristic-fallback rows and checked that no no tuned FlyDSL config warning appears.
3. End-to-end throughput. V4.1-Flash TP4, DSpark-5, ISL/OSL 1024/512, concurrency 8. Two runs per arm, alternating.
4. Accuracy. Routed-expert output on layers 0/20/39 against an fp32 reference of the checkpoint's MoE; GSM8K 1319 questions, 3-shot.

Test Result

1. Per-shape timing (this table ÷ heuristic):

Tokens	Ratio
1–2	0.77–0.94
32–64	0.89–1.02 (the v2-only table was 1.22–1.25)
16k–32k	0.80–0.99
7 fallback rows	1.00–1.02

2. Fallback rows: no no tuned FlyDSL config warnings for the 7 rows.

3. End to end: 2597 vs 2552 tok/s (heuristic), TPOT 8.39 vs 8.38 ms. Within run-to-run spread.

4. Accuracy:

• Routed-expert error vs the fp32 reference: 2.8% / 3.6% / 2.9% on layers 0/20/39, unchanged from before tuning.
• GSM8K: 0.9158, inside the A8W4 band of 0.9158–0.9287.

Notes for reviewers

• Tuner coverage. The tuner's default TUNE_ONLY=flydslv2 will produce slower-than-heuristic A8W4 tables at decode sizes unless the v1 family is also swept.
• Tuner vs production timing. At token 2048 on 384-expert shapes, the tuner's timing disagreed with the production op (1.14–1.18×). Those rows use the heuristic pair.
• test_moe_2stage.py import side effect. It runs test_output_buffer_contract (torch.compile) at module level, so importing it on a box whose Inductor can't build fails before any test runs.

[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
multigpu	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: GPT-OSS-120B, DeepSeek-R1-0528, Kimi-K2.5
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 6125 --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.