[fxmarty-amd]

Motivation

https://huggingface.co/amd/MiniMax-M3-MXFP4 served with AITER_MXFP4_BF16 MOE backend in vLLM yields wrong results when using gate_mode=GateMode.INTERLEAVE, which it does by default: https://github.com/vllm-project/vllm/blob/aba01ef1f73d5bf24952c8f7d61febc509bbfefc/vllm/model_executor/layers/fused_moe/experts/rocm_aiter_moe.py#L400

Cause is AITER cktile producing wrong outputs in this case, details below.

Technical Details

This dispatches into:

aiter/aiter/fused_moe.py

Lines 3568 to 3584 in 48b1365

	if (
	gate_mode != GateMode.SEPARATED
	and dtype in [dtypes.bf16, dtypes.fp16]
	and q_type == QuantType.per_1x32
	and activation == ActivationType.Swiglu
	and q_dtype_w != dtypes.fp8
	and cktile_mxfp4_ok
	):
	return MOEMetadata(
	functools.partial(
	cktile_moe_stage1,
	n_pad_zeros=intermediate_pad // 64 * 64 * (2 if use_g1u1 else 1),
	k_pad_zeros=hidden_pad // 128 * 128,
	activation=activation,
	split_k=1,
	dtype=dtype,
	),

that uses split_k=1, activation == ActivationType.Swiglu.

Additionally, Minimax does not use bias, and we wrongfully fall into:

aiter/csrc/ck_tile_gemm_moe_2stages/include/moe_cktile2stages.h

Lines 47 to 57 in 48b1365

	template <typename ADataType,
	typename BDataType,
	typename AccDataType,
	typename CDataType,
	int activation,
	bool kHasBias,
	int split_k>
	struct moe_gemm1_heuristic_dispatcher
	{
	static MoeKernel dispatch(int M, int N, int K, int block_m) {}
	};

as cktile_moe_gemm1 calls moe_dispatch with:

aiter/csrc/ck_tile_gemm_moe_2stages/moe_cktile2stages.cu

Line 284 in 48b1365

bool has_bias = exp_bias.has_value();

that is has_bias=False.

split_k=1, activation=ActivationType.Swiglu, has_bias=False does not have a specialization and falls into the default that actually uses silu and thus produces wrong outputs.

Note that MOEMetadata.has_bias is not used here. This seems to have a different meaning.

Test Plan

See added test.

I verified that test_mxfp4_bf16_swiglu_interleaved_matches_separated is indeed failing on aiter==v0.1.22.post1

Test Result

Passing on this branch

[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 6022 --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.

[fxmarty-amd]

Note: I'll add test on vLLM side to assert AITER_MXFP4_BF16 MOE backend matches the reference EMULATION backend. This is currently not the case.

[zufayu]

[P1] Please keep the existing SiLU+bias path working. This condition still skips generating the (Silu, has_bias=true) specialization, while the new primary-template check turns that existing path into a hard failure. The exact-head MI35X run reproduces it in op_tests/test_moe_2stage.py with activation=0, has_bias=1, split_k_gt_1=0: https://github.com/ROCm/aiter/actions/runs/36763567340/job/110085522039. Please generate the required specialization (including GEMM2), or change the caller contract so a supported SiLU+bias invocation cannot reach this dispatcher, and add regression coverage for that combination.

[zufayu]

@aiter-bot review

[aiter-bot]

ROCm/aiter PR #6022 — [CK] [bug fix] Fix wrong cktile 2stages moe_gemm1_heuristic_dispatcher dispatch in case a template is not built, add template specialization for (Swiglu, HasBias=false, split_k=1)

Generates the missing (Swiglu, has_bias=false) CK-Tile MoE 2-stage template specializations that bias-less Swiglu callers (MiniMax-M3-MXFP4 under vLLM's AITER_MXFP4_BF16 backend) dispatch into, and replaces the empty-body (undefined-behavior) dispatcher fallbacks with loud TORCH_CHECK errors.

Review (advisory): ⚠️ NEEDS WORK
Validation (deterministic): NOT RUN — triage required validation (runtime csrc surface changed; single test target op_tests/test_moe_mxfp4_inter_dim_dispatch.py) but the auto-run was disabled (REVIEW_AUTO_VALIDATE=0) and no report was supplied — environment gap, not a PR defect
Perf (advisory): NOT RUN — the triaged target op_tests/test_moe_mxfp4_inter_dim_dispatch.py exposes no benchmark entry point (no --scenario bench, no perftest/@benchmark harness), so no base-vs-head timing exists for the newly enabled specialization

⚠️ [verified] The patch no longer applies to main (all three csrc hunks reject on merge target 79634b1), and main already carries this exact fix — #6478 (2026-10-16, "Fix missing SwiGLU non-split no-bias CK-Tile MXFP4 dispatch") landed the same has_bias codegen axis in gen_instances.py, the same fail-loud TORCH_CHECK fallbacks in include/moe_cktile2stages.h and moe_cktile2stages.cu, plus a Relu2 dispatch branch this head lacks — so reviewing or merging this PR describes a tree that no longer exists. Author must rebase onto main — where the change already lives — or close the PR in favor of #6478.
⚠️ [verified] op_tests/test_moe_mxfp4_inter_dim_dispatch.py's new monkeypatch.setenv("GPTOSS_SWIGLU_MXFP4_BF16_BOUND", "2147483647") can never take effect — aiter/fused_moe.py captures that bound into a module-level constant at import (line 106 on the PR base) — and with TOKEN=32 both new setenvs are value-irrelevant (32 sits below the 256 defaults), so the test reaches cktile on the default bounds alone and silently measures something else if TOKEN is ever raised to 256 or more for prefill coverage (the SEPARATED side flips to fp4x2 activations). Author must patch fused_moe_module._SWIGLU_MXFP4_BF16_BOUND directly or drop the inert setenvs and state the TOKEN-below-256 assumption.
⚠️ [verified] The dispatch surface changed (new (Swiglu, has_bias=false) specializations plus new error paths) but no base-vs-head timing exists — op_tests/test_moe_mxfp4_inter_dim_dispatch.py exposes no benchmark entry point, so the latency of the newly enabled specialization that the INTERLEAVE/Swiglu/no-bias route now selects is unmeasured. Author must point at a timing harness for the new specialization or state that latency parity is out of scope for this fix.
📝 [verified] gen_instances.py now instantiates swiglu with both bias variants across every (a_type, c_dtype, split_k) — on gfx950 the top-level codegen combos grow from 6 to 9, raising JIT compile time and the module_moe_cktile2stages build size for every build. Reviewer should ask whether the added build cost is acceptable (main accepted the same growth in #6478).

[github-actions]

No activity for 15 days, so this is now labelled stale. A push or a comment clears it; keep-open exempts it. Nothing is closed automatically.