Skip to content

[ROCm][Perf] Fuse MXFP8 activation quant into the 1x32 dot_scaled linear for M <= 16 - #63915

Merged
aoshen02 merged 3 commits into
vllm-project:mainfrom
RuibinCheung:cursor/rocm-mxfp8-linear-fused-act-quant-45a8
Oct 29, 2026
Merged

aoshen02 merged 3 commits into
vllm-project:mainfrom
RuibinCheung:cursor/rocm-mxfp8-linear-fused-act-quant-45a8

Conversation

@RuibinCheung

Copy link
Copy Markdown
Contributor

Overview

On gfx950, RocmDotScaledMxfp8LinearKernel (1x32-scale MXFP8 dense layers: MiniMax-M3/M3.1, ModelOpt/INC MXFP8) runs a separate mxfp8_e4m3_quantize launch before every tl.dot_scaled GEMM. For bf16 activations with 1-16 rows, this PR quantizes the activation inside the GEMM instead, so each such linear is one launch and the output is bit-identical. Fused: K <= 1024 for M <= 16, and K = 2048 for M <= 4.

The fused GEMM quantizes only min(16, next_pow2(M)) activation rows per K step and repeats them to fill the 16-row MFMA tile; that keeps the per-N-tile quant cost proportional to the real rows and is what makes K = 2048 profitable.

