Repository navigation
feat(hipblaslt): add tuned gfx942 228-CU zgemm grid-based logic - #12611
Merged
NaveenElumalaiAMD merged 2 commits intoOct 1, 2026
Merged
NaveenElumalaiAMD merged 2 commits into
NaveenElumalaiAMD merged 2 commits into
Conversation
Adds the nine zgemm (Z_B) library-logic files for aquavanjaram gfx942_228cu/GridBased, covering every transA/transB pair: tuned directly NN, NT, TN, TT derived from T NC, CN, TC, CT, CC The four tuned files come from a benchmark run on MI300A. A benchmark run produces tuning only for N/T operands, but a 'C' operand is a conjugate transpose: it shares the transposed data layout of 'T' and differs only in ComplexConjugateA/B. The tuning -- solution set and grid table alike -- is therefore reusable verbatim, so each of the five conjugate files carries the same rows as its source 'T' file and differs only in the problem type, the embedded solution/kernel names, and the CC parameter overrides below. CC is doubly conjugated and needs those overrides. A complex MFMA is emulated with four instructions, and conjugating both operands checks out extra negation temporaries (ccVgprs) per iteration. On the tightest tiles that pushes the VGPR pool past MaxVgpr and checkResources() aborts codegen with overflowedResources == 1. AMD's shipped gfx942/GridBased CC file resolves this by flipping ClusterLocalRead, UseSgprForGRO and _UseSgprForGRO on exactly the affected solutions, and only in CC; the same overrides are applied here to the solutions matched by full parameter signature. Each file declares CUCount: 228 and UseKdTree: true, scoping the tables to the 228-CU part and opting them into the 2-D (M, N) kd-tree selection path. Eight files use the list header form; Ailk_Bljk is of the dict vintage and declares the same two keys as top-level entries. Both forms are handled by LibraryIO, and each file is left in the vintage it already used. KernArgsVersion is deliberately left unpinned. Every grid row in all nine tables has batch count 1 -- 8696 rows across 4220 distinct (M, N) cells, with no cell lacking a B == 1 row -- so the kd-tree's in-cell batch handling, which indexes on (M, N) only and carries no batch dimension, is never asked to resolve a cell that holds no row for the problem's batch count. Batched zgemm problems take the hierarchical path, which orders on B directly. Validation: scripts/run_tensile_logic_check.py over the nine files keeps 7030 / 7030 solutions, rejecting none. rocblas-test zgemm coverage passes 30296 / 30296.
✅ 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 #12611 +/- ##
===========================================
+ Coverage 70.43% 70.44% +0.01%
===========================================
Files 2933 2933
Lines 482478 482463 -15
Branches 69519 69515 -4
===========================================
+ Hits 339816 339850 +34
+ Misses 118813 118757 -56
- Partials 23849 23856 +7
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
babakpst
approved these changes
Sep 30, 2026
NaveenElumalaiAMD
deleted the
users/nelumala/hipBLASLt/add-zgemm-MI300A-liblogic
branch
October 1, 2026 00:34
assistant-librarian Bot
pushed a commit
to ROCm/hipBLASLt
that referenced
this pull request
Oct 1, 2026
feat(hipblaslt): add tuned gfx942 228-CU zgemm grid-based logic (#12611) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit JIRA ID : https://amd-hub.atlassian.net/browse/AIHPBLAS-3404 ## Motivation ZGEMM on MI300A has no tuned grid-based library logic for the 228-CU part and falls back to the generic gfx942 tables. This adds CU-scoped tuning for that part, covering all nine transA/transB combinations. ## Technical Details Nine ZGEMM (`Z_B`) library-logic files for `aquavanjaram/gfx942_228cu/GridBased` — `NN`, `NT`, `TN`, `TT` tuned directly on MI300A, and `NC`, `CN`, `TC`, `CT`, `CC` derived verbatim from the matching `T` file, since a `C` operand shares the transposed layout of `T` and differs only in `ComplexConjugateA/B`. `CC` additionally carries the `ClusterLocalRead` / `UseSgprForGRO` / `_UseSgprForGRO` overrides that AMD's own shipped `gfx942/GridBased` `CC` file uses to keep double conjugation's extra negation temporaries from overflowing the VGPR pool on the tightest tiles. Each file declares `CUCount: 228` and `UseKdTree: true`; eight use the list header form and `Ailk_Bljk` is of the dict vintage declaring the same two keys as top-level entries, both of which `LibraryIO` handles. Every grid row in all nine tables has batch count 1 (8,696 rows, 4,220 cells, zero cells lacking a `B == 1` row), so the kd-tree's 2-D `(M, N)` index is never asked to resolve a cell holding rows only for a different batch count. ## Test Plan Build hipBLASLt from this branch and rocBLAS against it (`ROCBLAS_USE_HIPBLASLT=1`), then validate the nine new logic files with `scripts/run_tensile_logic_check.py` and run the full ZGEMM surface of `rocblas-test` (input type `f64_c`). The batch coverage of all nine grid tables is also verified structurally by scanning their `[M, N, B, K]` keys. ## Test Result ```bash ROCBLAS_USE_HIPBLASLT=1 \ ./rocblas-test --gtest_filter='*blas3_tensile*f64_c*-*known_bug*:*graph_test*' ``` | Check | Result | |---|---| | `run_tensile_logic_check.py` over the nine new files | **7030 / 7030 solutions kept, 0 rejected** | | `rocblas-test` ZGEMM | **30296 / 30296 PASS** | | Batch coverage scan (9 tables) | 8696 rows, 4220 cells, `B ∈ {1}`, 0 cells without a `B == 1` row | ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level **Low.** Purely additive — nine new library-logic files, no runtime or codegen source touched — and scoped by an exact `{gfx942, CUCount: 228}` predicate, so no other architecture or CU configuration can select it.
Contributor
TheRock Submodule Bump ActivityNewest first
|
NaveenElumalaiAMD
added a commit
that referenced
this pull request
Oct 8, 2026
Adds the nine cgemm (C_B) library-logic files for aquavanjaram gfx942_228cu/GridBased, covering every transA/transB pair (N/T/C) analogous to the existing ZB (zgemm) set (#12611): tuned directly TT (Alik_Bjlk) derived from TT TC, CT, CC, and the remaining N-combinations Tuning basis: 3,185-shape CGEMM TT-orientation benchmark corpus on MI300A (gfx942, 228 CU), GA/geko-tuned, merged via TensileMergeLibrary (pre_merge_md5=a3b009e270d87128f44a7e407f140f32, tuned_md5=9a2b8e28e2fe075c8f2ac9e488714e35; 1,382 tuned rows merged into 2,103 sizes / 1,974 solutions). See docs/PR_cgemm_TT_CU228_tuned.md in the perf-tuning-cgemm-TT-MI300A tracking repo for full before/after results (both-pass confirmed: 1,793/3,185 -> 2,408/3,185 wins, +19.3pp). Each file declares CUCount: 228 and UseKdTree: true, scoping the tables to the 228-CU part. The ExactLogic perf column is zeroed (not used at runtime; selection is by nearest-size kd-tree match, not by perf comparison), matching the convention already used by the stock shipped gfx942/GridBased tables.
NaveenElumalaiAMD
added a commit
that referenced
this pull request
Oct 9, 2026
JIRA ID : https://amd-hub.atlassian.net/browse/AIHPBLAS-3403 ## Motivation CGEMM on MI300A has no tuned grid-based library logic for the 228-CU part and falls back to the generic gfx942 tables. ## Technical Details This contribution adds nine CGEMM library-logic files for the `aquavanjaram/gfx942_228cu/GridBased` directory, analogous to the ZGEMM set added in #12611. The TT combination (`Alik_Bjlk`) was GA/geko-tuned directly on MI300A hardware against a 3,185-shape CGEMM TT-orientation benchmark corpus; the remaining eight transA/transB combinations derive from it, following the same approach as #12611: a conjugate ('C') operand shares the transposed data layout of its 'T' counterpart and differs only in `ComplexConjugateA/B`, so the tuned solution set and grid table are reusable verbatim. Each file declares `CUCount: 228` and `UseKdTree: true`, scoping the tables to the 228-CU part. The `ExactLogic` perf column is zeroed in all nine files — it is not consulted at runtime (solution selection is by nearest-size kd-tree match, not perf comparison), matching the convention already used by the stock shipped `gfx942/GridBased` tables. Tuning basis: `pre_merge_md5=a3b009e270d87128f44a7e407f140f32` -> `tuned_md5=9a2b8e28e2fe075c8f2ac9e488714e35`; 1,382 tuned rows merged into 2,103 sizes / 1,974 solutions via `TensileMergeLibrary`. librocblas / libhipblaslt binaries are unchanged — only this tuned table differs from stock. ## Test Plan Build hipBLASLt and rocBLAS with `ROCBLAS_USE_HIPBLASLT=1`, then run full CGEMM coverage via `rocblas-test` with input type `f32_c`: ``` ROCBLAS_USE_HIPBLASLT=1 ./rocblas-test --gtest_filter='*blas3_tensile*f32_c*-*known_bug*:*graph_test*' ``` ## Test Result `ROCBLAS_USE_HIPBLASLT=1 ./rocblas-test --gtest_filter='*blas3_tensile*f32_c*-*known_bug*:*graph_test*'` 30,326 ran / 30,326 passed / 0 failed, identical in both arms.
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 : https://amd-hub.atlassian.net/browse/AIHPBLAS-3404
Motivation
ZGEMM on MI300A has no tuned grid-based library logic for the 228-CU part and falls back to the generic gfx942 tables. This adds CU-scoped tuning for that part, covering all nine transA/transB combinations.
Technical Details
Nine ZGEMM (
Z_B) library-logic files foraquavanjaram/gfx942_228cu/GridBased—NN,NT,TN,TTtuned directly on MI300A, andNC,CN,TC,CT,CCderived verbatim from the matchingTfile, since aCoperand shares the transposed layout ofTand differs only inComplexConjugateA/B.CCadditionally carries theClusterLocalRead/UseSgprForGRO/_UseSgprForGROoverrides that AMD's own shippedgfx942/GridBasedCCfile uses to keep double conjugation's extra negation temporaries from overflowing the VGPR pool on the tightest tiles.Each file declares
CUCount: 228andUseKdTree: true; eight use the list header form andAilk_Bljkis of the dict vintage declaring the same two keys as top-level entries, both of whichLibraryIOhandles. Every grid row in all nine tables has batch count 1 (8,696 rows, 4,220 cells, zero cells lacking aB == 1row), so the kd-tree's 2-D(M, N)index is never asked to resolve a cell holding rows only for a different batch count.Test Plan
Build hipBLASLt from this branch and rocBLAS against it (
ROCBLAS_USE_HIPBLASLT=1), then validate the nine new logic files withscripts/run_tensile_logic_check.pyand run the full ZGEMM surface ofrocblas-test(input typef64_c). The batch coverage of all nine grid tables is also verified structurally by scanning their[M, N, B, K]keys.Test Result
ROCBLAS_USE_HIPBLASLT=1 \ ./rocblas-test --gtest_filter='*blas3_tensile*f64_c*-*known_bug*:*graph_test*'run_tensile_logic_check.pyover the nine new filesrocblas-testZGEMMB ∈ {1}, 0 cells without aB == 1rowSubmission Checklist
Risk level
Low. Purely additive — nine new library-logic files, no runtime or codegen source touched — and scoped by an exact
{gfx942, CUCount: 228}predicate, so no other architecture or CU configuration can select it.