Repository navigation
perf(hipblaslt): retune DGEMM TN 8K square grid-point on MI300x. - #14348
Merged
Merged
Conversation
⏳ Waiting on Required Checks⏳ Required checks have not all concluded yet. This bot run is exiting without failing so those checks can report independently.
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
⏳ Waiting for required checks to start or finish:
|
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #14348 +/- ##
========================================
Coverage 75.87% 75.87%
========================================
Files 2883 2883
Lines 480529 480529
Branches 72254 72254
========================================
Hits 364589 364589
Misses 90897 90897
Partials 25043 25043
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
pkamd
approved these changes
Oct 26, 2026
babakpst
approved these changes
Oct 26, 2026
Contributor
TheRock Submodule Bump ActivityNewest first
|
NaveenElumalaiAMD
added a commit
that referenced
this pull request
Nov 3, 2026
rocHPL shows a DGEMM performance regression on MI325X on this branch. f64 on gfx942 currently routes to hipBLASLt on every device, so MI325X resolves against the generic gfx942 fp64 GridBased tables, which were tuned on MI300X. Narrow the fp64 default back to the 228-CU MI300A, restoring the gate introduced by PR #12724 and later widened to all gfx942 by PR #13292. fp16/bf16 keep their existing 304-CU gate; f32, int8 and both complex types keep the Tensile default on gfx942 (the trailing else preserves the behaviour the previous `!is_same<Ti, double>` arm provided). Effect, verified instantiation-by-instantiation over the seven types this template is instantiated with (half, bfloat16, float, double, float_complex, double_complex, int8_t): MI300A 228 CU -> hipBLASLt (unchanged) MI300X 304 CU -> Tensile (changed) MI325X 304 CU -> Tensile (changed) partitioned other -> Tensile (changed) all other types / arches (unchanged) Note the change is not limited to the 304-CU parts: any partitioned gfx942 (CPX/DPX/TPX) reports its partition's CU count and therefore also moves to Tensile. Those configurations ship no fp64 logic of their own and were previously falling back on the 304-CU tuning. This reverts the fp64 default on MI300X, where PR #13134 and PR #14348 measured hipBLASLt wins. Reviewers should weigh that trade-off; MI300X users can restore the previous behaviour with ROCBLAS_USE_HIPBLASLT=1, which bypasses this gate entirely. useHipBLASLt() also gates getAllSolutions(), so rocblas_gemm_ex_get_ solutions() now reports Tensile solution indices for fp64 on the affected devices. Query and dispatch stay consistent, but cached solution indices will not carry over.
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/ROCM-33102
Motivation
After setting hipBLASLt backend as the default backend of the rocBLAS for DGEMM on MI300x, performance regression is reported with ROCm10.2 release with DGEMM, TN, 8192 x 8192 x 8192 problem compared to the Tensile backend runs. The purpose of this PR is to fix the regression.
Technical Details
I've reproduced the issue locally on an MI300x machine, and retuned the originally assigned grid-point for the regressed problem size. I've retuned the grid-point by using
--algo_method alltuning method, and confirmed that the performances of both [8192, 8192, 1, 8192] and [8192, 8192, 1, 1024] (grid-point) problem sizes are improved by the newly assigned local SolutionIndex of the GridBased lib-logic yaml file.Test Result
The performance of
M=8192, N=8192, K=8192is measured through rocblas-bench with Tensile and hipBLASLt backend, which is shown below table:I've also checked functional correctness locally, and all f64 & TN related hipblaslt tests PASSED:
Submission Checklist