Repository navigation
[ROCm] Bump AITER to v0.1.26 - #62486
Conversation
![claude[bot]](https://avatars.githubusercontent.com/in/1236702?s=60&v=4){kind=small}
|
✅ @josh-schell-amd, CI is now available for this PR.
|
|
/ci run |
|
✅ Triggered Buildkite CI #95930 for commit ℹ️ This PR is 5 commits behind upstream |
|
/ci run all --allow-stale |
|
✅ CI is already running for this commit: https://buildkite.com/vllm/ci/builds/95930 |
|
/ci cancel |
|
✅ Requested cancellation of 1 CI build for |
|
/ci run all |
|
✅ Triggered Buildkite CI #95937 for commit |
Signed-off-by: josh schell <josh.schell@amd.com>
Signed-off-by: josh schell <josh.schell@amd.com>
…curacy test The test runs rocm_aiter_fused_experts with SWIGLUOAI and no bias but scored it against a silu(gate) * up reference. It only passed because AITER <= v0.1.25 had no CK-Tile instance for SwiGLU without bias and dispatched to a SiLU kernel instead. AITER v0.1.26 (ROCm/aiter#6478) adds the missing instance, so the kernel now computes gpt-oss SwiGLU and the silu reference no longer matches. Co-Authored-By: Claude <noreply@anthropic.com> Signed-off-by: josh schell <josh.schell@amd.com>
9ac1ad7 to
4784d8e
Compare
|
/ci run all |
|
✅ Triggered Buildkite CI #96064 for commit |
|
/ci retry |
|
✅ Triggered Buildkite CI #96201 for commit |
|
/amd-ci run all |
|
✅ Triggered Buildkite AMD CI #15056 for commit |
josh-schell-amd
changed the title
|
/ci retry |
|
✅ Triggered Buildkite CI #96256 for commit |
|
/ci retry |
|
✅ Triggered Buildkite CI #96274 for commit |
micah-wil
left a comment
micah-wil
left a comment
There was a problem hiding this comment.
LGTM. Failing CUDA test is unrelated and present on main
Conflict: docker/Dockerfile.rocm_72_base (modify/delete). Main's vllm-project#62486 bumps AITER to v0.1.26 in both base Dockerfiles; this branch removes the ROCm 7.2 stack, so the file stays deleted. The same bump in the ROCm 10 docker/Dockerfile.rocm_base, and its test change, merged cleanly. Co-Authored-By: Claude <noreply@anthropic.com> Signed-off-by: Andreas Karatzas <AndreasKaratzas@users.noreply.github.com>
Signed-off-by: josh schell <josh.schell@amd.com> Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: Micah Williamson <micah.williamson@amd.com>
Overview
Bump the ROCm base images to AITER v0.1.26, and fix a test bug that the bump exposed.
With only the version bump,
test_aiter_fused_moe_mi350_mxfp4_w4a16_accuracyfailed in this PR's CI (build 95937, step "(MI355 DPX) FP4 and INT4 MoE Kernels"). This is not a regression in v0.1.26. The test has always compared AITER's SwiGLU-OAI MoE against a SiLU reference. It only passed because of an AITER v0.1.25 dispatch bug that ran SiLU instead of SwiGLU. ROCm/aiter#6478 fixes that bug in v0.1.26, so the kernel's output is now correct and the wrong reference fails.Claims
AITER_BRANCHmoves fromv0.1.25tov0.1.26indocker/Dockerfile.rocm_baseanddocker/Dockerfile.rocm_72_base.ref_moe_forwardintests/kernels/moe/test_rocm_aiter_moe.pygains a"swigluoai"activation (alpha 1.702, clamp 7,up + 1, the same formula as CK-Tile'sck_tile::moe::Swiglu). The MI350 MXFP4 W4A16 accuracy test now uses it, so the test checks the activation it actually requests.Validation
All runs below used an MI355X (gfx950) with AITER v0.1.26 built from the tag.
pytest -v -s tests/kernels/moe/test_rocm_aiter_moe.py -k "mi350_mxfp4 or quant_scheme_support_matches_gfx950"pytest -v -s tests/kernels/moe/test_ocp_mx_moe.py::test_rocm_mxfp4_moe_oracle -k AITER_MXFP4pytest -v -s tests/kernels/moe/test_flydsl_moe.pyThe "original test" result for the first command reproduces the CI failure in build 95937, which reported 20.97% pass and mean_abs 0.412 on an MI350X.
To show which activation the kernel computes, I ran the test's exact case (32 tokens, hidden 512, inter 1024, 4 experts, top-2, seed 11) on both AITER versions. Each run used no bias (as the test does) and a zero bias. A zero bias leaves the math unchanged but selects AITER's with-bias CK-Tile kernel. Each cell is the pass rate at atol = rtol = 0.1:
On v0.1.25, a SwiGLU request without bias produced SiLU output. Every other combination, including v0.1.25 with a bias, produced SwiGLU-OAI.
Model evals. These are from this PR's CI (build 95937, commit
9707fa4, AITER v0.1.26 before the test fix; the test fix does not affect evals). Every eval step passed. A model passes when its measured score is at least the expected score minus the eval's tolerance.GPQA,
:amd: (MI355) GPQA Eval (GPT-OSS), 4/4 passed:GSM8K, `:amd: (MI355 DPX) LM Eval Small Models` (10 passed, 2 skipped) and `:amd: (MI300) LM Eval Small Models` (6 passed, 2 skipped)
Qwen3-30B-A3B-MXFP4A16 and gemma-4-E4B-it-qat-mobile-ct are skipped on ROCm on both GPUs. The last four configs are only in the MI355 list.
Details
Root cause. CK-Tile's
moe_dispatchselects a kernel by(activation, has_bias, split_k), andhas_biascomes from whether a bias tensor is passed. In AITER ≤ v0.1.25,gen_instances.pybuilt SwiGLU kernels only withhas_bias=true. A SwiGLU call with no bias therefore reachedmoe_gemm1_heuristic_dispatcher<..., Swiglu, false, ...>, which had no specialization. Its primary template isstatic MoeKernel dispatch(...) {}, an empty body with no return, which is undefined behavior. In practice it ran a SiLU kernel. The test callsrocm_aiter_fused_expertswithMoEActivation.SWIGLUOAIand no bias, so it was scored against that SiLU result and matched. ROCm/aiter#6478, which is in v0.1.26, generates the missing no-bias SwiGLU kernel and makes the unspecialized dispatcher throw. The kernel now computes gpt-oss SwiGLU, so the test's SiLU reference fails.Impact on real models. gpt-oss checkpoints have expert biases, so they take the with-bias kernel, which computed SwiGLU on v0.1.25 as well (zero-bias row above). Any caller that used SwiGLU without bias on AITER ≤ v0.1.25 was silently getting SiLU. This PR fixes that path for them.
Why the test fix is in this PR. The corrected reference fails on v0.1.25 (no-bias row above), so it has to land together with the bump.
Duplicate check. No open vLLM PR references this test or AITER v0.1.26 (
gh pr list --search "mi350_mxfp4_w4a16"and--search "aiter v0.1.26"). On the AITER side, #6478 and #6620 are both already in v0.1.26, and neither changes this test's no-bias path.Pull Request Checklist
I used vLLM's
/pr-checklistskill. (Mandatory for agents, optional for humans).AI assistance was used during the creation of this PR.
Design Fit: Minimizes impact on core components, reuses existing functionality, and justifies added complexity.
Testing and Validation: Validates the change and ensures any added tests are meaningful and reliable, with CI coverage or documented CI resource constraints and validation performed outside CI.
Code Quality and Style: Keeps code and comments clear and concise, and updates relevant documentation and examples.
Pull Request Contents: Includes a brief summary and relevant links, supports claims with evidence, explains root causes and implementation trade-offs, and follows the contributing guide.
AI assistance (Claude Code) was used to find the root cause and draft the test fix and this description. I reviewed every changed line and ran the tests above on MI355X hardware.
BEFORE SUBMITTING, PLEASE READ https://docs.vllm.ai/en/latest/contributing (anything written below this line will be removed by GitHub Actions)