[Config] [GEMM] [gfx950] Add MiniMax-M3 master BF16 GEMM table & tuned ViT linear projections by amd-bishwoadhikari · Pull Request #6235 · ROCm/aiter

amd-bishwoadhikari · 2026-10-07T21:57:03Z

Summary

This PR updates the original configuration from #5758 and expands coverage across TP=8, TP=4, and TP=2 for both language model projections and multimodal vision encoders.

Key Changes

Expanded Coverage Across Parallelism Degrees (112 → 300 Rows):
- Language Model (TP=8, TP=4, TP=2): Adds prefill shapes ($M \in [1024..66576]$) for QKV dense/sparse, out-proj, gate-up, down-proj, router gate, and LM head.
Elimination of hipBLASLt & Zero PyTorch Fallbacks:
- Retuned all 74 shapes using native, compiled backends (flydsl, asm, triton, opus) with GPU event synchronization and strict numerical accuracy (err_ratio <= 0.05).
- Distribution across all 300 rows: 254 FlyDSL (flydsl_hgemm_*), 28 Opus, 16 ASM, 2 Triton, 0 hipBLASLt, and 0 torch.
JIT Deduplication & 1:1 Companion Parity:
- Removed 38 primary key collisions (gfx, cu_num, M, N, K, bias, dtype, outdtype, scaleAB, bpreshuffle) shared with minimax_m3_eagle_bf16_tuned_gemm.csv to ensure clean global BF16 JIT auto-discovery in aiter/jit/core.py:506-548 without runtime exceptions.
- Synchronized companion minimax_m3_mxfp4_bf16_untuned_gemm.csv 1:1 to the 300 tuned shapes in order and dimensions.
Full Model Coverage Across Parallelism Degrees:
- Language Model Decode (TP=8, TP=4, TP=2): Comprehensive acceleration for core language projections (bias=False) across token decode batch sizes ($M \in [16..512]$), covering QKV dense, QKV sparse, out-proj, gate-up, down-proj, and router gate, plus selective anchor prefill shapes ($M \in [2048, 16384, 32768]$).
- 32-Layer Multimodal Vision Transformer (ViT): Complete coverage for image patch tokens ($M \in [256..66576]$) with native bias (bias=True) across ViT QKV, MLP, and Multimodal Projector dimensions: $(N=3840, K=1280)$, $(N=1280, K=1280)$, $(N=5120, K=1280)$, $(N=1280, K=5120)$, $(N=6144, K=1280)$, $(N=6144, K=24576)$, and $(N=6144, K=6144)$.

Reproduction / Validation Command:

python3 csrc/gemm_a16w16/gemm_a16w16_tune.py \
  --run_config aiter/configs/model_configs/minimax_m3_mxfp4_bf16_tuned_gemm.csv \
  -i aiter/configs/model_configs/minimax_m3_mxfp4_bf16_untuned_gemm.csv

Breaking changes

None.

Verification Checklist

Tested on AMD Instinct MI355X (gfx950).
Zero libtype: torch fallback rows.
Zero duplicate primary keys and zero sibling-table collisions.
Kernel naming compliance verified.
Looked over contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

Add gfx950 master BF16 GEMM tuned configuration table (5,449 kernel solutions) for MiniMax-M3 across text prefill, decode, multimodal composite token counts, and vision encoder linear projections. Covers ASM, Opus, and FlyDSL kernels with and without bias.

github-actions · 2026-10-07T21:58:05Z

🏷️ 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 6235 --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.

…ntuned companion - Replace 2,076 torch fallback rows with profiled FlyDSL/hipBLASLt winners - Eliminate 3,605 sibling-table primary key collisions - Synchronize untuned companion CSV to exact 338 rows (1:1 shape parity) - Enforce strict non-torch, non-colliding table across TP=8, TP=4, TP=2, and ViT Signed-off-by: Bishwo Adhikari <bishwo.adhikari@amd.com>

Copilot

🟡 Changes recommended

Several impossible timing outliers and benchmark shape/bias mismatches undermine the tuning evidence.

1 open finding

Invalid timing outlier may select the wrong kernel · New

What changed in this PR

Expands MiniMax-M3 MXFP4 BF16 GEMM coverage on gfx950 with synchronized tuned and untuned configurations.

Changes:

Adds 338 TP and ViT GEMM shapes.
Uses native FlyDSL, hipBLASLt, Opus, ASM, and Triton kernels.
Synchronizes tuned and untuned shape tables.

