Repository navigation
feat(hipblaslt): add tuned gfx942 228-CU cgemm grid-based logic - #13364
Merged
NaveenElumalaiAMD merged 1 commit intoOct 9, 2026
Merged
NaveenElumalaiAMD merged 1 commit into
NaveenElumalaiAMD merged 1 commit into
Conversation
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.
✅ 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 #13364 +/- ##
===========================================
- Coverage 72.16% 72.16% -0.00%
===========================================
Files 2874 2874
Lines 474873 474873
Branches 70141 70141
===========================================
- Hits 342674 342673 -1
Misses 108398 108398
- Partials 23801 23802 +1
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
pkamd
self-requested a review
October 9, 2026 19:01
pkamd
approved these changes
Oct 9, 2026
pkamd
left a comment
pkamd
left a comment
Contributor
There was a problem hiding this comment.
I ran the internal PR automation script and there are no errors that need to be addressed. @NaveenElumalaiAMD has already validated the kernels being added via rocblas-test to ensure the kernels are valid and in terms of performance the said library is adhering to the perf requirement.
NaveenElumalaiAMD
deleted the
users/nelumala/hipBLASLt/add-cgemm-tt-MI300A-liblogic
branch
October 9, 2026 19:08
assistant-librarian Bot
pushed a commit
to ROCm/hipBLASLt
that referenced
this pull request
Oct 9, 2026
feat(hipblaslt): add tuned gfx942 228-CU cgemm grid-based logic (#13364) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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.
Contributor
TheRock Submodule Bump ActivityNewest first
|
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-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/GridBaseddirectory, 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 inComplexConjugateA/B, so the tuned solution set and grid table are reusable verbatim.Each file declares
CUCount: 228andUseKdTree: true, scoping the tables to the 228-CU part. TheExactLogicperf 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 shippedgfx942/GridBasedtables.Tuning basis:
pre_merge_md5=a3b009e270d87128f44a7e407f140f32->tuned_md5=9a2b8e28e2fe075c8f2ac9e488714e35; 1,382 tuned rows merged into 2,103 sizes / 1,974 solutions viaTensileMergeLibrary. 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 viarocblas-testwith input typef32_c: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.