Repository navigation
[AMDGPU] Validate scale_sel in v_cvt_scale_* - #227475
Merged
Merged
Conversation
Contributor
Author
This stack of pull requests is managed by sgh. |
rampitec
marked this pull request as ready for review
September 29, 2026 21:16
llvmorg-github-actions
Bot
added
backend:AMDGPU
clang:frontend
|
@llvm/pr-subscribers-llvm-ir @llvm/pr-subscribers-backend-amdgpu Author: Stanislav Mekhanoshin (rampitec) ChangesThese instructions can be block16 or block32 depending on the target Re-enable the rest of the instructions in the strict mode but validate Full diff: https://github.com/llvm/llvm-project/pull/227475.diff 9 Files Affected:
The server is unavailable at this time. Please wait a few minutes before you try again. |
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
rampitec
force-pushed
the
users/rampitec/block16-validate-scale
branch
from
September 29, 2026 21:19
917a745 to
d8a9087
Compare
Contributor
Author
|
@scchan do you mind to double check actual OPSEL bit ranges used here? |
scchan
suggested changes
Sep 30, 2026
rampitec
force-pushed
the
users/rampitec/block16-cvt
branch
from
September 30, 2026 20:05
48bad19 to
32fdcc5
Compare
rampitec
force-pushed
the
users/rampitec/block16-validate-scale
branch
from
September 30, 2026 20:05
d8a9087 to
58e3ec4
Compare
scchan
approved these changes
Sep 30, 2026
These instructions can be block16 or block32 depending on the target and scale_sel bits. Block16 is not supported in strict mode. Re-enable the rest of the instructions in the strict mode but validate the scale selector.
rampitec
force-pushed
the
users/rampitec/block16-validate-scale
branch
from
October 1, 2026 19:08
58e3ec4 to
b7cf639
Compare
rampitec
enabled auto-merge (squash)
October 1, 2026 19:15
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/94/builds/21627 Here is the relevant piece of the build log for the reference |
Expertcoderz
pushed a commit
to Expertcoderz/llvm-project
that referenced
this pull request
Oct 3, 2026
These instructions can be block16 or block32 depending on the target and scale_sel bits. Block16 is not supported in strict mode. Re-enable the rest of the instructions in the strict mode but validate the scale selector.
sunhmy-amd
added a commit
to ROCm/rocm-systems
that referenced
this pull request
Oct 4, 2026
LCOMPILER-2841 established that all six v_cvt_scale_pk16_*_{fp6,bf6}
opcodes execute on MI450-A0. What A0 lacks is the block16 half of the
scale_sel range, not the opcodes themselves.
llvm/llvm-project#227475 moves them off block16-cvt-scale-insts onto
gfx1250-insts, which gfx1250-strict has, and validates scale_sel instead:
0-3 on strict, 0-7 with block16.
Flip HIP_ENABLE_GFX1250_BLOCK16_BUILTINS to cover both targets so the six
pk16 sites use hardware on strict instead of the software fallback. All 28
pk16 call sites pass scale_sel 0, which is block32 and inside the valid
range on both targets, so the half of the range strict lacks is never
requested.
The macro guards nothing but those six opcodes: all 16 use sites across
amd_hip_ocp_fp.hpp and amd_hip_fp6.h reduce to cvt_scale_pk16_* calls,
with no block16-specific logic. Its name is now misleading, since it does
not gate block16 mode; renaming it to HIP_ENABLE_GFX1250_PK16_SCALE_BUILTINS
to mirror PK8_SCALE is left to a follow-up to keep this diff reviewable.
Depends on a compiler carrying #227475, which is itself stacked on
#227426. Until then gfx1250-strict fails to build with "needs target
feature block16-cvt-scale-insts", so the dependency is enforced at compile
time rather than degrading silently.
Both macro copies are updated. They cannot be merged into one:
amd_hip_ocp_fp.hpp must stay self-contained because hipRTC compiles it
without amd_hip_mx_common.h, which redefines hipRoundMode under RTC.
Verified by attempting the deduplication and observing hipRTC fail.
NOT YET MEASURED on hardware, unlike the pk8 flip in the previous commit.
#227475 is open, based on a branch rather than main, and its scale_sel
validation applies MaxSel=8 to all nine pk8 opcodes where LCOMPILER-2841
expects 0-3 for the three fp4 variants, so its scope may still change.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
sunhmy-amd
added a commit
to ROCm/rocm-systems
that referenced
this pull request
Oct 5, 2026
LCOMPILER-2841 established that all six v_cvt_scale_pk16_*_{fp6,bf6}
opcodes execute on MI450-A0. What A0 lacks is the block16 half of the
scale_sel range, not the opcodes themselves.
llvm/llvm-project#227475 moves them off block16-cvt-scale-insts onto
gfx1250-insts, which gfx1250-strict has, and validates scale_sel instead:
0-3 on strict, 0-7 with block16.
Flip HIP_ENABLE_GFX1250_BLOCK16_BUILTINS to cover both targets so the six
pk16 sites use hardware on strict instead of the software fallback. All 28
pk16 call sites pass scale_sel 0, which is block32 and inside the valid
range on both targets, so the half of the range strict lacks is never
requested.
The macro guards nothing but those six opcodes: all 16 use sites across
amd_hip_ocp_fp.hpp and amd_hip_fp6.h reduce to cvt_scale_pk16_* calls,
with no block16-specific logic. Its name is now misleading, since it does
not gate block16 mode; renaming it to HIP_ENABLE_GFX1250_PK16_SCALE_BUILTINS
to mirror PK8_SCALE is left to a follow-up to keep this diff reviewable.
Depends on a compiler carrying #227475, which is itself stacked on
#227426. Until then gfx1250-strict fails to build with "needs target
feature block16-cvt-scale-insts", so the dependency is enforced at compile
time rather than degrading silently.
Both macro copies are updated. They cannot be merged into one:
amd_hip_ocp_fp.hpp must stay self-contained because hipRTC compiles it
without amd_hip_mx_common.h, which redefines hipRoundMode under RTC.
Verified by attempting the deduplication and observing hipRTC fail.
NOT YET MEASURED on hardware, unlike the pk8 flip in the previous commit.
#227475 is open, based on a branch rather than main, and its scale_sel
validation applies MaxSel=8 to all nine pk8 opcodes where LCOMPILER-2841
expects 0-3 for the three fp4 variants, so its scope may still change.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
1 task
jichangjichang
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 6, 2026
Follow-up to [#11777](#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-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`](ROCm/llvm-project@7667f52dd687) cherry-picks [llvm/llvm-project#227861](llvm/llvm-project#227861) ("Disable e5m3 conversions on gfx1250-strict", LCOMPILER-2845). It moves `FeatureFP8E5M3Insts` out of the common gfx1250 feature set, so `gfx1250-strict` (v0) no longer has the e5m3 conversion instructions. `hipblaslt_e5m3.h` and `DataTypes_E5M3.hpp` still 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_PATH` is now enabled for `__gfx1250__` only, in both copies of the header: - `projects/hipblaslt/library/include/hipblaslt/hipblaslt_e5m3.h` - `projects/hipblaslt/tensilelite/include/Tensile/DataTypes_E5M3.hpp` ```diff -#if (defined(__gfx1250__) || defined(__gfx1250_strict__)) && __HIP_DEVICE_COMPILE__ +// gfx1250-strict (v0) has no e5m3 conversion instructions; it takes the software path. +#if defined(__gfx1250__) && __HIP_DEVICE_COMPILE__ #define HIP_E5M3_CVT_FAST_PATH 1 ``` - **gfx1250-strict:** takes the existing software conversion. This is also correct on nightlies up to 20261003. Their compiler still accepted the e5m3 builtins on strict, but v0 has no such instructions. - **gfx1250 (v1):** unchanged. It keeps the hardware `v_cvt_pk_fp8_f32`, and the condition evaluates the same for it as before. - **The remaining guards from #11777** (`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. - **Tensile kernels:** not affected. TensileLite emits no e5m3 conversion, and the strict library tree assembles on 20261004 to 20261006. - **Not covered:** the other changes in the same cherry-pick ([#227426](llvm/llvm-project#227426), [#227475](llvm/llvm-project#227475)) re-enable `v_cvt_scale_pk8/pk16` on 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. - **Still not addressed, as listed in #11777:** on a v0 part with `HSA_DISABLE_GFX12_STRICT` unset or `1`, the device reports `gfx1250`. 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 the `gfx1250-strict` builds. ## Test Plan - Compile the real client TU (`hipblaslt_init_device.cpp`, flags from `compile_commands.json`) for `gfx1250` and `gfx1250-strict`, with the original and the new headers, on every nightly from 20260930 to 20261006. - Clean builds of hipBLASLt and hipSPARSELt, clients included (`install.sh -c -a "gfx1250;gfx1250-strict"`), on every nightly from 20260930 to 20261006, then verify the library layout. - Full test pass on gfx1250 (v1), ROCm 10.2.0a20261005, following `gfx1250-strict-testing.md`: - hipblaslt-test (all, and `*quick*`) - FP4 bench and mxf4 smoke - hipsparselt-test - TensileLite unit, arch-stepping and packaging tests - tox `Tests/common` - pre-commit (hipBLASLt and repo root configs) ## Test Result - **Client TU compile:** - Original headers fail on `gfx1250-strict` for 20261004, 20261005 and 20261006. - With this change, both targets build on every nightly from 20260930 to 20261006. - **Clean builds, 20260930 to 20261006:** rc=0 on every nightly. - hipBLASLt: 400 / 397 / 400 files for `gfx1250` / `gfx1250v0` / `gfx1250-strict`, layout OK. - hipSPARSELt: 131 files in each subtree. - **gfx1250 (v1), ROCm 10.2.0a20261005:** | Suite | Result | |---|---| | hipblaslt-test all | 67569 passed, 0 failed, 34 skipped | | hipblaslt-test `*quick*` | 18043 passed | | FP4 bench / mxf4 smoke | 345 solutions, no nan / passed | | hipsparselt-test | 165405 / 165405 passed | | TensileLite unit, arch-stepping, packaging | passed | | tox `Tests/common` | 9642 passed, 1 failed | | pre-commit, both configs | passed | The 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 `d116965bed9` plus this change. The branch is rebased onto develop `9db04770489`. The 3 commits in between touch only rocALUTION and rocSPARSE. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
sunhmy-amd
added a commit
to ROCm/rocm-systems
that referenced
this pull request
Oct 6, 2026
The macro gates the pk16 fp6/bf6 scaled converts, not block16 mode. gfx1250-strict has those opcodes (LCOMPILER-2841, llvm/llvm-project#227475) but not the block16 half of the scale_sel range, so a gate named BLOCK16 enabled under __gfx1250_strict__ reads as a contradiction and invites a block16-mode call site that would silently alias to scale_sel 0-3 on A0. Rename to HIP_ENABLE_GFX1250_PK16_SCALE_BUILTINS, matching the existing HIP_ENABLE_GFX1250_PK8_SCALE_BUILTINS. Identifier rename only; no change to gating conditions or generated code. Addresses review feedback from Siu Chi Chan on #12070.
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.
These instructions can be block16 or block32 depending on the target
and scale_sel bits. Block16 is not supported in strict mode.
Re-enable the rest of the instructions in the strict mode but validate
the scale selector.