Claims

  • Per-call speedup on MI355X: 1.14x-1.64x for K <= 1024 at M = 1-16, and 1.03x-1.27x for K = 2048 at M = 1-4 (1.15x-1.27x for MiniMax-M3.1 o_proj at TP4, 6144 x 2048). Every other shape and M keeps the existing path (1.00x).
  • Output is bitwise equal to the existing path (separate quant + _select_cfg tile), including where the fused tile has a different BLOCK_K/BLOCK_N.
  • MiniMax-M3.1-Flash TP4 (on [Perf][FlashAttention] FA4: pass no FP8 descale for a unit per-tensor scale #63232), AgentX concurrency 1: o_proj and shared-expert down_proj run fused with no separate quant launch. Over baseline, in two runs with the arm order swapped: P90 TPOT 1.011x in both, P90 interactivity 1.011-1.012x.
  • No change under VLLM_BATCH_INVARIANT, for fp16 activations, for M > 16, for K = 2048 at M > 4, or for K > 2048.

Validation

Microbenchmark: "separate" is mxfp8_e4m3_quantize + GEMM (current behaviour); "fused" is this PR. Speedup over separate, MI355X:

shape (N x K) M=1 M=2 M=3 M=4 M=5 M=8 M=16 M=32
6144 x 384 (M3.1 shared down, TP8) 1.64x 1.46x 1.45x 1.46x 1.45x 1.47x 1.45x 1.01x*
8192 x 512 1.48x 1.49x 1.49x 1.51x 1.44x 1.46x 1.45x 0.99x*
6144 x 768 (M3.1 shared down, TP4) 1.37x 1.29x 1.29x 1.29x 1.26x 1.26x 1.26x 1.00x*
6144 x 1024 (M3.1 o_proj, TP8) 1.52x 1.48x 1.39x 1.41x 1.34x 1.36x 1.30x 1.00x*
4096 x 1024 1.47x 1.37x 1.29x 1.30x 1.22x 1.25x 1.14x 1.01x*
6144 x 2048 (M3.1 o_proj, TP4) 1.27x 1.20x 1.16x 1.15x 1.00x* 1.00x* 1.00x* 1.00x*
4096 x 2048 1.17x 1.11x 1.05x 1.03x 1.00x* 1.00x* 1.00x* 1.00x*
2560 x 6144 (M3.1 qkv, TP4) 0.99x* 1.00x* 1.00x* 1.00x* 1.00x* 1.00x* 1.00x* 1.00x*
1536 x 6144 (M3.1 shared gate_up, TP4) 0.99x* 1.00x* 1.00x* 1.00x* 1.00x* 1.00x* 1.00x* 1.00x*

* not fused at this shape; the column times the fallback, i.e. the noise floor.

End to end: MiniMax-M3.1-Flash-Quark-MXFP4, TP4, MI355X, mono decode off, FULL_DECODE_ONLY CUDA graphs, on #63232 (8d0f05a0c, image vllm/vllm-openai-rocm:minimax-m31-flash) with this PR applied on top; baseline is the same image unchanged. AgentX trace replay (InferenceX), concurrency 1, 600 s per run, max-num-seqs 2; run 1 in the order baseline, this PR, run 2 reversed. This PR / baseline (> 1 is better; TPOT as baseline / this PR; P90 interactivity is InferenceX's intvty.p90 = 1 / P90 TPOT):

Run P90 TPOT P90 interactivity Total throughput / GPU
1 1.011x 1.012x 1.004x
2 1.011x 1.011x 1.020x

GSM8K (lm-eval, 5-shot, full test set, greedy), measured with K <= 1024 fused: baseline 96.51 / this PR 96.21 strict-match (stderr 0.51 / 0.53), within one stderr; vLLM is not batch invariant when requests are batched together, so greedy outputs vary between runs even with bit-identical kernels. The fused output is bitwise equal to the existing path for every tested shape, including K = 2048.

Details

  • Row trimming. BLOCK_M = 16 is the smallest MFMA tile, but at M < 16 only the first M rows are stored. The kernel takes a constexpr QUANT_ROWS = min(16, next_pow2(M)), loads and quantizes that many rows, and repeats them with broadcast_to + reshape. At BLOCK_K = 128 the repeat costs more than the saved quant, so those tiles (K = 384, 640, ...) keep quantizing 16 rows. The unfused path passes QUANT_ROWS = BLOCK_M and generates the same code as before.
  • Why only M <= 16, and K = 2048 only up to M = 4. The fused GEMM requantizes the activation once per N-tile. At small M the separate launch (plus the round trip of the fp8 activation and its scales) dominates; with more rows or larger K the redundant quant work wins (table above). The gate does not look at N; cutting K = 2048 at M = 4 keeps every measured N at or above 1.03x.
  • K becomes a tl.constexpr. Without it the fused loop is not unrolled/pipelined and the fused GEMM is ~0.8x even at K = 512. [Perf][FlashAttention] FA4: pass no FP8 descale for a unit per-tensor scale #63232 makes the same change.
  • Compiles. Triton already specializes M on == 1 and % 16 == 0; QUANT_ROWS adds up to three more variants per fused shape (6 instead of 3 for M = 1-16 at BLOCK_K >= 256; 5 at K = 2048, where M = 5-16 stays unfused). For MiniMax-M3.1 at TP4 that is 17 instead of 12 _mxfp8_linear_kernel compiles over M = 1-16 (15 instead of 12 at TP8). They are compiled in the eager warm-up before CUDA graph capture; vLLM's capture sizes include every M from 1 to 16, so decode never compiles at runtime, and graph capture time was unchanged in the end-to-end runs.
  • bf16 only. The in-kernel fp8 cast makes the compiler enable the FP16 overflow clamp in the MODE register for the whole kernel, so an fp16 output that overflows saturates at +-65504 where the unfused GEMM produces inf. bf16 output is unaffected and bit-identical; fp16 activations keep the separate quant.
  • Batch invariance. Fused and unfused use different tiles (and therefore K-reduction orders), so under VLLM_BATCH_INVARIANT ([Feature] Support batch invariance on ROCm #52231) every M stays on the unfused path.
  • Limitations. Only gfx950 (dot_scaled path). Shapes are tuned on MI355X.

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.

…ear for M <= 16

On gfx950, RocmDotScaledMxfp8LinearKernel quantized every activation with a
separate mxfp8_e4m3_quantize launch before the tl.dot_scaled GEMM. For bf16
activations with 1-16 rows, the GEMM now loads the bf16 tile and quantizes it
in registers with the same fp32 operations, so each linear is one launch and
the output is bit-identical.

Every N-tile requantizes its activation, so the in-kernel work grows with K
while the saved launch does not. With BLOCK_K >= 256 the kernel quantizes
only min(16, next_pow2(M)) rows and repeats them to fill the 16-row MFMA
tile (rows >= M are never stored); at BLOCK_K=128 repeating the rows costs
more than it saves. The fused path is taken for K <= 1024 at M <= 16
(1.14-1.64x on MI355X) and for K=2048 at M <= 4 (1.03-1.27x); larger M or K
keep the separate quant. K becomes a constexpr, which the fused loop needs
to be profitable.

fp16 is not fused, because the in-kernel fp8 cast enables the FP16 overflow
clamp for the whole kernel. VLLM_BATCH_INVARIANT keeps the unfused path.

Adds test_rocm_mxfp8_fused_act_quant_bitwise_eq_unfused: the fused path is
taken and its output is bitwise equal to mxfp8_e4m3_quantize plus the
_select_cfg tile, for every quantized-row count, with and without padding
rows, including K=2048 at M=1-4.

Signed-off-by: RuibinCheung <ruibzhan@amd.com>

@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.

@mergify mergify Bot added quantization rocm Related to AMD ROCm labels Oct 27, 2026
@github-project-automation github-project-automation Bot moved this to Todo in AMD Oct 27, 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.

🚀

@mergify

mergify Bot commented Oct 28, 2026

Copy link
Copy Markdown
Contributor

This pull request has merge conflicts that must be resolved before it can be
merged. Please rebase the PR, @RuibinCheung.

https://docs.github.com/en/pull-requests/collaborating-with-pull-requests/working-with-forks/syncing-a-fork

@mergify mergify Bot added the needs-rebase label Oct 28, 2026
@shen-shanshan shen-shanshan added verified Run pre-commit for new contributors without triggering other tests ready ONLY add when PR is ready to merge/full CI is needed labels Oct 28, 2026
@github-actions

Copy link
Copy Markdown

✅ @RuibinCheung, CI is now available for this PR.

  • /ci run starts upstream CI; /amd-ci run starts AMD CI only.
  • Run commands test your branch with every commit currently on its upstream target branch. If it is behind, the bot merges the target branch into it first (pull before pushing again). Append --allow-stale to a run command to test an outdated branch as is, 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.

@RuibinCheung

Copy link
Copy Markdown
Contributor Author

/ci run

@github-actions

Copy link
Copy Markdown

❌ This PR is 263 commits behind upstream main. The bot could not merge upstream main into this branch (API returned 422: merge conflict between base and head). No new CI build was started. Either merge or rebase onto the latest main yourself and rerun /ci run, or use /ci run --allow-stale to test this branch as is.

Signed-off-by: RuibinCheung <ruibzhan@amd.com>
@mergify mergify Bot removed the needs-rebase label Oct 28, 2026
@RuibinCheung

Copy link
Copy Markdown
Contributor Author

/ci run

@github-actions

Copy link
Copy Markdown

✅ Triggered Buildkite CI #97942 for commit d138cf7aacff.

ℹ️ This PR is 1 commit behind upstream main. Merged upstream main into this branch as d138cf7aacff before starting CI. Pull the branch before pushing more changes.

@benenzhu benenzhu 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.

LGTM! Improvements under low concurrency GEMMs.

@aoshen02
aoshen02 merged commit bcbcd02 into vllm-project:main Oct 29, 2026
104 checks passed
@github-project-automation github-project-automation Bot moved this from Todo to Done in AMD Oct 29, 2026
@RuibinCheung
RuibinCheung deleted the cursor/rocm-mxfp8-linear-fused-act-quant-45a8 branch October 29, 2026 02:47
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

quantization ready ONLY add when PR is ready to merge/full CI is needed rocm Related to AMD ROCm verified Run pre-commit for new contributors without triggering other tests

Projects

Status: Done

Development

Successfully merging this pull request may close these issues.

5 participants