Skip to content

fix(tensilelite): validate SourceSwap MI in effective space - #11877

Merged
mengzcai merged 2 commits into
ROCm:developfrom
mengzcai:users/mengzcai/fix_hipbaslt_f4_SS_build
Sep 14, 2026
Merged

mengzcai merged 2 commits into
ROCm:developfrom
mengzcai:users/mengzcai/fix_hipbaslt_f4_SS_build

Conversation

@mengzcai

@mengzcai mengzcai commented Sep 9, 2026 •

Copy link
Copy Markdown
Contributor

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

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.

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.

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

Submission Checklist

@therock-pr-bot

therock-pr-bot Bot commented Sep 9, 2026 •

Copy link
Copy Markdown

✅ All Policy Checks Passed

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —

🎉 All policy checks passed!

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

🙋 Wish to Override Policy?

@therock-pr-bot

therock-pr-bot Bot commented Sep 9, 2026

Copy link
Copy Markdown

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

@mengzcai
mengzcai force-pushed the users/mengzcai/fix_hipbaslt_f4_SS_build branch from bd6e28e to bd920d4 Compare September 9, 2026 17:58
@mengzcai mengzcai changed the title Users/mengzcai/fix hipbaslt f4 ss build fix(tensilelite): validate SourceSwap MI in effective space Sep 9, 2026
@codecov

codecov Bot commented Sep 9, 2026 •

Copy link
Copy Markdown

Codecov Report

❌ Patch coverage is 85.71429% with 1 line in your changes missing coverage. Please review.

Files with missing lines Patch % Lines
...hipblaslt/tensilelite/Tensile/BenchmarkProblems.py 50.00% 1 Missing ⚠️
Additional details and impacted files
@@             Coverage Diff             @@
##           develop   #11877      +/-   ##
===========================================
+ Coverage    69.42%   70.27%   +0.85%     
===========================================
  Files         2828     2810      -18     
  Lines       470140   462487    -7653     
  Branches     69425    68089    -1336     
===========================================
- Hits        326353   324992    -1361     
+ Misses      120081   113939    -6142     
+ Partials     23706    23556     -150     
Flag Coverage Δ *Carryforward flag
TensileLite-CPP 46.45% <ø> (+7.93%) ⬆️
TensileLite-Unit 76.07% <85.71%> (+0.03%) ⬆️
hipBLAS 90.62% <ø> (ø) Carriedforward from b496fe0
hipBLASLt 35.27% <ø> (+0.05%) ⬆️
hipCUB 82.68% <ø> (ø) Carriedforward from b496fe0
hipDNN 87.02% <ø> (ø) Carriedforward from b496fe0
hipFFT 42.66% <ø> (ø) Carriedforward from b496fe0
hipRAND 76.12% <ø> (ø) Carriedforward from b496fe0
hipSOLVER 68.92% <ø> (ø) Carriedforward from b496fe0
hipSPARSE 86.99% <ø> (ø) Carriedforward from b496fe0
rocBLAS 48.31% <ø> (ø) Carriedforward from b496fe0
rocFFT 47.78% <ø> (ø) Carriedforward from b496fe0
rocRAND 56.90% <ø> (ø) Carriedforward from b496fe0
rocSOLVER 76.83% <ø> (ø) Carriedforward from b496fe0
rocSPARSE 74.61% <ø> (ø) Carriedforward from b496fe0
rocThrust 91.60% <ø> (ø) Carriedforward from b496fe0

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

Files with missing lines Coverage Δ
...le/SolutionStructs/Validators/MatrixInstruction.py 93.16% <100.00%> (+1.05%) ⬆️
...ilelite/Tensile/TensileLogic/HandleCustomKernel.py 100.00% <ø> (ø)
...hipblaslt/tensilelite/Tensile/BenchmarkProblems.py 80.53% <50.00%> (ø)

... and 28 files 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.

@jichangjichang jichangjichang left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Please also check there is no error of building hipblaslt library

@mengzcai
mengzcai force-pushed the users/mengzcai/fix_hipbaslt_f4_SS_build branch from bd920d4 to 67191e0 Compare September 10, 2026 11:17
mengzcai added 2 commits September 11, 2026 17:30
JIRA ID : AIHPBLAS-835

matrixInstructionToMIParameters now emits effective (SourceSwap-transposed)
MatrixInstM/N so the benchmark path matches the serialized-logic convention;
validateMIParameters derives the expected pair via effectiveMatrixInstMN and
drops the orientation-agnostic set check.
Drives the --cpu-only benchmark flow on a config whose only candidates are
32x16 WMMA + SourceSwap:[true]. With no square / SS-false fallback, an
effective-vs-physical MatrixInstM/N regression rejects every candidate, so the
group yields 0 valid solutions and printExit fails the run instead of the drop
being silently masked. No LibraryLogic step: the contract under test fails
inside BenchmarkProblems before any logic yaml is analyzed, so the run asserts
on benchmark data and skips that step's getCUCount device probe.

JIRA ID : AIHPBLAS-835
@mengzcai
mengzcai force-pushed the users/mengzcai/fix_hipbaslt_f4_SS_build branch from 67191e0 to a489af4 Compare September 11, 2026 09:34
@mengzcai

Copy link
Copy Markdown
Contributor Author

