Repository navigation
Add tuned Equality library logic for five gfx1250 GEMM - #11746
Merged
Merged
Conversation
validateMIParameters asserted MatrixInstM/N against MatrixInstruction[0]/[1] unconditionally; with SourceSwap they are transposed (F4 32x16 WMMA)
Add 5 tuned Equality logic yamls (BBS, F8F8S, F4BS MXAE8B32, F8BS MXAE8B32, F4BS MXAF8B16)
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
geotseng-amd
approved these changes
Sep 7, 2026
Serge45
approved these changes
Sep 7, 2026
Codecov Report❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## develop #11746 +/- ##
========================================
Coverage 70.24% 70.24%
========================================
Files 2809 2809
Lines 462408 462412 +4
Branches 68093 68094 +1
========================================
+ Hits 324780 324788 +8
+ Misses 114064 114062 -2
+ Partials 23564 23562 -2
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
jichangjichang
enabled auto-merge (squash)
September 7, 2026 11:28
assistant-librarian Bot
pushed a commit
to ROCm/hipBLASLt
that referenced
this pull request
Sep 7, 2026
Add tuned Equality library logic for five gfx1250 GEMM
(#11746)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
JIRA ID : AIHPBLAS-835
## Motivation
Add tuned Equality library logic for five gfx1250 GEMM types (BBS,
F8F8S, MXF8,
MXF4, NVF4). Building the device library with these yamls surfaced a
latent
validator bug: a solution using SourceSwap (F4 32x16 WMMA) failed
MatrixInstruction validation and aborted `TensileCreateLibrary` with a
blank
`Validation failed: (index: 0)`. Both changes are needed for the new
logic to
build.
## Technical Details
**Commit 1 — feat: 5 new Equality logic yamls**
Under `Logic/asm_full/gfx1250/gfx1250/Equality/`:
`BBS_BH`, `F8F8S_BH`, `F4BS_MXAE8B32_MXBE8B32` (MXF4),
`F8BS_MXAE8B32_MXBE8B32`
(MXF8), `F4BS_MXAF8B16_MXBF8B16` (NVF4).
**Commit 2 — fix: honor SourceSwap in MatrixInst M/N validation**
(`SolutionStructs/Validators/MatrixInstruction.py`).
`validateMIParameters`
asserted `MatrixInstM/N == MatrixInstruction[0]/[1]` unconditionally.
With
`SourceSwap` on a non-square instruction the M/N extents are transposed
(e.g. `[32,16,128,1]` with `MatrixInstM=16, MatrixInstN=32`). The
validator is
**dual-use**: solution generation keeps the instruction order
(`M=mi4[0],
N=mi4[1]`) while the serialized logic yaml stores the swapped order — so
pinning
a single orientation fixes one path and breaks the other. The fix
accepts
**either orientation** under SourceSwap (`{MatrixInstM, MatrixInstN} ==
{mi4[0],
mi4[1]}`), keeps the strict order otherwise, and adds `elineno()`
messages so the
assert is no longer blank.
## Test Plan
- hipblaslt-test
## Test Result
- hipblaslt-test test pass
## Submission Checklist
- [ ] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
1 task done
jichangjichang
added a commit
that referenced
this pull request
Sep 8, 2026
#11773) JIRA ID : AIHPBLAS-835 ## Motivation #11746 caused the issue. Building the gfx1250 device library and benchmarking F4 SourceSwap solutions failed on the MatrixInstruction validator. `validateMIParameters` is dual-use: - **Solution generation** (`matrixInstructionToMIParameters`) sets `MatrixInstM = mi[0]`, `MatrixInstN = mi[1]` in the instruction's native operand order (SourceSwap not applied). - **Serialized LibraryLogic yaml** stores the SourceSwap-transposed order, so `MatrixInstM`/`MatrixInstN` come in swapped. The prior code asserted *only* the transposed orientation when `SourceSwap` was set. That broke the un-swapped path for non-square instructions such as the F4 32×16 WMMA ## Technical Details When `SourceSwap` is set, compare `MatrixInstM`/`MatrixInstN` against `mi4[0]`/`mi4[1]` as an **unordered set** instead of asserting a fixed transposed mapping. This accepts both conventions (native benchmark order and swapped logic order) while still rejecting genuinely mismatched M/N. The non-SourceSwap branch and the K/B asserts are unchanged. Test Plan - Run the previously-failing F4 yamls through the full BenchmarkProblems → LibraryLogic pipeline - Confirm zero Error processing permutation and a valid best solution. - Confirm non-SourceSwap and square-instruction paths still validate. Test Result Pass Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
Contributor
TheRock Submodule Bump ActivityNewest first
|
1 task done
mengzcai
added a commit
that referenced
this pull request
Sep 14, 2026
JIRA ID : AIHPBLAS-835 ## Motivation `validateMIParameters` is dual-use but was fed MatrixInstM/N in two conventions: the tuning path (`matrixInstructionToMIParameters`) produced the physical instruction order, while serialized logic solutions carry the effective (SourceSwap-transposed) order. Related works: #11746 , #11773 <!-- Explain the purpose of this PR and the goals it aims to achieve. --> ## Technical Details - `matrixInstructionToMIParameters` gains `sourceSwap` (default False, so existing callers are unchanged) and stores the *effective* extents via `effectiveMatrixInstMN` as its final step. Internal derivations and MIBlock keep the physical mi[0]/mi[1] for the opcode. Its three callers pass SourceSwap. - `validateMIParameters` derives the expected M/N from the physical MatrixInstruction with the same helper and compares in effective space, replacing both prior variants and again rejecting a genuinely transposed M/N. - No consumer change: Solution.py already re-derives M/N from MIBlock, and physical-sensitive code reads MIBlock. <!-- Explain the changes along with any relevant GitHub links. --> ## Test Plan New GPU-less guard (`test_fp4_SourceSwap_nonsquare_guard.py`): runs the `--cpu-only` benchmark flow on a config whose only candidates are 32x16 WMMA + SourceSwap:[true]. With no square / SS-false fallback a regression rejects every candidate -> 0 valid solutions -> `printExit` fails the run, the failure a mixed candidate list would silently mask. gfx1250 is named explicitly (gfx1250v0 has no fp4 32x16 opcode). Plus a negative control and existing SourceSwap:false paths as regression. <!-- Explain any relevant testing done to verify this PR. --> ## Test Result - `tox -e unit -- Tensile/Tests/unit/test_fp4_SourceSwap_nonsquare_guard.py` -> **1 passed in ~3.8s**. - Negative control (producer forced back to physical M/N) -> `Tensile::FATAL: 0 valid solutions` -> **1 failed**. - gfx1250 hardware tuning: passed - gfx1250 build and hipblaslt-test passed <!-- Briefly summarize test outcomes. --> ## 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>
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.
JIRA ID : AIHPBLAS-835
Motivation
Add tuned Equality library logic for five gfx1250 GEMM types (BBS, F8F8S, MXF8,
MXF4, NVF4). Building the device library with these yamls surfaced a latent
validator bug: a solution using SourceSwap (F4 32x16 WMMA) failed
MatrixInstruction validation and aborted
TensileCreateLibrarywith a blankValidation failed: (index: 0). Both changes are needed for the new logic tobuild.
Technical Details
Commit 1 — feat: 5 new Equality logic yamls
Under
Logic/asm_full/gfx1250/gfx1250/Equality/:BBS_BH,F8F8S_BH,F4BS_MXAE8B32_MXBE8B32(MXF4),F8BS_MXAE8B32_MXBE8B32(MXF8),
F4BS_MXAF8B16_MXBF8B16(NVF4).Commit 2 — fix: honor SourceSwap in MatrixInst M/N validation
(
SolutionStructs/Validators/MatrixInstruction.py).validateMIParametersasserted
MatrixInstM/N == MatrixInstruction[0]/[1]unconditionally. WithSourceSwapon a non-square instruction the M/N extents are transposed(e.g.
[32,16,128,1]withMatrixInstM=16, MatrixInstN=32). The validator isdual-use: solution generation keeps the instruction order (
M=mi4[0], N=mi4[1]) while the serialized logic yaml stores the swapped order — so pinninga single orientation fixes one path and breaks the other. The fix accepts
either orientation under SourceSwap (
{MatrixInstM, MatrixInstN} == {mi4[0], mi4[1]}), keeps the strict order otherwise, and addselineno()messages so theassert is no longer blank.
Test Plan
Test Result
[----------] Global test environment tear-down
[==========] 61106 tests from 36 test suites ran. (653987 ms total)
[ PASSED ] 61085 tests.
hipBLASLt version: 100401
hipBLASLt git version: b91169d5c1b
command line: ./hipblaslt-test
Submission Checklist