Repository navigation
[hipblaslt] Equality tuning for BF16, MXF8, MXF4, F8 MAF - #10250
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 #10250 +/- ##
===========================================
- Coverage 68.93% 68.78% -0.16%
===========================================
Files 2762 2741 -21
Lines 455034 451375 -3659
Branches 67122 66507 -615
===========================================
- Hits 313673 310442 -3231
+ Misses 118430 118285 -145
+ Partials 22931 22648 -283
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
|
Merging this with Both attempt-2 re-runs are done —
What I am not signing off on: Why not just wait: |
[hipblaslt] Equality tuning for BF16, MXF8, MXF4, F8 MAF (#10250) Equality tuning for 4K x 4K x 65536 shape. Added 4 liblogic yaml files for each BBS_TN, F8F8S_TN, MX_F8BS_TN, MX_F4BS_TN. Tested with tensilelite and hipblaslt-bench verification run. Merged with --admin for the BKC release code cut-off. Math CI Summary and pre-commit green; the reds in TheRock CI Summary are ROCm/rocm-libraries#10522 (rocjitsu libomp.so, fixed on develop after this run was created) and ROCm/TheRock#7170 (30-min amd-llvm clone timeout on gfx125X, with a same-run green gfx950 build as control). Rationale in the PR thread. JIRA ID : AIHPBLAS-4151
|
Closing the loop on the one gap I flagged when merging this: gfx1250 now has a real build. The So the Also confirmed: the subrepo mirror landed in |
…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.
Motivation
Measure MAF for various datatypes
Technical Details
Equality tuning for 4K x 4K x 65536 shape
Added 4 liblogic yaml files for each BBS_TN, F8F8S_TN, MX_F8BS_TN, MX_F4BS_TN
Test Plan
Tested with tensilelite and hipblaslt-bench verification run
Test Result
Passed
Submission Checklist
JIRA ID : AIHPBLAS-4151