[AMD] MiniMax-M3 MXFP8 MI355X vLLM: nightly + AITER-on TP4 + emulatin linear / MiniMax-M3 MXFP8 MI355X vLLM：升级 nightly + 启用 AITER TP4 + emulation linear by hongxiayang · Pull Request #2003 · SemiAnalysisAI/InferenceX

hongxiayang · 2026-07-03T14:41:09Z

Bump minimaxm3-fp8-mi355x-vllm to nightly-09663abde
enable AITER for TP-only explicitly via --moe-backend aiter,
use --linear-backend emulation (beats stock native MXFP8 linear),
--max-num-batched-tokens 32768, and
a single TP4 conc 1-512 sweep for 1k1k and 8k1k.

Thanks for @chunfangamd to help inferenceX setup

中文说明

将 minimaxm3-fp8-mi355x-vllm 镜像升级至最新 nightly，显式启用 AITER（--moe-backend aiter），使用 --linear-backend emulation（优于原生 MXFP8 linear），设置 --max-num-batched-tokens 32768，并配置 TP4 单节点 1k1k/8k1k 并发 1-512 扫描。

github-actions · 2026-07-03T14:41:19Z

Thanks for the contribution! For vLLM & SGLang, please ensure that your recipes is similar to the official vLLM recipes and/or the SGLang cookbook

If it is not, please create a PR first before we can merge your single node PR into the master branch. Let's ensure that the documentation is first class such that the entire ML community can benefit from your hard work! Thank you

PR authors are responsible for ensuring that after merging, all GitHub Action jobs fully pass. A lot of the time, failures are just flakes and simply re-running the failed jobs will fix it. If re-running failed jobs is attempted, PR authors are responsible for ensuring it passes. See GitHub's docs on re-running failed jobs: https://docs.github.com/en/actions/how-tos/manage-workflow-runs/re-run-workflows-and-jobs#re-running-failed-jobs-in-a-workflow

As a rule of thumb, generally, PR authors should request a review & get a PR approval from the respective companies' CODEOWNERS before requesting a review from core maintainers.

If additional help is needed, PR authors can reach out to core maintainers over Slack.

感谢你的贡献！对于 vLLM 与 SGLang，请确保你的 recipe 与官方 vLLM recipes 和/或 SGLang cookbook 保持一致

如果不一致，请先创建一个 PR，之后我们才能将你的单节点 PR 合并到 master 分支。让我们确保文档保持一流水准，使整个 ML 社区都能从你的辛勤工作中受益！谢谢

PR 作者有责任确保合并后所有 GitHub Action 任务完全通过。 很多时候失败只是偶发抖动（flake），重新运行失败的任务即可解决。如果选择重新运行失败的任务，PR 作者有责任确保其最终通过。参见 GitHub 关于重新运行失败任务的文档：https://docs.github.com/en/actions/how-tos/manage-workflow-runs/re-run-workflows-and-jobs#re-running-failed-jobs-in-a-workflow

一般而言，PR 作者应先向相应公司的 CODEOWNERS 请求审阅并获得 PR 批准，然后再请求核心维护者审阅。

如需更多帮助，PR 作者可通过 Slack 联系核心维护者。

hongxiayang · 2026-07-03T20:13:23Z

https://github.com/SemiAnalysisAI/InferenceX/actions/runs/28667371208

claude · 2026-07-03T20:13:36Z

Claude finished @hongxiayang's task in 6m 3s —— View job

Review complete

Gather context (PR diff, changed files)
Review benchmark script changes (verified --moe-backend aiter and --linear-backend emulation are valid backend choices in the vLLM nightly at commit 09663abde)
Validate master config & perf-changelog rules (changelog entry present, appended at end of file, image is public Docker Hub)
Post review feedback

Summary: One 🔴 blocking issue — the new perf-changelog entry's pr-link is a placeholder pull/XXX and should be pull/2003 (inline suggestion posted). One 🟡 warning — the AITER master-switch gate was inverted so the (now-unused) EP branch sets VLLM_ROCM_USE_AITER=0 while --moe-backend aiter is hardcoded, an untested combination if an ep: search-space entry is ever re-added; please confirm intent or note it in the script comment. Everything else checks out: config sweep simplification matches the changelog, the image is publicly accessible, and the serve command follows the repo's formatting and EP-conditional patterns.

claude · 2026-07-03T20:19:25Z

+    - "Add --max-num-batched-tokens 32768 (env MAX_NUM_BATCHED_TOKENS) to enlarge the per-step prefill budget and improve TP4 throughput at high concurrency"
+    - "Enable the AITER master switch for TP-only (no-EP) runs via --moe-backend aiter: the earlier degenerate-output issue that forced it off for TP-only is fixed by vllm-project/vllm#47158, so TP4 uses the AITER_MXFP8 MoE path (verified GSM8K 0.9613 flex / 0.9621 strict on this nightly)"
+    - "Simplify both search spaces to a single TP4 conc 1-512 sweep for 1k1k and 8k1k (drop TP8 and TP4/EP4: TP8 has poor throughput/GPU and plain TP4 matches or beats TP4/EP4 at high concurrency)"
+  pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/XXX


🔴 BLOCKING: pr-link is a placeholder (pull/XXX) instead of the actual PR number.
Why it matters: The perf-changelog is the record tying config/perf changes to their PRs; every other entry has a real link, and this one will be a dead link once merged.
Fix:

Suggested change

pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/XXX

pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2003

claude · 2026-07-03T20:19:31Z

