Repository navigation
feat(hipblaslt): add MAF/MAB kernels for MC system - #11274
Merged
Merged
Conversation
✅ 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. |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #11274 +/- ##
========================================
Coverage 69.65% 69.65%
========================================
Files 2799 2799
Lines 461048 461048
Branches 67958 67958
========================================
Hits 321103 321103
Misses 116703 116703
Partials 23242 23242
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
minsukim-amd
changed the title
babakpst
approved these changes
Aug 27, 2026
tony-davis
added a commit
that referenced
this pull request
Aug 28, 2026
…r CI gaps `test_PlaceholderMerge.py` and `test_GpuRevisionTarget.py` validate the production logic YAML corpus for cross-file naming and metadata invariants, a third test category the document did not name alongside unit and characterization. Add a subsection for it, with the gfx942 placeholder-merge incident (PR #6946) as the origin story. The `_needs_logic_dir` xfail that gates the corpus-dependent tests in both files is an eighth known-bug mechanism, uncounted in the existing "seven different places" tally. Add it to that table and classify it: it is conditioned on an environment, not a bug, and in TheRock CI's installed-artifact layout that environment is permanent, so the tests it guards can never run for real there. Add two gaps to CI visibility and gating: Math CI's `preliminary` job appears to skip the unit/characterization suite entirely on YAML-only diffs, and the xfail above is unconditional in TheRock CI. Both were confirmed live via the 2026-08-27 `develop` break, where a YAML-only directory restructure (#11274) surfaced a months-old DeviceNames mismatch that neither CI system's own PR run caught. Add a roadmap item to fold the two data-validating tests here (plus the four in test_GpuRevisionTarget.py) into `TensileLogic --check-all`, which runs unconditionally wherever kernels are generated and would close both gaps at once; the three code-level tests in test_PlaceholderMerge.py stay where they are. Co-authored-by: Cursor <cursoragent@cursor.com>
1 task done
tony-davis
pushed a commit
that referenced
this pull request
Aug 29, 2026
…ers (#11442) JIRA ID : AIHPBLAS-4531 GitHub ID: #11397 ## Motivation `Tensile/Tests/unit/test_PlaceholderMerge.py::test_logic_yaml_sibling_device_names_consistent` is currently **failing on `develop`**, turning CI red for unrelated hipBLASLt PRs (observed on #10941 among others). This PR fixes the underlying data divergence so the shared CI lane goes green again. The test asserts that library-logic YAMLs sharing a basename within one architecture directory all declare the same `DeviceNames` header. Four gfx1250 `Equality/` basenames declared `- [Device 73f0, Device 0073, Device 75c1]` while their `GridBased/` and `Origami/` siblings declared `[Device 73f0]`. Provenance: the 3-ID list was introduced by `3a08f1b4d94` (#10250, a tuning PR whose description never mentions device IDs, admin-merged with partial CI) and never propagated to the sibling trees — 208 of the 212 files in each affected arch directory use `[Device 73f0]`. The divergence has existed since then but was invisible to CI. `134466bcb4a` (#11274) later re-nested the tree (`asm_full/gfx1250/Equality/` -> `asm_full/gfx1250/gfx1250/Equality/`), which caused the test's `_iter_arch_dirs()` to start descending into that directory and **surfaced** the pre-existing drift. #11274 is not the cause — the divergence is verifiably present at `134466bcb4a^`. ## Technical Details **Commit 1 — normalize the data (8 files).** Four basenames x the `gfx1250` and `gfx1250v0` trees = 8 `Equality/` YAMLs narrowed from `- [Device 73f0, Device 0073, Device 75c1]` to `- [Device 73f0]`, matching their siblings. Narrowing 8 files was chosen over widening the 416 siblings to two additional IDs of unknown provenance (see the open question below). **Why this is functionally safe.** `DeviceNames` entries are PCIe chip IDs consumed to emit a runtime `PciChipId` hardware predicate — but only where `supportsChipIdPredicate()` is true, which is `gfx950` and nothing else: - `tensilelite/Tensile/Common/Architectures.py:180-186` - mirrored in C++ at `tensilelite/include/Tensile/AMDGPUPredicates.hpp:53-56` On gfx1250 the field is parsed and then dropped: no `PciChipId` predicate is emitted and no `_ID` filename suffix is generated. It is inert for kernel selection on this architecture, so no GPU sees a different kernel set as a result of this change. The value does remain live in the Python solution de-duplication key (`Tensile/SolutionStructs/Solution.py:6193-6205`, which is not arch-gated) — which is exactly why keeping it consistent across siblings matters. **Commit 2 — close a test blind spot (separate concern, 1 file).** `_DEVICE_NAMES_RE` in `test_PlaceholderMerge.py` matched only the positional list form (`- [Device ...]`) and silently skipped the mapping form (`DeviceNames: [Device ...]`) used by Origami logic files. The regex now matches both. This is a genuine coverage gap, not cosmetic: one of the four basenames (`*_BBS_BH_Bias_HA_S_SAV_UserArgs.yaml`) diverges only against a mapping-form Origami sibling and was therefore invisible to the old test. Parse coverage over the `asm_full/` tree goes from **1840/1936 to 1936/1936** files. ## Test Plan Regression coverage is the pre-existing shared-CI unit test itself — this is the test that is red on `develop`, so the fix is validated by the same lane that caught it (`test_PlaceholderMerge.py` matches the `test_*` pattern picked up by the Libraries PR Bot unit-test check). Run locally with `rocisa` built via `tlbuild` into `build_tmp` and executed through a `tlrun`-style `PYTHONPATH` (nothing pip-installed): ```bash pytest projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_PlaceholderMerge.py pytest projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_PlaceholderMerge.py::test_logic_yaml_sibling_device_names_consistent ``` Control run: the same test executed against the pre-fix YAML content, to confirm it actually detects the condition rather than passing vacuously. ## Test Result - `test_logic_yaml_sibling_device_names_consistent` — **PASSED**. - Full `test_PlaceholderMerge.py` — **22 passed**. - Control run against pre-fix data — **failed as expected**, with `AssertionError: Divergent sibling DeviceNames` listing both `*_BBS_*` pairs. Confirms the assertion is live and the widened regex is what surfaces the second pair. - Post-rebase scan of the whole `asm_full/` tree (replicating `_iter_arch_dirs` + `_read_device_names` + the grouping assert): **1936 YAML files walked, 1936 with `DeviceNames` parsed, 0 violations.** Re-run after rebasing onto current `origin/develop` to confirm no newly-landed YAMLs reintroduce divergence. - Branch rebased onto current `origin/develop` (clean, no conflicts) per the `hipblaslt-pr-quality` stale-base rule for validator-vs-validated-data changes. - Automated GitHub checks are pending. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. ## Risk level **Low (1/5).** Eight data files change a field that is provably inert on the only architecture they apply to (`supportsChipIdPredicate()` is gfx950-only), plus one widened regex in a host-side unit test. No kernel source, no build output, no host library code, no public API. The change makes 8 files agree with the 208 already-consistent files in the same directory.
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.
Motivation
Add MAF/MAB kernels for MC system
Technical Details
Restructured gfx1250 library like below:
All MC MAF/MAB kernels are added into
*_192cuand_96cuEquality libraries.Test Plan
Run tensilelite verification test for new MAF/MAB kernels
Run hipblaslt-bench verification test for new MAF/MAB kernels
Run hipblaslt-test
Test Result
Passed tensilelite verification test for new MAF/MAB kernels
Passed hipblaslt-bench verification test for new MAF/MAB kernels
Passed hipblaslt-test
Submission Checklist
JIRA ID : AIHPBLAS-4351