Merged. The remaining CI failures (ffm gfx1250) are unrelated to this PR — they also fail on other PRs (e.g. PR #11980, which doesn't touch tox at all).

@mengzcai
mengzcai merged commit b99d0fb into ROCm:develop Sep 14, 2026
115 of 121 checks passed
@mengzcai
mengzcai deleted the users/mengzcai/fix_hipbaslt_f4_SS_build branch September 14, 2026 04:52
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

tony-davis pushed a commit to ROCm/TheRock that referenced this pull request Sep 15, 2026
## Summary

Test selection returns **no tests** for a change confined to
TensileLite, because the nested subtree path
`projects/hipblaslt/tensilelite` has no entry in
`_EXTERNAL_SUBTREE_ALIASES`. This adds the alias.

## Motivation

ROCm/rocm-libraries#11785 taught rocm-libraries' change detection to
report `projects/hipblaslt/tensilelite` (longest-prefix match) instead
of collapsing TensileLite files into `projects/hipblaslt`.
`determine_rocm_test_dependencies.py` has no alias for that string, so
it falls through to the `projects/` strip, produces the non-existent
graph key `hipblaslt/tensilelite`, warns, and selects nothing — while
exiting `0`.

Verified against `df3d451` (the commit rocm-libraries pins today):

```console
$ python test_tools/determine_rocm_test_dependencies.py \
    --changed-projects projects/hipblaslt/tensilelite
Warning: unrecognized project(s) ['hipblaslt/tensilelite']; no tests will be selected for them
["hipblaslt/tensilelite"]
$ echo $?
0
```

The consequence is visible in rocm-libraries. TensileLite-only PRs ran a
sanity check and nothing else:

| PR | files | multi-arch test components |
|---|---|---|
| ROCm/rocm-libraries#11518 | `tensilelite/src` + 2 new
`tensilelite/tests/*.cpp` | no `Test tensilelite` |
| ROCm/rocm-libraries#11896 | 7, all under
`projects/hipblaslt/tensilelite/**` | sanity only |
| ROCm/rocm-libraries#11877 | 6, all under
`projects/hipblaslt/tensilelite/**` | sanity only |

ROCm/rocm-libraries#11518 merged that way and was reverted by
ROCm/rocm-libraries#12022 after failing `Test tensilelite` on gfx94X and
gfx950 in this repo's CI — after merge, blocking the submodule bump in
#8183.

## Technical Details

One alias entry. `tensilelite` is already a synthetic node at level 3
(`test_policies.toml`), so selection also reaches `hipblaslt`, `rocblas`
and `hipblas` transitively — the same set `shared/stinkytofu` already
gets.

After the change, on the real committed graph:

```console
$ python test_tools/determine_rocm_test_dependencies.py \
    --changed-projects projects/hipblaslt/tensilelite
["hipblas", "hipblaslt", "hipblasltprovider", "hipdnn-samples", "hipsolver",
 "hipsparse", "hipsparselt", "miopen", "miopenprovider", "rocalution",
 "rocblas", "rocsolver", "rocsparse", "rocwmma", "tensilelite"]
```

## Risk Assessment

Risk 2. Test-selection only; no product or device behaviour changes. It
strictly widens selection for one input that currently selects nothing,
so the failure mode is longer CI on TensileLite PRs, not reduced
coverage.

## Related

- Work tracking: ROCm/rocm-libraries#11784
- Introduced by: ROCm/rocm-libraries#11785 (rocm-libraries half of the
nested-subtree fix)
- Companion PRs: #7998 (TensileLite as a synthetic node), #7999
(longest-prefix match)
- Fallout: ROCm/rocm-libraries#11518, reverted by
ROCm/rocm-libraries#12022, blocking #8183

## Testing Summary

- Unit: new regression test in
`determine_rocm_test_dependencies_test.py`, plus the existing suite.
- Manual: CLI run against the real committed graph, before and after.

## Testing Checklist

- [x] New regression test — `python -m unittest
test_tools.tests.determine_rocm_test_dependencies_test -k
nested_tensilelite` — Status: Passed
- [x] Verified it fails without the fix (stashed the alias, test fails)
— Status: Passed
- [x] Full selector suite — `python -m unittest
test_tools.tests.determine_rocm_test_dependencies_test` — 70 tests —
Status: Passed
- [ ] PR CI — Status: Pending

The new test declares `[synthetic.tensilelite]` in its fixture,
mirroring the real `test_policies.toml`, so it exercises the walk past
the alias rather than stubbing it. Clean `stderr` is asserted
deliberately: the regression was a warning plus an empty selection, not
a non-zero exit.

## Device / Architecture Coverage

None. Change-detection logic only; no compiled code and no device
behaviour. PR CI is sufficient.

## Adjacent Tests Considered

`shared/stinkytofu`, `shared/origami`, `shared/mxdatagenerator` and
`shared/tensile` all resolve through the same alias table and are
covered by `test_shared_blas_prefixes_mapped`; all still pass.

## Follow-ups not in this PR

1. **`projects/hipblaslt` still does not select `tensilelite`.** A
non-TensileLite hipBLASLt change selects `[hipblas, hipblaslt,
hipblasltprovider, hipsparselt, miopen, rocblas]`. rocm-libraries'
`therock_matrix.py` maps `projects/hipblaslt` → `blas` → `[hipblaslt,
rocblas, hipblas, tensilelite]`. Closing that gap means adding
`tensilelite` as a consumer of `hipblaslt` in the consumer graph, which
is a larger change than an alias and wants its own review.
2. **Unmapped `projects/` paths fail open.** `_EXTERNAL_ONLY_NAMESPACES`
makes an unmapped `shared/` path raise, but an unmapped `projects/` path
warns and exits `0`, yielding a green run with no tests. Extending the
fail-closed behaviour would have surfaced this on the day
ROCm/rocm-libraries#11785 merged rather than via a revert a week later.

Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants