Repository navigation
fix(hipblaslt): take the e5m3 software path on gfx1250-strict - #13136
Merged
jichangjichang merged 1 commit intoOct 6, 2026
Merged
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Contributor
Author
|
All tests passed and have been verified on the v1 machine. |
boringmorning
approved these changes
Oct 6, 2026
## Problem
The client build for gfx1250-strict fails on ROCm nightlies from
10.2.0a20261004:
error: '__builtin_amdgcn_cvt_pk_fp8_f32_e5m3' needs target feature
fp8e5m3-insts
ROCm/llvm-project 7667f52dd687 (llvm-project #227861, "Disable e5m3
conversions on gfx1250-strict") moved FeatureFP8E5M3Insts out of the
gfx1250 common feature set, so gfx1250-strict (v0) no longer has the e5m3
conversion instructions. hipblaslt_e5m3.h and DataTypes_E5M3.hpp enabled
the hardware fast path for both __gfx1250__ and __gfx1250_strict__.
## Fix
Enable HIP_E5M3_CVT_FAST_PATH for __gfx1250__ only. gfx1250-strict takes
the existing software conversion, which is also correct on older nightlies
(up to 20261003), whose compiler still accepted the e5m3 builtins on strict
but where v0 has no such instructions. gfx1250 (v1) code is unchanged and
keeps the hardware v_cvt_pk_fp8_f32.
## Testing
- Compile of the real client TU (hipblaslt_init_device.cpp) for gfx1250
and gfx1250-strict: the original headers fail on strict for 20261004,
20261005 and 20261006; with this change both targets build on every
nightly from 20260930 to 20261006.
- Clean hipBLASLt and hipSPARSELt builds with -a "gfx1250;gfx1250-strict"
on nightlies 20260930 to 20261006: rc=0, library layout OK.
- On gfx1250 (v1), ROCm 10.2.0a20261005:
- hipblaslt-test all: 67569 passed, 0 failed, 34 skipped
- hipblaslt-test *quick*: 18043 passed
- FP4 bench: 345 solutions, no nan; mxf4 smoke passed
- hipsparselt-test: 165405 / 165405 passed
- TensileLite unit, arch-stepping and packaging tests passed
- tox Tests/common: 9642 passed, 1 failed (gsu_gfx1250, BF16 GSU; flaky,
passes on rerun, and the failing kernel differs between runs)
- pre-commit (hipBLASLt and repo root configs): passed
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #13136 +/- ##
===========================================
+ Coverage 72.12% 72.14% +0.02%
===========================================
Files 2861 2861
Lines 472089 471234 -855
Branches 69803 69773 -30
===========================================
- Hits 340449 339944 -505
+ Misses 107910 107578 -332
+ Partials 23730 23712 -18
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
geotseng-amd
force-pushed
the
users/geotseng/develop-e5m3
branch
from
October 6, 2026 07:13
8a1b00d to
79ab6b3
Compare
geotseng-amd
enabled auto-merge (squash)
October 6, 2026 07:14
jichangjichang
approved these changes
Oct 6, 2026
jichangjichang
disabled auto-merge
October 6, 2026 07:24
Member
|
Even though there was no update posted by the bot, the PR was included in a recent submodule bump and is part of TheRock. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Follow-up to #11777 ("hipblaslt: add support for gfx1250-strict build").
Motivation
#11777 made eight source guards that tested only
__gfx1250__accept__gfx1250_strict__too, because compiling for gfx1250-strict predefines only the latter. Two of those guards enable the hardware e5m3 conversion. That was valid with the compiler #11777 was tested on (10.2.0a20260930), but the compiler has since changed.The hipBLASLt client build for
gfx1250-strictfails on ROCm nightlies from 10.2.0a20261004:ROCm/llvm-project
7667f52dd687cherry-picks llvm/llvm-project#227861 ("Disable e5m3 conversions on gfx1250-strict", LCOMPILER-2845). It movesFeatureFP8E5M3Instsout of the common gfx1250 feature set, sogfx1250-strict(v0) no longer has the e5m3 conversion instructions.hipblaslt_e5m3.handDataTypes_E5M3.hppstill enable the hardware fast path for both__gfx1250__and__gfx1250_strict__, as #11777 left them.This PR makes the strict build work again on current nightlies and keeps it working on older ones.
Technical Details
HIP_E5M3_CVT_FAST_PATHis now enabled for__gfx1250__only, in both copies of the header:projects/hipblaslt/library/include/hipblaslt/hipblaslt_e5m3.hprojects/hipblaslt/tensilelite/include/Tensile/DataTypes_E5M3.hppv_cvt_pk_fp8_f32, and the condition evaluates the same for it as before.memory_gfx.h, the float4/bfloat6/float6 headers, mxDataGenerator) are unchanged. They still accept both names, and the client and libraries build with them on every nightly from 20260930 to 20261006.v_cvt_scale_pk8/pk16on strict in their block32 forms. hipBLASLt and hipSPARSELt do not use these instructions or the HIP FP4/FP6 headers, so they need no change here.HSA_DISABLE_GFX12_STRICTunset or1, the device reportsgfx1250. The HIP kernels compiled into libhipblaslt and the clients then load their gfx1250 builds, which keep the hardware e5m3 path. This change does not alter that. It covers only thegfx1250-strictbuilds.Test Plan
hipblaslt_init_device.cpp, flags fromcompile_commands.json) forgfx1250andgfx1250-strict, with the original and the new headers, on every nightly from 20260930 to 20261006.install.sh -c -a "gfx1250;gfx1250-strict"), on every nightly from 20260930 to 20261006, then verify the library layout.gfx1250-strict-testing.md:*quick*)Tests/commonTest Result
gfx1250-strictfor 20261004, 20261005 and 20261006.gfx1250/gfx1250v0/gfx1250-strict, layout OK.*quick*Tests/commonThe one tox failure is
gsu_gfx1250(BF16 GSU, size 512x512x1x32767, Relu). It is flaky and unrelated to e5m3. Of three reruns on an idle GPU, two passed and one failed, and the failing kernel differs between runs.The tests ran on develop
d116965bed9plus this change. The branch is rebased onto develop9db04770489. The 3 commits in between touch only rocALUTION and rocSPARSE.Submission Checklist