File	Description
`minimax_m3_mxfp4_bf16_tuned_gemm.csv`	Adds gfx950 kernel selections and measurements.
`minimax_m3_mxfp4_bf16_untuned_gemm.csv`	Mirrors the expanded tuning shape set.

🧠 Review effort: Balanced

Give feedback about Copilot approvals in this survey to enter a drawing for a $150 gift card.

Copilot

🟡 Changes recommended

The tuned table retains a PyTorch fallback and introduces primary-key collisions that cause merged-config generation to fail.

2 open findings

Remove duplicate shape keys from merged BF16 tables · New
Replace PyTorch fallback with a verified native GEMM kernel · New

1 resolved since last review

Invalid timing outlier may select the wrong kernel

🧠 Review effort: Balanced

- Re-tune 3 physical outlier shapes ((240,4608,6144), (112,4992,6144), (128,4992,6144)) on MI355X (gfx950) with explicit event synchronization (use_cuda_event=True). - Brings latencies to realistic 45-89 us range and TFLOPS within physical boundaries (<2614 TFLOPS, <8000 GB/s): * (240,4608,6144): 89.34 us, 152.11 TFLOPS (FlyDSL solidx=534, splitK=3) * (112,4992,6144): 46.49 us, 147.79 TFLOPS (Torch native) * (128,4992,6144): 45.65 us, 172.00 TFLOPS (ASM pf3_splitk, splitK=1)

…licate BF16 tables - Replace PyTorch fallback on shape (112, 4992, 6144) with verified native hipblaslt kernel (solidx 440108, 23.0212 us, 298.43 TFLOPS, 2772.93 GB/s, err_ratio 0.0), guaranteeing zero unaccelerated PyTorch fallback rows. - Remove 38 duplicate shape keys colliding with minimax_m3_eagle_bf16_tuned_gemm.csv to ensure clean global BF16 JIT config auto-discovery without RuntimeError. - Synchronize companion minimax_m3_mxfp4_bf16_untuned_gemm.csv 1:1, reducing both tables to 300 clean, non-colliding rows on gfx950 (187 FlyDSL, 72 hipBLASLt, 24 Opus, 16 ASM, 1 Triton).

Copilot

🟡 Changes recommended

One existing shape regresses by roughly 80%, and the committed row totals conflict with the verification claims.

2 open findings

ASM replacement regresses exact-shape winner latency by ~80% · New
Committed table row counts conflict with PR coverage claims · New

2 resolved since last review

Remove duplicate shape keys from merged BF16 tables
Replace PyTorch fallback with a verified native GEMM kernel

🧠 Review effort: Balanced

…nner - Re-tune shape (128,4992,6144) on MI355X (gfx950) with event synchronization across all candidate backends, replacing the 45.65 us ASM entry with the winning native hipblaslt kernel (solidx 440108, 23.4287 us, 335.13 TFLOPS, 2739.91 GB/s, err_ratio 0.0), improving latency by 1.95x and outperforming the prior 25.35 us baseline. - Update backend distribution across the 300 non-colliding rows: 187 FlyDSL, 73 hipBLASLt, 24 Opus, 15 ASM, 1 Triton, 0 torch fallback.

Copilot

🔵 Needs a closer look

The stated coverage and benchmark metadata conflict with the tables, and one configuration introduces a severe latency outlier.

1 open finding

Committed table row counts conflict with PR coverage claims

1 resolved since last review

ASM replacement regresses exact-shape winner latency by ~80%

🧠 Review effort: Balanced

Copilot

🔵 Needs a closer look

The documented backend distribution and highlighted hipBLASLt winners contradict the committed table.

1 open finding

Update stated FlyDSL and hipBLASLt row counts

🧠 Review effort: Balanced

Copilot

🟡 Changes recommended

Several latency sequences appear extrapolated, contradicting the stated per-shape event-synchronized benchmarking.

1 open finding

Replace extrapolated benchmark rows with independently timed measurements · New

1 resolved since last review

Update stated FlyDSL and hipBLASLt row counts

🧠 Review effort: Balanced

…y timed measurements - Benchmark all 44 previously extrapolated shapes with GPU event-synchronized tuner execution across native backends (flydsl, opus, asm, triton) on MI355X (gfx950). - Eliminate all identical TFLOPS sequences across lines 224-243 and 255-301, establishing measured per-shape winners (FlyDSL vs Opus). - Maintain 100% physical sanity (<= 2614 TFLOPS, <= 8000 GB/s) and numerical accuracy (err_ratio <= 0.05). - Final table distribution across 300 rows: 254 FlyDSL, 28 Opus, 16 ASM, 2 Triton, 0 hipBLASLt, 0 torch.

