feat(hipblaslt): add gfx1250 Aux, FP8 AmaxD and MXFP8 logic for TransformerEngine by mengzcai · Pull Request #13300 · ROCm/rocm-libraries

mengzcai · 2026-10-08T06:11:17Z

FP8, MXFP8, GELU_AUX support
JIRA ID : https://amd-hub.atlassian.net/browse/ROCM-32461

MXBF8 support
JIRA ID : https://amd-hub.atlassian.net/browse/AIHPBLAS-5158

Motivation

TransformerEngine core GEMM tests on gfx1250 fail with "no algorithm" because hipBLASLt has no matching kernel for:

bf16/fp32 GELU_AUX(_BIAS) (only fp16 Aux ships);
per-tensor FP8 -> FP8/BF8, because TE always requests AmaxD and no gfx1250 kernel outputs it;
MXFP8 with mixed e4m3/e5m2 inputs, or fp8 x fp8 -> fp16 (also used by DqGEMM).

gfx1250 also has no MX kernel with bf8 inputs, so the bf8 rows of matmul_fast_check_mx_fp8{,_large_k} (#13031) fail with "NO solution found" and were limited to gfx950 by #12983 (AIHPBLAS-5158).

Technical Details

New GridBased logic for gfx1250 and gfx1250-strict:

Aux: BF16 and FP32 GEMM + Bias + Act with Aux output (UseE), all 4 transposes. Uses the existing non-Aux BF16/FP32 solutions.
FP8 AmaxD: per-tensor {F8, B8, F8B8, B8F8} -> {F8, B8} with OutputAmaxD, all 4 transposes. Uses the existing ScaleCD solutions.
MXFP8 (TN): F8xB8 and B8xF8 -> BF16/F16/F32, F8xF8 -> F16, and B8xB8 -> F32. Uses the existing MX solution.

All new solutions build exactly the kernels of the solutions they are derived from. Logic was generated from Tensile configs with LibraryLogic.

Fixes:

PackData_BF16 (tensilelite): on HasBF16CVT archs, the BF16 Aux store converted its own destination register instead of the input values. This produced a garbage Aux output on full tiles with gwvw > 1. It also changes 2 of the existing gfx950 BBS Aux kernels.
Test client: for fp8 inputs with an fp8/bf8/f16 output, a bf16 bias was put in an f32 buffer while the descriptor declared bf16.

New gtests (gpu_arch: 1250):

matmul_bias_gelu_aux_bf16_fp32{,_large}
matmul_fp8_dst_fp8_amaxD_1250
matmul_mxfp8_mix_dst_f16_bf16_f32_1250

The bf8 rows of matmul_fast_check_mx_fp8{,_large_k} run on (950|1250) again.

Test Plan

hipblaslt-bench on TE's shapes (768x4096x768, 2304x4096x768, 256^3) and hipblaslt-test (new gtests + existing *matmul_real_1b_dst_f8_SCD*, *1250_smoke*, *matmul_fast_check_mx_fp8*), each on both library subtrees (gfx1250 and gfx1250-strict code objects)
Functional validation on the gfx1250 FFM and B0

Test Result

FFM, both subtrees (gfx1250 and gfx1250-strict):

Tensile: every new logic file passes validation (D, Aux and AmaxD outputs) — bf16/fp32 Aux, FP8 AmaxD, mixed MXFP8, BF8xBF8 MX.
hipblaslt-bench: all TE shapes pass — bf16/fp32 GELU_AUX_BIAS (4 transposes), FP8/BF8 output with AmaxD (all 8 input/output combinations, plus fp16 and bf16 bias), MXFP8 mixed and ->f16, DqGEMM 256^3. Every case selects the new kernel (previously 0 solutions).
hipblaslt-test: *matmul_bias_gelu_aux* 392/392, new FP8 gtests 440/440, regression set 4664/4664 — each on both subtrees.
hipblaslt-test: *matmul_fast_check_mx_fp8*-*placement* 12/12 on each subtree, including the 8 bf8 tests that failed with "NO solution found" before.

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

therock-pr-bot · 2026-10-08T06:15:27Z

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

therock-pr-bot · 2026-10-08T06:15:30Z

🎉 All checks passed! This PR is ready for review.

codecov · 2026-10-08T11:18:32Z

Codecov Report

✅ All modified and coverable lines are covered by tests.

❌ Your project check has failed because the head coverage (43.30%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage.

Additional details and impacted files

@@             Coverage Diff             @@
##           develop   #13300      +/-   ##
===========================================
- Coverage    74.65%   74.64%   -0.00%     
===========================================
  Files         2881     2881              
  Lines       479786   479786              
  Branches     71117    71117              
===========================================
- Hits        358139   358132       -7     
- Misses       97406    97410       +4     
- Partials     24241    24244       +3

Flag	Coverage Δ		*Carryforward flag
TensileLite-CPP	`49.69% <ø> (ø)`		Carriedforward from 52ddd47
TensileLite-Unit	`83.55% <ø> (ø)`		Carriedforward from 52ddd47
hipBLAS	`90.62% <ø> (ø)`		Carriedforward from 52ddd47
hipBLASLt	`43.30% <ø> (-0.03%)`	⬇️
hipCUB	`82.68% <ø> (ø)`		Carriedforward from 52ddd47
hipDNN	`86.80% <ø> (ø)`		Carriedforward from 52ddd47
hipFFT	`42.00% <ø> (ø)`		Carriedforward from 52ddd47
hipRAND	`76.12% <ø> (ø)`		Carriedforward from 52ddd47
hipSOLVER	`69.42% <ø> (ø)`		Carriedforward from 52ddd47
hipSPARSE	`86.99% <ø> (ø)`		Carriedforward from 52ddd47
rocBLAS	`68.73% <ø> (ø)`		Carriedforward from 52ddd47
rocFFT	`45.70% <ø> (ø)`		Carriedforward from 52ddd47
rocRAND	`57.59% <ø> (ø)`		Carriedforward from 52ddd47
rocSOLVER	`76.83% <ø> (ø)`		Carriedforward from 52ddd47
rocSPARSE	`77.05% <ø> (ø)`		Carriedforward from 52ddd47
rocThrust	`93.15% <ø> (ø)`		Carriedforward from 52ddd47

*This pull request uses carry forward flags. Click here to find out more.

Files with missing lines	Coverage Δ
...slt/tensilelite/tensilelite/Components/PackData.py	`84.41% <ø> (ø)`

... and 1 file with indirect coverage changes

🚀 New features to boost your workflow:

❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

Add GEMM+Bias+Activation with Aux output (UseE) for BF16 and FP32, all four transposes, for gfx1250 and gfx1250-strict. Kernels are the existing non-Aux BF16/FP32 GridBased solutions; sizes add the gfx950 Aux sizes and 768x4096x768. Add matching gfx1250 gtests. Also fix PackData_BF16 on HasBF16CVT archs: when tmpVgpr was passed (the E-store destination), it converted tmpVgpr instead of the real inputs, so the BF16 Aux output was garbage on full tiles with gwvw > 1.

For fp8/bf8 inputs with an fp8/bf8/f16 output, the test client mapped a bf16 bias to an f32 buffer while the matmul desc still declared bf16, so the GPU read the bias with the wrong type and validation failed.

Add per-tensor FP8/BF8 (incl. mixed) -> FP8/BF8 with OutputAmaxD, all four transposes, for gfx1250 and gfx1250-strict, from the existing ScaleCD solutions. Add matching gfx1250 gtests.

Add MXFP8 TN F8xB8 / B8xF8 -> BF16/F16/F32, F8xF8 -> F16 (also used by DqGEMM) and B8xB8 -> F32 for gfx1250 and gfx1250-strict, from the existing MX solution. Add matching gfx1250 gtests, and run the bf8 rows of matmul_fast_check_mx_fp8{,_large_k} on gfx1250 again (AIHPBLAS-5158).

assistant-librarian · 2026-10-16T13:05:10Z

TheRock Submodule Bump Activity

Newest first

2026-10-16 — Included in TheRock via ROCm/TheRock#9032.

#15040) JIRA ID : AIHPBLAS-5448 ## Motivation `matmul_fp8_dst_fp8_amaxD_1250` fails intermittently on gfx1250 HW CI: the GPU AmaxD is low by a fixed amount per shape (TT 768x4096x768 bf8 x f8: CPU 211540, GPU 207480). Root cause: in `amax_inter_wave_reduction`, each tree step has upper waves store to slot `w - Offset`, a barrier, then lower waves load slot `w`. Every step reuses slots from 0, and nothing orders step i's loads before step i+1's stores. With 4 waves, w1 can store max(w1, w3) to slot 0 in step 2 before w0 has read w2's value from it in step 1, so w2's partial max is lost. The code is old and shared by all archs; it surfaced after #13300 made gfx1250 ship AmaxD kernels.  ## Technical Details - Wave `w` stores to slot `w`; lower wave `w` loads slot `w + Offset`. Each wave stores once per reduction, so no slot is reused and no barrier is added. Slot 0 stays free for the broadcast; LDS reservation is unchanged. - Reject `OutputAmaxD` with a non-power-of-two `NumWaves`: the tree halves the wave count, so e.g. 3 waves would drop w2. Shipped AmaxD solutions use 1, 2 or 4 waves. - Tests: unit tests for the slot addressing (every power-of-two wave count a valid WorkGroup allows, wave32 and wave64) and the reject; the gfx1250 AmaxD tox YAML ran only 1-wave workgroups, so it now also runs 2 and 4 waves.  ## Test Plan On gfx1250 (MI450X B0), develop d1dd5a8 vs this branch: hipblaslt-test `*amaxD_1250*`, the tox YAML, and the new unit tests. To make the race deterministic, a debug-only build delays wave 0 with `s_sleep` before its inter-wave load.  ## Test Result | test | develop | this PR | |---|---|---| | `*amaxD_1250*` x200 | pass (race too rare in one process) | pass | | `*amaxD_1250*` + forced delay | 80/384 fail every repeat; same values as CI (211540 vs 207480) | pass (x20) | | tox `amax_gfx1250.yaml` + forced delay | 7 of 30 4-wave runs fail; 1/2 waves pass | pass | | tox `amax_gfx1250.yaml` | - | pass | | new unit tests | fail | pass | | TensileLite unit + characterization, pre-commit | - | pass | gfx942 shares the codegen and is covered by the wave64 unit test only.  ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. --------- Co-authored-by: mengzcai <mengzcai@amd.com>

mengzcai requested a review from a team as a code owner October 8, 2026 06:11

github-actions Bot added project: hipblaslt project: hipsparselt ci:hipsparselt-fast labels Oct 8, 2026

therock-pr-bot Bot added the Not ready to Review label Oct 8, 2026

assistant-librarian Bot added the organization: ROCm label Oct 8, 2026

mengzcai force-pushed the users/mengzcai/add_bf16_fp32_aux_logic_yamls branch from ed562dd to 9d41a93 Compare October 8, 2026 07:38

mengzcai force-pushed the users/mengzcai/add_bf16_fp32_aux_logic_yamls branch from 9d41a93 to 823fb4a Compare October 12, 2026 07:11

mengzcai requested a review from jichangjichang October 12, 2026 07:32

mengzcai force-pushed the users/mengzcai/add_bf16_fp32_aux_logic_yamls branch from 823fb4a to cc3b173 Compare October 12, 2026 09:32

therock-pr-bot Bot removed the Not ready to Review label Oct 12, 2026

mengzcai force-pushed the users/mengzcai/add_bf16_fp32_aux_logic_yamls branch 2 times, most recently from 43336ba to a8eeca9 Compare October 13, 2026 06:31

mengzcai changed the title ~~feat(hipblaslt): add gfx1250 BF16/FP32 Aux GridBased logic and fix BF16 Aux store~~ feat(hipblaslt): add gfx1250 Aux, FP8 AmaxD and MXFP8 logic for TransformerEngine Oct 13, 2026

mengzcai added 3 commits October 14, 2026 03:43

fix(hipblaslt): honor bf16 bias for fp8 inputs in matmul tests

9deb788

For fp8/bf8 inputs with an fp8/bf8/f16 output, the test client mapped a bf16 bias to an f32 buffer while the matmul desc still declared bf16, so the GPU read the bias with the wrong type and validation failed.

feat(hipblaslt): add gfx1250 FP8 AmaxD GridBased logic yamls

9eb0707

Add per-tensor FP8/BF8 (incl. mixed) -> FP8/BF8 with OutputAmaxD, all four transposes, for gfx1250 and gfx1250-strict, from the existing ScaleCD solutions. Add matching gfx1250 gtests.

mengzcai force-pushed the users/mengzcai/add_bf16_fp32_aux_logic_yamls branch from a8eeca9 to fe2b4ba Compare October 14, 2026 03:58

mengzcai force-pushed the users/mengzcai/add_bf16_fp32_aux_logic_yamls branch from fe2b4ba to 9efe159 Compare October 14, 2026 06:18

geotseng-amd approved these changes Oct 14, 2026

View reviewed changes

jichangjichang approved these changes Oct 15, 2026

View reviewed changes

jichangjichang merged commit 001c232 into ROCm:develop Oct 15, 2026
291 of 298 checks passed

mengzcai deleted the users/mengzcai/add_bf16_fp32_aux_logic_yamls branch October 15, 2026 08:49

mengzcai mentioned this pull request Nov 4, 2026

fix(tensilelite): AmaxD inter-wave reduction LDS write-after-read race #15040

Merged

1 task done

Navigation Menu

Repository navigation

feat(hipblaslt): add gfx1250 Aux, FP8 AmaxD and MXFP8 logic for TransformerEngine - #13300

feat(hipblaslt): add gfx1250 Aux, FP8 AmaxD and MXFP8 logic for TransformerEngine#13300
jichangjichang merged 4 commits into
ROCm:developfrom
mengzcai:users/mengzcai/add_bf16_fp32_aux_logic_yamls

mengzcai commented Oct 8, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Oct 8, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Oct 8, 2026 •

edited

Loading

Uh oh!

codecov Bot commented Oct 8, 2026 •

edited

Loading

Uh oh!

Uh oh!

assistant-librarian Bot commented Oct 16, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

mengzcai commented Oct 8, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Uh oh!

therock-pr-bot Bot commented Oct 8, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

therock-pr-bot Bot commented Oct 8, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

codecov Bot commented Oct 8, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Codecov Report

Uh oh!

Uh oh!

assistant-librarian Bot commented Oct 16, 2026

TheRock Submodule Bump Activity

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

mengzcai commented Oct 8, 2026 •

edited

Loading

therock-pr-bot Bot commented Oct 8, 2026 •

edited

Loading

therock-pr-bot Bot commented Oct 8, 2026 •

edited

Loading

codecov Bot commented Oct 8, 2026 •

edited

Loading