Repository navigation
feat(rocBLAS): default to hipBLASLt for ZGEMM/CGEMM/SGEMM/fp16/bf16 on MI300A (gfx942, 228 CU) - #13045
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. |
TorreZuk
left a comment
There was a problem hiding this comment.
You can ignore code coverage as known issue to be fixed separately.
f16/bf16 fail you could back out and investigate separately as well.
This PR:#12609 has the fix, waiting for review on this PR to merge and then this issue should be resolved. |
|
This overlaps #13292 right? |
|
I guess fix was merged to hipblaslt but develop artifacts with fix likely needed, maybe this afternoon can rebase ? |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #13045 +/- ##
========================================
Coverage 76.04% 76.04%
========================================
Files 2867 2867
Lines 474973 474973
Branches 71126 71126
========================================
Hits 361148 361148
Misses 89279 89279
Partials 24546 24546
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
a472b4e to
475e6cf
Compare
…fx942 (#13292) JIRA ID: https://amd-hub.atlassian.net/browse/AIROCBLAS-1268 JIRA ID: https://amd-hub.atlassian.net/browse/AIROCBLAS-1267 JIRA ID: https://amd-hub.atlassian.net/browse/AIHPBLAS-3019 ## Motivation rocBLAS currently defaults to the Tensile Solution Library for FP16, BF16, and FP64 GEMM on gfx942. As part of the GEMM backend unification initiative, the default backend is flipped to hipBLASLt for these three input types, following the performance parity threshold of within 5% of Tensile or acceptable performance. This extends the two preceding flips — DGEMM on gfx950 (#12424) and DGEMM on the 228-CU MI300A (#12724) — as follows: * FP64 now defaults to hipBLASLt on every gfx942 device (MI300X and MI300A alike), rather than MI300A (228 CUs) only. * FP16 and BF16 now default to hipBLASLt on the 304-CU MI300X only. MI300A (228 CUs) continues to default to Tensile for FP16/BF16 in this PR, but it will be changed by #13045. * ## Technical Details `useHipBLASLt()` (`library/src/tensile_host.cpp`) previously routed only `double` to hipBLASLt on gfx942, and only when `multiProcessorCount == 228` (MI300A). The gfx942 branch for MI300x now applies FP16, BF16, and FP64. No change is needed in `isDefaultHipBLASLtArch()` (`handle.hpp`): `942` was already added there by #12724, and that function only gates eager hipBLASLt handle construction, not per-datatype routing. `docs/conceptual/rocblas-design-notes.rst` is also updated. It previously stated that hipBLASLt is the default "for problems on the gfx12 architecture", which was already stale after #12424 and #12724 flipped FP64 on gfx950 and MI300A. It now carries a table of which architectures default to hipBLASLt and for which data types. ## Test Plan 1. Confirm backend routing per datatype on gfx942 with `ROCBLAS_LAYER=8`. 2. Run the GEMM gtest suite on MI300X for regressions. ## Test Result ### Backend routing (MI300X, gfx942, `ROCBLAS_LAYER=8`) | Input type | Backend selected | |---|---| | `f16_r` | `rocblas_gemm_hipblaslt_backend` | | `bf16_r` (via `gemm_ex`, compute `f32_r`) | `rocblas_gemm_hipblaslt_backend` | | `f64_r` | `rocblas_gemm_hipblaslt_backend` | | `f32_r` | `rocblas_gemm_tensile_backend` (unchanged) | ### gtest `rocblas-test --gtest_filter='*hipblaslt*:*non_hipblaslt*-*known_bug*'`: | PASS | FAIL | SKIP | total | TEST_EXIT | |---|---|---|---|---| | 2,208 | 0 | 0 | 2,208 | 0 | `./rocblas-test --gtest_filter=*gemm*f64_r*-*known_bug*`: | PASS | FAIL | SKIP | total | TEST_EXIT | |---|---|---|---|---| | 163,672 | 0 | 44 | 163,716 | 0 | ` ./rocblas-test --gtest_filter=*gemm*f16_r*:*gemm*bf16_r*-*known_bug*` | PASS | FAIL | SKIP | total | TEST_EXIT | |---|---|---|---|---| | 497,739 | 19 | 192 | 497,950 | 0 | Those failed tests are all stress_gemm with ILP tests, which requires appropriate ILP library link, so after I rebuilt the library, I reran those failed tests as below, and all the tests PASSED. `./rocblas-test --gtest_filter=_/gemm.*stress*ILP64*f16_r*:_/gemm_ex.*stress*ILP64*f16_r*:_/gemm_ex.*stress*ILP64*bf16_r*` | PASS | FAIL | SKIP | total | TEST_EXIT | |---|---|---|---|---| | 23 | 0 | 0 | 23 | 0 | ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
81bf64b to
84bb7e1
Compare
TorreZuk
left a comment
There was a problem hiding this comment.
If all tests passing LGTM
…fx942 (ROCm#13292) JIRA ID: https://amd-hub.atlassian.net/browse/AIROCBLAS-1268 JIRA ID: https://amd-hub.atlassian.net/browse/AIROCBLAS-1267 JIRA ID: https://amd-hub.atlassian.net/browse/AIHPBLAS-3019 ## Motivation rocBLAS currently defaults to the Tensile Solution Library for FP16, BF16, and FP64 GEMM on gfx942. As part of the GEMM backend unification initiative, the default backend is flipped to hipBLASLt for these three input types, following the performance parity threshold of within 5% of Tensile or acceptable performance. This extends the two preceding flips — DGEMM on gfx950 (ROCm#12424) and DGEMM on the 228-CU MI300A (ROCm#12724) — as follows: * FP64 now defaults to hipBLASLt on every gfx942 device (MI300X and MI300A alike), rather than MI300A (228 CUs) only. * FP16 and BF16 now default to hipBLASLt on the 304-CU MI300X only. MI300A (228 CUs) continues to default to Tensile for FP16/BF16 in this PR, but it will be changed by ROCm#13045. * ## Technical Details `useHipBLASLt()` (`library/src/tensile_host.cpp`) previously routed only `double` to hipBLASLt on gfx942, and only when `multiProcessorCount == 228` (MI300A). The gfx942 branch for MI300x now applies FP16, BF16, and FP64. No change is needed in `isDefaultHipBLASLtArch()` (`handle.hpp`): `942` was already added there by ROCm#12724, and that function only gates eager hipBLASLt handle construction, not per-datatype routing. `docs/conceptual/rocblas-design-notes.rst` is also updated. It previously stated that hipBLASLt is the default "for problems on the gfx12 architecture", which was already stale after ROCm#12424 and ROCm#12724 flipped FP64 on gfx950 and MI300A. It now carries a table of which architectures default to hipBLASLt and for which data types. ## Test Plan 1. Confirm backend routing per datatype on gfx942 with `ROCBLAS_LAYER=8`. 2. Run the GEMM gtest suite on MI300X for regressions. ## Test Result ### Backend routing (MI300X, gfx942, `ROCBLAS_LAYER=8`) | Input type | Backend selected | |---|---| | `f16_r` | `rocblas_gemm_hipblaslt_backend` | | `bf16_r` (via `gemm_ex`, compute `f32_r`) | `rocblas_gemm_hipblaslt_backend` | | `f64_r` | `rocblas_gemm_hipblaslt_backend` | | `f32_r` | `rocblas_gemm_tensile_backend` (unchanged) | ### gtest `rocblas-test --gtest_filter='*hipblaslt*:*non_hipblaslt*-*known_bug*'`: | PASS | FAIL | SKIP | total | TEST_EXIT | |---|---|---|---|---| | 2,208 | 0 | 0 | 2,208 | 0 | `./rocblas-test --gtest_filter=*gemm*f64_r*-*known_bug*`: | PASS | FAIL | SKIP | total | TEST_EXIT | |---|---|---|---|---| | 163,672 | 0 | 44 | 163,716 | 0 | ` ./rocblas-test --gtest_filter=*gemm*f16_r*:*gemm*bf16_r*-*known_bug*` | PASS | FAIL | SKIP | total | TEST_EXIT | |---|---|---|---|---| | 497,739 | 19 | 192 | 497,950 | 0 | Those failed tests are all stress_gemm with ILP tests, which requires appropriate ILP library link, so after I rebuilt the library, I reran those failed tests as below, and all the tests PASSED. `./rocblas-test --gtest_filter=_/gemm.*stress*ILP64*f16_r*:_/gemm_ex.*stress*ILP64*f16_r*:_/gemm_ex.*stress*ILP64*bf16_r*` | PASS | FAIL | SKIP | total | TEST_EXIT | |---|---|---|---|---| | 23 | 0 | 0 | 23 | 0 | ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
|
@evedovelli it looks like get_solutions only, are all the launches guarded by can_solve checks now? |
do we need to run the TF32 get_solutions test for MI300A? @TorreZuk @evedovelli? |
For gemm_ex, we are iterating through all the solutions and trying to solve them. Some of them are not solvable (e.g. only works with alpha 1 but our problem requires another value for alpha) so we just ignore them, but for the ones that provide a result, we do validate those against the golden result. |
…BLASLt defaults and the xf32 math mode fix
deff505 to
5b817e2
Compare
5b817e2 to
deff505
Compare
… 228-CU MI300A and the 304-CU MI300X
TheRock Submodule Bump ActivityNewest first
|
…x950 (#13901) ## Motivation This PR extends hipBLASLt backend support for F32 (SGEMM) operations on both gfx942 and gfx950 architectures. Previously: - **gfx950**: F32 was explicitly disabled and defaulted to Tensile - **gfx942**: F32 was only enabled on the 228-CU MI300A configuration This change enables F32 GEMM to leverage the hipBLASLt backend for improved performance on: - All gfx950 devices (MI350 series) - Both MI300A (228 CU) and MI300X (304 CU) gfx942 devices This follows the pattern established in #13045, which defaulted various data types to hipBLASLt on specific architectures. ## Technical Details ### Changes to `tensile_host.cpp` The `useHipBLASLt()` function has been modified: 1. **gfx950 architecture (lines 1295-1301)**: - **Before**: Early return `false` for F32 types, forcing Tensile backend - **After**: All types including F32 now use hipBLASLt via `prob.handle->tryHipBLASLt()` - Removed the TODO comment about SGEMM support 2. **gfx942 architecture (lines 1303-1333)**: - **Before**: F32 was in the `only_MI300A` category, restricting it to 228 CU devices only - **After**: F32 moved to the `MI300A_or_MI300X` category alongside fp16, bf16, and fp64 - This enables F32 on both 228 CU (MI300A) and 304 CU (MI300X) configurations - Complex types (complex-fp64, complex-fp32) remain in the `only_MI300A` category ### Changes to `CHANGELOG.md` Updated the rocBLAS 5.8.0 changelog to reflect the new defaults: - **gfx950**: Updated from "every input type except real `f32_r`" to "all input types" - **gfx942**: Added `f32_r` to the list of types enabled on both MI300A and MI300X ### Key Implementation Details ```cpp // gfx950: Now enables hipBLASLt for all types including F32 if(arch == 950) { return prob.handle->tryHipBLASLt(); } // gfx942: F32 now included in MI300A_or_MI300X category constexpr bool MI300A_or_MI300X = sizeof(Ti) == 2 || std::is_same<Ti, double>::value || std::is_same<Ti, float>::value; ``` ## Impact ### Architectures Affected - **gfx950 (MI350 series)**: All devices now use hipBLASLt for F32 - **gfx942 (MI300 series)**: - MI300A (228 CU): Already supported, no change - MI300X (304 CU): Now enabled for F32 (previously Tensile only) ### Backward Compatibility - The `ROCBLAS_USE_HIPBLASLT` environment variable continues to force or disable the hipBLASLt backend, providing an override mechanism Co-authored-by: Madhusoodhanan Prabha <amadhuso@ctr2-alola-ctrl-01.amd.com> Co-authored-by: Claude Sonnet 4 <noreply@anthropic.com>
JIRA ID : https://amd-hub.atlassian.net/browse/AIHPBLAS-3404
JIRA ID : https://amd-hub.atlassian.net/browse/AIHPBLAS-3402
JIRA ID : https://amd-hub.atlassian.net/browse/AIROCBLAS-1334
JIRA ID : https://amd-hub.atlassian.net/browse/AIHPBLAS-3403
JIRA ID : https://amd-hub.atlassian.net/browse/AIHPBLAS-3400
Motivation
Follow-up to #12724, which defaulted DGEMM to hipBLASLt on the 228-CU MI300A. This extends the same gate to ZGEMM, fp16 and bf16 so they also default to hipBLASLt on MI300A, while MI300X and every other gfx942
variant stay on Tensile for all types.
Technical Details
useHipBLASLt()(tensile_host.cpp) previously let onlydoublethrough thearch == 942branch; the predicate is widened tosizeof(Ti) == 2 || double || rocblas_double_complex, covering fp16, bf16, DGEMM and ZGEMM. ThemultiProcessorCount != 228guard is unchanged, so MI300X is unaffected, andisDefaultHipBLASLtArch()already lists942.Also plumbs the handle's math mode into the hipBLASLt backend (HIPBLAS_COMPUTE_32F_FAST_TF32), mirroring the Tensile path's setF32XdlMathOp(XFloat32). Without it, flipping SGEMM to hipBLASLt would silently ignore rocblas_set_math_mode(xf32) — which is what gemm_ex_get_solutions_alpha_beta caught.
Test Plan
ROCBLAS_USE_HIPBLASLT=1 ./rocblas-test --gtest_filter='blas3_tensilef64_c*-known_bug:graph_test'ROCBLAS_USE_HIPBLASLT=1 ./rocblas-test --gtest_filter='blas3_tensilef16_r*-known_bug:graph_test'Test Result
ZGEMM — logic validation, gtest and batch coverage all clean:
run_tensile_logic_check.pyover the nine new filesrocblas-testZGEMMROCBLAS_USE_HIPBLASLT=1 ./rocblas-test --gtest_filter='blas3_tensilef32_c*-known_bug:graph_test'
30,326 ran / 30,326 passed / 0 failed, identical in both arms.
fp16 — gtest clean, and at or above Tensile parity in all configs: