Repository navigation
fix(tensilelite): accept either M/N orientation for SourceSwap MI val… - #11773
Merged
Merged
Conversation
…idation validateMIParameters is dual-use: solution generation (matrixInstructionToMIParameters) keeps the instruction's native operand order (M=mi[0], N=mi[1]), while the serialized LibraryLogic yaml stores the SourceSwap-transposed order. The previous strict-swap assert only accepted the transposed orientation, so the benchmark/solution-generation path failed for non-square SourceSwap instructions (e.g. F4 32x16 WMMA)
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. |
|
🎉 All checks passed! This PR is ready for review. |
geotseng-amd
approved these changes
Sep 8, 2026
aazz44ss
approved these changes
Sep 8, 2026
assistant-librarian Bot
pushed a commit
to ROCm/hipBLASLt
that referenced
this pull request
Sep 8, 2026
=?UTF-8?q?fix(tensilelite):=20accept=20either=20M/N=20ori?=
=?UTF-8?q?entation=20for=20SourceSwap=20MI=20val=E2=80=A6=20(#11773)?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
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.
Codecov Report❌ Patch coverage is
❌ Your project status has failed because the head coverage (35.22%) 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 #11773 +/- ##
========================================
Coverage 69.41% 69.41%
========================================
Files 2828 2828
Lines 470121 470120 -1
Branches 69418 69418
========================================
Hits 326311 326311
+ Misses 120105 120104 -1
Partials 23705 23705
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
1 task
1 task done
Contributor
TheRock Submodule Bump ActivityNewest first
|
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
#11746 caused the issue.
Building the gfx1250 device library and benchmarking F4 SourceSwap solutions
failed on the MatrixInstruction validator.
validateMIParametersis dual-use:matrixInstructionToMIParameters) setsMatrixInstM = mi[0],MatrixInstN = mi[1]in the instruction's nativeoperand order (SourceSwap not applied).
so
MatrixInstM/MatrixInstNcome in swapped.The prior code asserted only the transposed orientation when
SourceSwapwas set. That broke the un-swapped path for non-square instructions such as
the F4 32×16 WMMA
Technical Details
When
SourceSwapis set, compareMatrixInstM/MatrixInstNagainstmi4[0]/mi4[1]as an unordered set instead of asserting a fixedtransposed 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
LibraryLogic pipeline
Test Result
Pass
Submission Checklist