+# https://github.com/vllm-project/vllm/pull/47158 fix this
 if printf '%s\n' "${PARALLEL_ARGS[@]}" | grep -qxF -- '--enable-expert-parallel'; then
-    export VLLM_ROCM_USE_AITER=1
-else
    export VLLM_ROCM_USE_AITER=0
+else
+    export VLLM_ROCM_USE_AITER=1
 fi


🟡 WARNING: The EP branch now sets VLLM_ROCM_USE_AITER=0 while --moe-backend aiter is hardcoded unconditionally on the serve command (line 73) — the exact inverse of the previously validated EP configuration (which relied on VLLM_ROCM_USE_AITER=1 with the aiter fused-MoE path auto-selected).
Why it matters: No current search-space entry uses EP (this PR drops tp4/ep4), so nothing breaks in CI today. But the script still accepts EP_SIZE/DP_ATTENTION from configs, and anyone re-adding an ep: entry (which existed until this PR) will silently run an untested combination: forced aiter MoE backend with the AITER master switch off. Per vLLM's rocm_aiter_ops, is_fused_moe_enabled() requires VLLM_ROCM_USE_AITER=1, so master-switch-gated aiter paths (shared-experts fusion, attention/norm ops) behave differently from what was validated. If EP is intentionally unsupported now, please say so in the comment; the current # vllm-project/vllm#47158 fix this comment doesn't explain why EP flips the switch off.

claude · 2026-07-03T20:20:45Z

+# https://github.com/vllm-project/vllm/pull/47158 fix this
 if printf '%s\n' "${PARALLEL_ARGS[@]}" | grep -qxF -- '--enable-expert-parallel'; then
-    export VLLM_ROCM_USE_AITER=1
-else
    export VLLM_ROCM_USE_AITER=0
+else
+    export VLLM_ROCM_USE_AITER=1
 fi

+# Larger per-step prefill token budget to improve TP4 throughput at high