Copilot

🔵 Needs a closer look

The 300 hardware-specific configurations need MI355X validation, and their backend totals conflict with the PR description.

1 open finding

Backend distribution disagrees with PR description · New

1 resolved since last review

Replace extrapolated benchmark rows with independently timed measurements

🧠 Review effort: Balanced

Copilot

🔵 Needs a closer look

The 300 hardware-specific kernel selections and performance claims require final MI355X validation.

0 open findings

1 resolved since last review

Backend distribution disagrees with PR description

🧠 Review effort: Balanced

yzhou103 · 2026-10-12T03:24:44Z

please update you branch

Copilot

🔵 Needs a closer look

The 300 hardware-specific configurations require final MI355X correctness and performance verification.

0 open findings

🧠 Review effort: Balanced

Copilot

🔵 Needs a closer look

The 300 hardware-specific dispatch choices require final human review because no head-matched GPU validation report was available.

0 open findings

🧠 Review effort: Balanced

Copilot

🟡 Changes recommended

The tables do not contain the language-model prefill range claimed in the PR description.

1 open finding

Language-model prefill coverage does not match benchmark tables · New

🧠 Review effort: Balanced

…past buffer limits The gfx950 A16W16 epilogue computed C offsets as int32 `row * n + col`, so any GEMM with M * N >= 2^31 stored to a wrapped address. Tuned BF16 rows are looked up by padded M, and getPaddedM(gl=1) maps every M > 8192 with N > 4096 to the M=8192 row, so a row tuned at M=8192 serves arbitrarily large M. With the MiniMax-M3 ViT rows from ROCm#6235, vLLM's encoder-cache profile run (M=768000, N=6144, K=1280) hits this and faults with "Write access to a read-only page" on MI355X (reproduced standalone via tgemm.mm at M>=360000). - Compute the non-split-K C offsets in int64 in both the full-tile and half-tile-interleaved epilogues. - Reject shapes whose A or B exceeds the 2 GiB buffer-resource range in make_gemm_a16w16_param_and_validate (also prunes them from tuning). - Skip a FlyDSL tuned row in get_GEMM_A16W16_tuned_config when the kernel cannot run the actual shape, falling back like an unrecognized kernel. - Add the large-M MiniMax-M3 ViT shape and a shape-limit check to op_tests/test_flydsl_hgemm.py. Co-Authored-By: Claude <noreply@anthropic.com>

amd-bishwoadhikari added 2 commits October 8, 2026 12:19

Merge branch 'main' into config/minimax-m3-mxfp4-bf16-gemm

d63438d

amd-bishwoadhikari force-pushed the config/minimax-m3-mxfp4-bf16-gemm branch from da33891 to 18e64e3 Compare October 8, 2026 13:57

Merge branch 'main' into config/minimax-m3-mxfp4-bf16-gemm

108e3f0

amd-bishwoadhikari marked this pull request as ready for review October 8, 2026 15:14

amd-bishwoadhikari requested review from a team and a balanced review from Copilot October 8, 2026 15:14

github-actions Bot added the Config label Oct 8, 2026

Copilot started reviewing on behalf of amd-bishwoadhikari October 8, 2026 15:15

View session

Copilot AI reviewed Oct 8, 2026

View reviewed changes

Comment thread aiter/configs/model_configs/minimax_m3_mxfp4_bf16_tuned_gemm.csv Resolved

Copilot AI balanced review requested due to automatic review settings October 8, 2026 17:00

Copilot started reviewing on behalf of amd-bishwoadhikari October 8, 2026 17:01

View session

Copilot AI reviewed Oct 8, 2026

View reviewed changes

Comment thread aiter/configs/model_configs/minimax_m3_mxfp4_bf16_tuned_gemm.csv Resolved

Comment thread aiter/configs/model_configs/minimax_m3_mxfp4_bf16_tuned_gemm.csv Resolved

amd-bishwoadhikari force-pushed the config/minimax-m3-mxfp4-bf16-gemm branch from 6da9c17 to 3d68ce4 Compare October 8, 2026 21:30

Copilot AI balanced review requested due to automatic review settings October 8, 2026 22:04