🟡 The EP branch of the AITER master-switch gate silently flipped from VLLM_ROCM_USE_AITER=1 to 0 with no explanation (only the TP-only flip is justified by #47158 in the PR description). The current search space is TP4-only so the EP branch is dead code and there is no CI impact, but this is a latent trap for future EP re-introduction and the sibling minimaxm3_fp8_mi355x_mtp.sh still carries the original EP=1/TP=0 gate with the detailed rationale comment, so the two scripts now silently disagree on AITER gating policy. Suggest either deleting the now-dead EP branch here or restoring EP=1 with a comment explaining the divergence; the replacement comment '# vllm-project/vllm#47158 fix this' is also ungrammatical and no longer explains WHY EP gets 0.

Extended reasoning...

Nit / maintainability — no CI impact today.

The AITER master-switch conditional at benchmarks/single_node/fixed_seq_len/minimaxm3_fp8_mi355x.sh:54-58 was flipped on both branches by this PR:

Branch Before After

EP-on VLLM_ROCM_USE_AITER=1 VLLM_ROCM_USE_AITER=0

TP-only VLLM_ROCM_USE_AITER=0 VLLM_ROCM_USE_AITER=1

The PR description and the new perf-changelog entry justify only the TP-only flip — vllm-project/vllm#47158 fixed the earlier degenerate-output issue that had forced AITER off for TP-only. The EP branch flip from 1→0 has no explanation, and the replacement comment # https://github.com/vllm-project/vllm/pull/47158 fix this is ungrammatical and only speaks to the TP-only side.

Why this is only a nit, not a bug: the refutation is correct that the current sweep never exercises the EP branch. configs/amd-master.yaml:2487-2494 was simplified to a single { tp: 4, conc-start: 1, conc-end: 512 } for both 1k1k and 8k1k, with no ep field and no dp-attn. So both DP_ATTENTION and EP_SIZE > 1 are false at runtime, --enable-expert-parallel never lands in PARALLEL_ARGS, and every sweep row falls through the else branch (VLLM_ROCM_USE_AITER=1). Additionally, --moe-backend aiter is now set unconditionally on the vllm serve command line, so the AITER MoE path is forced regardless of the master switch value.

Why it is still worth flagging:

The EP conditional is still physically present in the script (lines 55-57). Any future re-introduction of an EP row to the search space, or a manual EP_SIZE=4/DP_ATTENTION=true run, would silently take the regressed EP=0 setting instead of the previously validated EP=1.

The sibling benchmarks/single_node/fixed_seq_len/minimaxm3_fp8_mi355x_mtp.sh still carries the original logic (EP=1, else=0) with the detailed rationale comment intact at lines 88-96. The MTP recipe is on the older nightly-4559c43a image and has not yet moved to the nightly-09663abde that carries #47158's fix — but the divergence is now silent (no comment explains why the base recipe gates AITER differently on EP), which makes future maintenance harder.

The new comment is ungrammatical (fix this) and no longer documents why EP gets AITER=0. The old comment explicitly named the invariant: with EP the fused AITER MoE path is the auto-selected backend, and with TP-only AITER produced degenerate MiniMax-M3 output.

Step-by-step proof this is latent-only today:

configs/amd-master.yaml:2487-2494 — search space is - { tp: 4, conc-start: 1, conc-end: 512 } with no ep/dp-attn field.

In minimaxm3_fp8_mi355x.sh, DP_ATTENTION is unset → the if [ "$DP_ATTENTION" = "true" ] branch is false. EP_SIZE defaults to 1 → the elif [ "$EP_SIZE" -gt 1 ] branch is false. PARALLEL_ARGS = (--tensor-parallel-size 4).

printf ... | grep -qxF -- --enable-expert-parallel → no match → else → VLLM_ROCM_USE_AITER=1.

So all TP4 rows get AITER=1 as intended; EP=0 branch is unreachable this sweep.

If a future PR re-adds { tp: 4, ep: 4, ... } (which existed in the pre-PR search space!), that row would now set VLLM_ROCM_USE_AITER=0 — different from the pre-PR behavior for the same config, with no comment to warn the maintainer.

Suggested fix (pick one):

Delete the now-dead EP branch entirely: export VLLM_ROCM_USE_AITER=1 unconditionally, since the sweep no longer includes EP configs and #47158 makes AITER=1 safe for TP-only.

Or restore VLLM_ROCM_USE_AITER=1 on the EP branch and update the comment to explain that both branches now use AITER post-#47158 (matching the sibling MTP script's rationale, which will presumably follow when it upgrades to the new nightly).

Either fixes the latent trap and the stale/ungrammatical comment together.

github-actions · 2026-07-03T22:28:46Z

see unofficial run visualizer at https://inferencex.semianalysis.com/inference?unofficialRun=28683732009
see unofficial run visualizer at https://inferencex.semianalysis.com/evaluation?unofficialRun=28683732009

github-actions · 2026-07-03T23:49:27Z

see unofficial run visualizer at https://inferencex.semianalysis.com/inference?unofficialRun=28686016620
see unofficial run visualizer at https://inferencex.semianalysis.com/evaluation?unofficialRun=28686016620

…n linear Bump minimaxm3-fp8-mi355x-vllm to nightly-09663abde..., enable AITER for TP-only (vllm#47158 fix) via --moe-backend aiter, use --linear-backend emulation (beats stock native MXFP8 linear), --max-num-batched-tokens 32768, and a single TP4 conc 1-512 sweep for 1k1k and 8k1k. Co-authored-by: Cursor <cursoragent@cursor.com>

Signed-off-by: Hongxia Yang <hongxia.yang@amd.com>

…ecipe Match minimaxm3_fp8_mi355x.sh: add --moe-backend aiter, --linear-backend emulation, and --max-num-batched-tokens 32768 to the EAGLE3 MTP recipe (keeping --speculative-config and the EAGLE3 in-place patch). Make DRAFT_MODEL env-overridable for local testing. Verified locally on nightly-09663abde (TP4 conc512, local eagle3 draft): 5120/5120 completed, eagle3 patch a no-op (nightly natively supports SupportsEagle3). Co-authored-by: Cursor <cursoragent@cursor.com>

Make the ported MTP recipe changes take effect: bump the minimaxm3-fp8-mi355x-vllm-mtp image to nightly-09663abde... and simplify its search space to a single TP4 conc 1-512 sweep (drop TP8/EP layouts, matching the non-MTP config). Add the changelog entry. Co-authored-by: Cursor <cursoragent@cursor.com>

github-actions · 2026-07-04T04:23:41Z

see unofficial run visualizer at https://inferencex.semianalysis.com/inference?unofficialRun=28690689417
see unofficial run visualizer at https://inferencex.semianalysis.com/evaluation?unofficialRun=28690689417

hongxiayang · 2026-07-04T04:40:46Z

@functionstackx error again Error: File was unable to be removed Error: EACCES: permission denied, rmdir '/it-share/gharunners2/gharunner06/actions-runner/_work/InferenceX/InferenceX/benchmark_logs/logs/slurm_job-32929'

functionstackx · 2026-07-04T04:50:19Z

@hongxiayang here is the claude command to fix it. this error happens when someone from AMD runs an github workflow that generates files as ROOT

https://github.com/SemiAnalysisAI/InferenceX/blob/main/.claude/commands/clean-amd-mi355-runner-root-files.md

functionstackx · 2026-07-04T04:55:08Z

@hongxiayang The root-owned files on gharunner06 (/it-share/gharunners2/gharunner06/actions-runner/_work/InferenceX/InferenceX/benchmark_logs/logs/slurm_job-32929) have been cleaned up — I ran the /clean-amd-mi355-runner-root-files Claude command to scan and verify all MI355X TW runners. The scan came back clean (zero root-owned files across all gharunners*/gharunner* workspaces), so the offending files were already removed since the last failure. I've rerun the failed 8k1k jobs from run 28690689417.

Root cause reminder: multi-node disagg benchmarks submit Slurm jobs whose containers write logs as root into the runner workspace (benchmark_logs/logs/slurm_job-*). When a job is cancelled, teardown is skipped and root-owned dirs get stranded — the runner user can't delete them, which breaks actions/checkout workspace cleanup for every subsequent job on that runner.

Could you please advise your coworkers (cc @billishyahao @seungrokj @YukioZzz) to avoid generating files as root in the runner workspace? If Slurm containers must run as root, the cleanup/teardown script should chown or rm those files before the job exits. This would prevent the EACCES errors from recurring.

functionstackx · 2026-07-04T04:57:44Z

@hongxiayang adding as an reminder to CONTRIBUTING.md see #2043

Never write as root into the runner workspace — use /tmp or a dedicated staging path instead.
If root writes are unavoidable, add a trap cleanup EXIT that removes root-owned files before the job exits.
Test by cancelling a running benchmark mid-flight and verifying no root files remain.

One stranded root-owned directory bricks every subsequent job on that runner and blocks the entire AMD MI355X sweep queue for all contributors.

seungrokj · 2026-07-06T02:24:30Z

m3 vllm recipe PR,
https://github.com/vllm-project/recipes/pull/609/changes

seungrokj

lgtm

functionstackx · 2026-07-06T02:38:14Z

@seungrokj codeowner signoff did not pass. u need create an new PR checklist comment that follow the template exactly https://github.com/SemiAnalysisAI/InferenceX/blob/main/docs/PR_REVIEW_CHECKLIST.md

If you don't follow the PR CHECKLIST template inclduing the phrase As a PR reviewer and CODEOWNER, I have reviewed this and have in ur approval comment, it wouldn't trigger our codeowner signoff CI verifcation https://github.com/SemiAnalysisAI/InferenceX/blob/main/.github/workflows/codeowner-signoff-verify.yml

seungrokj

As a PR reviewer and CODEOWNER, I have reviewed this and have:

Additional detail section:

insert any additional info here
m3 vllm recipe PR,
https://github.com/vllm-project/recipes/pull/609/changes

Signed: seungrokj

Klaud-Cold · 2026-07-06T02:42:15Z

❌❌❌ REJECTED ❌❌❌

@functionstackx — the audited sign-off comment (4888636211) is not a checklist sign-off at all (it is a rejection notice addressed to @seungrokj), and its author is not an AMD CODEOWNER for configs/amd-master.yaml. A fresh sign-off following the exact template, posted by an AMD CODEOWNER with the recipe link in the additional detail section, is required.

❌ Check 0 (CODEOWNER): FAIL — configs/amd-master.yaml is owned by @billishyahao @chunfangamd @seungrokj @yctseng0211 @1am9trash; signer @functionstackx is not among them.
✅ Check 1 (sweep on in-PR commit): PASS — commit c696685 (in this PR) has green single-node 1k1k/8k1k / and eval / check-runs: https://github.com/SemiAnalysisAI/InferenceX/actions/runs/28690689417
✅ Check 2 (evals pass): PASS — GSM8K em_strict 0.952–0.960 across all 4 lanes (non-MTP + MTP, conc 128/512), on the same vllm/vllm-openai-rocm:nightly-09663abde… image as this PR's config.
❌ Check 3 (recipe link): FAIL — the sign-off has no additional detail section and no recipe link. (The recipes link posted elsewhere in the thread, vllm-project/recipes#609, covers the MXFP4 TP8 variant and does not reflect this PR's MXFP8 TP4 major args such as --linear-backend emulation.)
✅ Check 4 (reuse command): PASS — /reuse-sweep-run posted by @seungrokj (COLLABORATOR).
❌ Check 5 (latest template): FAIL — the sign-off contains zero checklist items from PR_REVIEW_CHECKLIST.md; it is not a sign-off body.
✅ Check 6 (upstream image / engine-first): PASS — both entries use upstream vllm/vllm-openai-rocm on MI355X; framework is vLLM itself.
✅ Check 7 (no architecture hacks): PASS — no --hf-overrides or FLOP-reducing knobs; --moe-backend aiter, --linear-backend emulation, --max-num-batched-tokens are kernel/scheduler choices with passing evals.
✅ Check 8 (spec-decode chat template): PASS — MTP benchmark client uses --use-chat-template (minimaxm3_fp8_mi355x_mtp.sh:223).

seungrokj

As a PR reviewer and CODEOWNER, I have reviewed this and have:

Additional detail section:

insert any additional info here
m3 vllm recipe PR,
https://github.com/vllm-project/recipes/pull/609/changes

Signed: seungrokj

Klaud-Cold · 2026-07-06T02:47:11Z

❌❌❌ REJECTED ❌❌❌

@seungrokj — blocking on Check 3: the linked recipe vllm-project/recipes#609 does not cover this PR's major server args. --linear-backend emulation (kernel selection; this PR's changelog credits it with ~+26% tput / -21% TPOT at 8k1k conc1 vs the stock native MXFP8 linear) appears in neither the recipe PR nor the published MiniMax-M3 recipe — and #609 updates the amd/MiniMax-M3-MXFP4 serve command, while this PR benchmarks MiniMaxAI/MiniMax-M3-MXFP8. Please add the MXFP8 MI355X command (including --linear-backend emulation) to the recipe and re-link. (This supersedes the earlier verification comment for this SHA, which audited a non-checklist comment rather than this review.)

✅ Check 0 (CODEOWNER): PASS — @seungrokj is a listed owner of configs/amd-master.yaml; the other changed paths are catch-all-owned.
✅ Check 1 (sweep on in-PR commit): PASS — green executed single-node 1k1k/8k1k and eval / check-runs on in-PR commit c696685 via run 28690689417.
✅ Check 2 (evals pass): PASS — GSM8K em_strict 0.952–0.960 (n_eff 1319) across non-MTP and MTP lanes, run on this PR's image vllm/vllm-openai-rocm:nightly-09663abde….
❌ Check 3 (recipe linked and complete): FAIL — major kernel-selection arg --linear-backend emulation is missing from the linked recipe, which also targets the MXFP4 variant rather than the MXFP8 one benchmarked here. (--moe-backend aiter / VLLM_ROCM_USE_AITER=1 are covered by Add official InferenceMAX disclaimer to benchmark summaries #609; --max-num-batched-tokens 32768 is InferenceX sweep tuning — informational only, not a blocker.)
✅ Check 4 (reuse command): PASS — /reuse-sweep-run posted by @seungrokj (COLLABORATOR).
✅ Check 5 (latest template): PASS — all current-template items are present and checked in the sign-off.
✅ Check 6 (upstream image / engine-first): PASS — upstream vllm/vllm-openai-rocm image on established MI355X; the framework is vLLM itself, so ordering does not apply.
✅ Check 7 (no architecture hacks): PASS — no --hf-overrides or FLOPs-reducing changes; the kernel-backend switches run the same computation.
✅ Check 8 (spec-decode chat template): PASS — the MTP benchmark drives the server with --use-chat-template (minimaxm3_fp8_mi355x_mtp.sh:224).

seungrokj · 2026-07-06T02:48:58Z

@functionstackx will update the recipe first..

seungrokj

As a PR reviewer and CODEOWNER, I have reviewed this and have:

Additional detail section:

insert any additional info here
m3 vllm recipe PR,
vllm-project/recipes#581

Signed: seungrokj

Klaud-Cold · 2026-07-06T05:07:47Z

✅✅✅ Verdict: PASS ✅✅✅

Head is still ed4a985 (pinned SHA). Conclusion changed vs the prior REJECTED verifications for this SHA: the new sign-off (review 4632991312) links vllm-project/recipes#581, which now carries the previously missing major args.

✅ Check 0 (CODEOWNER): PASS — @seungrokj is a listed owner of configs/amd-master.yaml; remaining paths are catch-all.
✅ Check 1 (sweep on in-PR commit): PASS — commit c696685 (in this PR) has green executed single-node 1k1k/8k1k / and eval / check-runs in run 28690689417; the two later commits only merge main and edit perf-changelog.
✅ Check 2 (evals): PASS — GSM8K em_strict 0.9538–0.9598 across base and MTP lanes (n_eff 1319), run on the PR's exact image vllm/vllm-openai-rocm:nightly-09663abde....
✅ Check 3 (recipe): PASS — recipes#581 matches the major args: model MiniMax-M3-MXFP8 on MI355X, --linear-backend emulation, VLLM_ROCM_USE_AITER=1 (the --moe-backend aiter flag is the explicit form of that path), --attention-backend TRITON_ATTN, block-size 128, TP4 ("MXFP8 variant runs from TP=4"), and --kv-cache-dtype fp8 in the guide. Informational only: --max-num-batched-tokens 32768, --no-enable-prefix-caching, --language-model-only are InferenceX sweep/harness tuning.
✅ Check 4 (reuse command): PASS — /reuse-sweep-run posted by @seungrokj (COLLABORATOR).
✅ Check 5 (latest template): PASS — all 10 items of the current docs/PR_REVIEW_CHECKLIST.md present and checked.
✅ Check 6 (upstream image / engine-first): PASS — both entries use upstream vllm/vllm-openai-rocm:nightly-09663abde...; framework is vLLM so engine-first ordering is satisfied.
✅ Check 7 (no architecture hacks): PASS — no --hf-overrides/model-config edits; --linear-backend emulation runs the MXFP8 linears via bf16 hipblasLT (same FLOPs, not fewer).
✅ Check 8 (spec-decode via chat template): PASS — the MTP script benchmarks through run_benchmark_serving ... --use-chat-template.

seungrokj · 2026-07-06T05:08:50Z

@functionstackx @Oseltamivir can you plz merge this ?

… for 8k1k high-conc On the sparse-PA nightly the dense-linear backend crossover reversed vs #2003: the native Triton MXFP8 GEMM now wins in the memory-bound low-concurrency regime, while --linear-backend emulation (bf16 hipBLASLT) still wins in the compute-bound high-concurrency regime. Measured on gfx950 MXFP8 + sparse PA: emulation ~+3-5% at 8k1k conc 64-512 but ~-3% at conc 1-8. Gate emulation to isl>=8192 && conc>=64; native everywhere else (all 1k1k, 8k1k conc<64). Overridable via LINEAR_BACKEND. 中文：在 sparse-PA nightly 上，稠密 linear 后端的性能拐点相比 #2003 发生反转：原生 Triton MXFP8 GEMM 现在在访存受限的低并发区间更快，而 --linear-backend emulation（bf16 hipBLASLT）仍在计算受限的高并发区间更快。gfx950 MXFP8 + sparse PA 实测：emulation 在 8k1k conc 64-512 约 +3-5%，在 conc 1-8 约 -3%。因此将 emulation 限定在 isl>=8192 且 conc>=64，其余（所有 1k1k、8k1k conc<64）使用原生路径。可通过 LINEAR_BACKEND 覆盖。 Co-authored-by: Cursor <cursoragent@cursor.com>

…high-conc AITER page-16 sparse PA is a long-context/high-batch optimization: it wins at 8k1k conc>=64 but adds overhead at short context (1k1k) and low concurrency. Gate both VLLM_ROCM_SHUFFLE_KV_CACHE_LAYOUT (sparse PA) and --linear-backend emulation behind a single MM3_HIGH_CONC_FASTPATH condition (isl>=8192 && conc>=64); everything else falls back to the #2003 path (non-shuffled Triton attention + native linear) so there is no regression outside the target regime. Net vs #2003: 8k1k tput/gpu conc128 +3.3%, conc256 +2.5%, conc512 +4.7% (peak), neutral elsewhere. Overridable via MM3_HIGH_CONC_FASTPATH / LINEAR_BACKEND. 中文：AITER page-16 稀疏分页注意力(sparse PA)是长上下文/大批量优化：在 8k1k conc>=64 有收益，但在短上下文(1k1k)和低并发下带来额外开销。将 VLLM_ROCM_SHUFFLE_KV_CACHE_LAYOUT（sparse PA）与 --linear-backend emulation 统一用 MM3_HIGH_CONC_FASTPATH 条件（isl>=8192 且 conc>=64）门控；其余情况回退到 #2003 路径（非 shuffle 的 Triton 注意力 + 原生 linear），从而在目标区间之外不产生回退。相对 #2003：8k1k tput/gpu conc128 +3.3%、conc256 +2.5%、conc512 +4.7%（峰值），其余基本持平。可通过 MM3_HIGH_CONC_FASTPATH / LINEAR_BACKEND 覆盖。 Co-authored-by: Cursor <cursoragent@cursor.com>

… concurrencies Switch VLLM_ROCM_QUICK_REDUCE_QUANTIZATION from INT6 to INT4 unconditionally (#2003 used INT6); accuracy is guarded by the 8k1k evals. The AITER quick all-reduce tuning knobs (CAST_BF16_TO_FP16=0, QUANTIZATION_MIN_SIZE_KB=256) remain in the 8k1k high-conc fast path. Aimed at lifting the low/mid-conc points (e.g. the 8k1k conc8 dip) via cheaper all-reduces. 中文：将 VLLM_ROCM_QUICK_REDUCE_QUANTIZATION 从 INT6 无条件改为 INT4（#2003 用的是 INT6），准确性由 8k1k evals 保障。AITER quick all-reduce 调优开关 (CAST_BF16_TO_FP16=0、QUANTIZATION_MIN_SIZE_KB=256) 仍保留在 8k1k 高并发快速路径中。目的是通过更廉价的 all-reduce 提升低/中并发点（例如 8k1k conc8 的回退）。 Co-authored-by: Cursor <cursoragent@cursor.com>

…-conc + native linear default / MiniMax-M3 MXFP8 MI355X vLLM：AITER 稀疏分页注意力门控至 8k1k 高并发 + 默认原生 linear (#2187) * [AMD] MiniMax-M3 MXFP8 MI355X vLLM: enable AITER sparse PA + drop emulation linear Bump the minimaxm3-fp8-mi355x-vllm image to nightly-9e57de7197f234f9d9187715d96e07e007048c0f, which carries the merged AITER page-16 sparse paged-attention path (vllm-project/vllm#47287). Enable it via VLLM_ROCM_SHUFFLE_KV_CACHE_LAYOUT=1 (with VLLM_ROCM_USE_AITER=1 and fp8 KV cache on TP4, num_kv_heads == 1 per rank) plus the recipe's quick all-reduce knobs VLLM_ROCM_QUICK_REDUCE_CAST_BF16_TO_FP16=0 and VLLM_ROCM_QUICK_REDUCE_QUANTIZATION_MIN_SIZE_KB=256. Drop --linear-backend emulation now that the Triton MXFP8 dense-linear GEMM is optimized. Deliberately do NOT add the #47269 --hf-overrides use_index_cache/index_topk_freq indexer skip: it reduces model-architecture FLOPs, disallowed by docs/PR_REVIEW_CHECKLIST.md. Sparse PA is a kernel-level speedup only. Verified locally on MI355X (gfx950) TP4 8k1k: conc1 10/10 and conc512 5120/5120 completed with zero failures (conc512 ~28.6k tok/s total, ~7.1k tok/s/GPU). 中文：将 minimaxm3-fp8-mi355x-vllm 镜像升级到 nightly-9e57de7197f234f9d9187715d96e07e007048c0f，该镜像已包含合并进主分支的 AITER page-16 稀疏分页注意力(sparse paged attention, vllm-project/vllm#47287)。通过 VLLM_ROCM_SHUFFLE_KV_CACHE_LAYOUT=1（配合 VLLM_ROCM_USE_AITER=1、fp8 KV cache 及 TP4，每个 rank num_kv_heads == 1）启用，并补上 recipe 的 quick all-reduce 调优开关 VLLM_ROCM_QUICK_REDUCE_CAST_BF16_TO_FP16=0 与 VLLM_ROCM_QUICK_REDUCE_QUANTIZATION_MIN_SIZE_KB=256。由于 Triton MXFP8 稠密 linear GEMM 已优化，移除 --linear-backend emulation。刻意不添加 #47269 的 --hf-overrides use_index_cache/index_topk_freq indexer 跳层开关：它会减少模型架构 FLOPs，违反 docs/PR_REVIEW_CHECKLIST.md。sparse PA 仅为 kernel 级加速。已在 MI355X(gfx950) TP4 8k1k 本地验证：conc1 10/10、conc512 5120/5120 全部成功（conc512 约 28.6k tok/s 总吞吐，约 7.1k tok/s/GPU）。 Co-authored-by: Cursor <cursoragent@cursor.com> * [AMD] MiniMax-M3 MXFP8 MI355X vLLM: concurrency-gate emulation linear for 8k1k high-conc On the sparse-PA nightly the dense-linear backend crossover reversed vs #2003: the native Triton MXFP8 GEMM now wins in the memory-bound low-concurrency regime, while --linear-backend emulation (bf16 hipBLASLT) still wins in the compute-bound high-concurrency regime. Measured on gfx950 MXFP8 + sparse PA: emulation ~+3-5% at 8k1k conc 64-512 but ~-3% at conc 1-8. Gate emulation to isl>=8192 && conc>=64; native everywhere else (all 1k1k, 8k1k conc<64). Overridable via LINEAR_BACKEND. 中文：在 sparse-PA nightly 上，稠密 linear 后端的性能拐点相比 #2003 发生反转：原生 Triton MXFP8 GEMM 现在在访存受限的低并发区间更快，而 --linear-backend emulation（bf16 hipBLASLT）仍在计算受限的高并发区间更快。gfx950 MXFP8 + sparse PA 实测：emulation 在 8k1k conc 64-512 约 +3-5%，在 conc 1-8 约 -3%。因此将 emulation 限定在 isl>=8192 且 conc>=64，其余（所有 1k1k、8k1k conc<64）使用原生路径。可通过 LINEAR_BACKEND 覆盖。 Co-authored-by: Cursor <cursoragent@cursor.com> * [AMD] MiniMax-M3 MXFP8 MI355X vLLM: gate sparse-PA fast path to 8k1k high-conc AITER page-16 sparse PA is a long-context/high-batch optimization: it wins at 8k1k conc>=64 but adds overhead at short context (1k1k) and low concurrency. Gate both VLLM_ROCM_SHUFFLE_KV_CACHE_LAYOUT (sparse PA) and --linear-backend emulation behind a single MM3_HIGH_CONC_FASTPATH condition (isl>=8192 && conc>=64); everything else falls back to the #2003 path (non-shuffled Triton attention + native linear) so there is no regression outside the target regime. Net vs #2003: 8k1k tput/gpu conc128 +3.3%, conc256 +2.5%, conc512 +4.7% (peak), neutral elsewhere. Overridable via MM3_HIGH_CONC_FASTPATH / LINEAR_BACKEND. 中文：AITER page-16 稀疏分页注意力(sparse PA)是长上下文/大批量优化：在 8k1k conc>=64 有收益，但在短上下文(1k1k)和低并发下带来额外开销。将 VLLM_ROCM_SHUFFLE_KV_CACHE_LAYOUT（sparse PA）与 --linear-backend emulation 统一用 MM3_HIGH_CONC_FASTPATH 条件（isl>=8192 且 conc>=64）门控；其余情况回退到 #2003 路径（非 shuffle 的 Triton 注意力 + 原生 linear），从而在目标区间之外不产生回退。相对 #2003：8k1k tput/gpu conc128 +3.3%、conc256 +2.5%、conc512 +4.7%（峰值），其余基本持平。可通过 MM3_HIGH_CONC_FASTPATH / LINEAR_BACKEND 覆盖。 Co-authored-by: Cursor <cursoragent@cursor.com> * chore: set perf-changelog pr-link to #2187 中文：将 perf-changelog 的 pr-link 设置为 #2187。 Co-authored-by: Cursor <cursoragent@cursor.com> * minor tuning * [AMD] MiniMax-M3 MXFP8 MI355X vLLM: use INT4 quick all-reduce for all concurrencies Switch VLLM_ROCM_QUICK_REDUCE_QUANTIZATION from INT6 to INT4 unconditionally (#2003 used INT6); accuracy is guarded by the 8k1k evals. The AITER quick all-reduce tuning knobs (CAST_BF16_TO_FP16=0, QUANTIZATION_MIN_SIZE_KB=256) remain in the 8k1k high-conc fast path. Aimed at lifting the low/mid-conc points (e.g. the 8k1k conc8 dip) via cheaper all-reduces. 中文：将 VLLM_ROCM_QUICK_REDUCE_QUANTIZATION 从 INT6 无条件改为 INT4（#2003 用的是 INT6），准确性由 8k1k evals 保障。AITER quick all-reduce 调优开关 (CAST_BF16_TO_FP16=0、QUANTIZATION_MIN_SIZE_KB=256) 仍保留在 8k1k 高并发快速路径中。目的是通过更廉价的 all-reduce 提升低/中并发点（例如 8k1k conc8 的回退）。 Co-authored-by: Cursor <cursoragent@cursor.com> --------- Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: seungrokj <144636725+seungrokj@users.noreply.github.com> Co-authored-by: Alec Ibarra <93070681+adibarra@users.noreply.github.com>

…weep / MiniMax-M3 FP4 MI355X vLLM STP：INT4 all-reduce + fp8 KV，TP4 扫描 (#2104) * [AMD] MiniMax-M3 FP4 MI355X vLLM STP: INT4 all-reduce + fp8 KV, TP4 sweep Drop --hf-overrides indexer-skipping (disallowed per PR_REVIEW_CHECKLIST). Keep INT4 quick-reduce env knobs and --kv-cache-dtype fp8; bump to latest nightly and TP4-only conc 1-512 search space (#2003). 中文：MiniMax-M3 FP4 MI355X vLLM STP 保留 INT4 all-reduce 与 fp8 KV，去掉不允许的 indexer --hf-overrides；镜像升级至最新 nightly，扫描范围收窄为 TP4 conc 1-512。 Co-authored-by: Cursor <cursoragent@cursor.com> * fix(changelog): set pr-link to #2104 中文：将 changelog pr-link 更新为 #2104。 Co-authored-by: Cursor <cursoragent@cursor.com> * Repair perf-changelog: append #2104 entry at end The minimaxm3-fp4-mi355x-vllm changelog entry was inserted mid-file by a prior merge, which clobbered the neighboring kimik2.5-fp4-mi355x-atom config-keys header and left the file non-append-only. Rebuild the file as origin/main byte-for-byte with the single #2104 entry appended at the end so the append-only perf-changelog gate passes. Co-authored-by: Cursor <cursoragent@cursor.com> --------- Co-authored-by: Fangzhou Ai <Fangzhou-Ai@users.noreply.github.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Chun Fang <chun.fang@amd.com>

github-project-automation Bot added this to InferenceMAX Board Jul 3, 2026

hongxiayang marked this pull request as ready for review July 3, 2026 20:13

hongxiayang requested a review from a team July 3, 2026 20:13

hongxiayang requested review from 1am9trash, billishyahao, chunfangamd, seungrokj and yctseng0211 as code owners July 3, 2026 20:13

claude Bot reviewed Jul 3, 2026

View reviewed changes

hongxiayang added the full-sweep-enabled label Jul 3, 2026

hongxiayang force-pushed the hy/mm3-mxfp8-configs branch from 17583d7 to e0974d5 Compare July 3, 2026 22:28

hongxiayang and others added 4 commits July 4, 2026 01:24

address feedback

8699b80

Signed-off-by: Hongxia Yang <hongxia.yang@amd.com>

hongxiayang force-pushed the hy/mm3-mxfp8-configs branch from e0974d5 to c696685 Compare July 4, 2026 01:25

functionstackx changed the title ~~[AMD] MiniMax-M3 MXFP8 MI355X vLLM: nightly + AITER-on TP4 + emulatin linear~~ [AMD] MiniMax-M3 MXFP8 MI355X vLLM: nightly + AITER-on TP4 + emulatin linear / MiniMax-M3 MXFP8 MI355X vLLM：升级 nightly + 启用 AITER TP4 + emulation linear Jul 4, 2026

functionstackx mentioned this pull request Jul 4, 2026

[Klaud Cold] CONTRIBUTING: warn against root-owned files on AMD runners / 贡献指南：警告勿在 AMD runner 工作区留下 root 所属文件 #2043

Merged

2 tasks

seungrokj added 2 commits July 6, 2026 11:25

Merge branch 'main' into hy/mm3-mxfp8-configs

87d024e

Update perf-changelog.yaml

ed4a985

seungrokj approved these changes Jul 6, 2026

View reviewed changes

seungrokj self-requested a review July 6, 2026 02:42

seungrokj approved these changes Jul 6, 2026

View reviewed changes

seungrokj self-requested a review July 6, 2026 05:02

seungrokj approved these changes Jul 6, 2026

View reviewed changes

adibarra approved these changes Jul 6, 2026

View reviewed changes

adibarra merged commit 17046f4 into main Jul 6, 2026
36 checks passed

github-project-automation Bot moved this to Done in InferenceMAX Board Jul 6, 2026

adibarra deleted the hy/mm3-mxfp8-configs branch July 6, 2026 16:36

hongxiayang mentioned this pull request Jul 10, 2026

[AMD] MiniMax-M3 MXFP8 MI355X vLLM disagg: TP4 worker-ratio + serve sync / MiniMax-M3 MXFP8 MI355X vLLM 分离式：TP4 工作单元配比 + serve 参数对齐 #2144

Merged

2 tasks

hongxiayang mentioned this pull request Jul 14, 2026

[AMD] MiniMax-M3 MXFP8 MI355X vLLM: gate AITER sparse PA to 8k1k high-conc + native linear default / MiniMax-M3 MXFP8 MI355X vLLM：AITER 稀疏分页注意力门控至 8k1k 高并发 + 默认原生 linear #2187

Merged

5 tasks

	pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/XXX
	pr-link: https://github.com/SemiAnalysisAI/InferenceX/pull/2003

Branch	Before	After
EP-on	`VLLM_ROCM_USE_AITER=1`	`VLLM_ROCM_USE_AITER=0`
TP-only	`VLLM_ROCM_USE_AITER=0`	`VLLM_ROCM_USE_AITER=1`

Uh oh!

Conversation

hongxiayang commented Jul 3, 2026 • edited by functionstackx Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

中文说明

Uh oh!

github-actions Bot commented Jul 3, 2026

Uh oh!

hongxiayang commented Jul 3, 2026

Uh oh!

claude Bot commented Jul 3, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Review complete

Uh oh!

claude Bot Jul 3, 2026

Choose a reason for hiding this comment

Uh oh!

claude Bot Jul 3, 2026

Choose a reason for hiding this comment

Uh oh!

claude Bot Jul 3, 2026

Choose a reason for hiding this comment

Uh oh!

github-actions Bot commented Jul 3, 2026

Uh oh!

github-actions Bot commented Jul 3, 2026

Uh oh!

github-actions Bot commented Jul 4, 2026

Uh oh!

hongxiayang commented Jul 4, 2026

Uh oh!

functionstackx commented Jul 4, 2026

Uh oh!

functionstackx commented Jul 4, 2026

Uh oh!

functionstackx commented Jul 4, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

seungrokj commented Jul 6, 2026

Uh oh!

seungrokj left a comment

Choose a reason for hiding this comment

Uh oh!

functionstackx commented Jul 6, 2026

Uh oh!

seungrokj left a comment

Choose a reason for hiding this comment

Additional detail section:

Uh oh!

Klaud-Cold commented Jul 6, 2026

❌❌❌ REJECTED ❌❌❌

Uh oh!

seungrokj left a comment

Choose a reason for hiding this comment

Additional detail section:

Uh oh!

Klaud-Cold commented Jul 6, 2026

❌❌❌ REJECTED ❌❌❌

Uh oh!

seungrokj commented Jul 6, 2026

Uh oh!

seungrokj left a comment

Choose a reason for hiding this comment

Additional detail section:

Uh oh!

Klaud-Cold commented Jul 6, 2026

✅✅✅ Verdict: PASS ✅✅✅

Uh oh!

seungrokj commented Jul 6, 2026

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

hongxiayang commented Jul 3, 2026 •

edited by functionstackx

Loading

claude Bot commented Jul 3, 2026 •

edited

Loading

functionstackx commented Jul 4, 2026 •

edited

Loading