Copilot started reviewing on behalf of amd-bishwoadhikari October 8, 2026 22:05

View session

Copilot AI reviewed Oct 8, 2026

View reviewed changes

Comment thread aiter/configs/model_configs/minimax_m3_mxfp4_bf16_tuned_gemm.csv Resolved

Comment thread aiter/configs/model_configs/minimax_m3_mxfp4_bf16_tuned_gemm.csv Resolved

Copilot AI balanced review requested due to automatic review settings October 8, 2026 23:50

Copilot started reviewing on behalf of amd-bishwoadhikari October 8, 2026 23:51

View session

Copilot AI reviewed Oct 8, 2026

View reviewed changes

zufayu requested a review from yzhou103 October 9, 2026 01:56

zufayu assigned yzhou103 Oct 9, 2026

Merge branch 'main' into config/minimax-m3-mxfp4-bf16-gemm

d0499ba

Copilot AI balanced review requested due to automatic review settings October 9, 2026 07:14

Copilot started reviewing on behalf of amd-bishwoadhikari October 9, 2026 07:15

View session

Copilot AI reviewed Oct 9, 2026

View reviewed changes

Merge branch 'main' into config/minimax-m3-mxfp4-bf16-gemm

ac0549f

Copilot AI balanced review requested due to automatic review settings October 9, 2026 12:23

Copilot started reviewing on behalf of amd-bishwoadhikari October 9, 2026 12:24

View session

Copilot AI reviewed Oct 9, 2026

View reviewed changes

Comment thread aiter/configs/model_configs/minimax_m3_mxfp4_bf16_tuned_gemm.csv Resolved

Copilot AI balanced review requested due to automatic review settings October 9, 2026 12:43

Copilot started reviewing on behalf of amd-bishwoadhikari October 9, 2026 12:44

View session

Copilot AI reviewed Oct 9, 2026

View reviewed changes

Comment thread aiter/configs/model_configs/minimax_m3_mxfp4_bf16_tuned_gemm.csv Resolved

Merge branch 'main' into config/minimax-m3-mxfp4-bf16-gemm

4db9fe7

Copilot AI balanced review requested due to automatic review settings October 9, 2026 12:52

Copilot started reviewing on behalf of amd-bishwoadhikari October 9, 2026 12:53

View session

Copilot AI reviewed Oct 9, 2026

View reviewed changes

Merge branch 'main' into config/minimax-m3-mxfp4-bf16-gemm

6fe5786

Copilot AI balanced review requested due to automatic review settings October 12, 2026 06:14

Copilot started reviewing on behalf of amd-bishwoadhikari October 12, 2026 06:14

View session

Copilot AI reviewed Oct 12, 2026

View reviewed changes

Merge branch 'main' into config/minimax-m3-mxfp4-bf16-gemm

6c92165

Copilot AI balanced review requested due to automatic review settings October 12, 2026 14:50

Copilot started reviewing on behalf of amd-bishwoadhikari October 12, 2026 14:51

View session

Copilot AI reviewed Oct 12, 2026

View reviewed changes

Merge branch 'main' into config/minimax-m3-mxfp4-bf16-gemm

326c230

Copilot AI balanced review requested due to automatic review settings October 12, 2026 20:29

Copilot started reviewing on behalf of amd-bishwoadhikari October 12, 2026 20:30

View session

Copilot AI reviewed Oct 12, 2026

View reviewed changes

Comment thread aiter/configs/model_configs/minimax_m3_mxfp4_bf16_tuned_gemm.csv Resolved

yzhou103 approved these changes Oct 14, 2026

View reviewed changes

yzhou103 merged commit 614ecd5 into ROCm:main Oct 14, 2026

josh-schell-amd mentioned this pull request Oct 23, 2026

[FlyDSL] [Bugfix] gfx950 A16W16: fix int32 C-offset overflow #6853

Open

4 of 10 tasks

Conversation

amd-bishwoadhikari commented Oct 7, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Key Changes

Breaking changes

Verification Checklist

Uh oh!

github-actions Bot commented Oct 7, 2026

🏷️ CI Guide

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🔵 Needs a closer look

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🔵 Needs a closer look

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🔵 Needs a closer look

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🔵 Needs a closer look

Uh oh!

yzhou103 commented Oct 12, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🔵 Needs a closer look

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🔵 Needs a closer look

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

amd-bishwoadhikari commented Oct 7, 2026 •

edited

Loading