[hipBLASLt] Fix JIRA ROCM-3139: use full FP32 on gfx942 for FAST_TF32 requests#8868
Draft
sajandhy wants to merge 1 commit into
Draft
[hipBLASLt] Fix JIRA ROCM-3139: use full FP32 on gfx942 for FAST_TF32 requests#8868sajandhy wants to merge 1 commit into
sajandhy wants to merge 1 commit into
Conversation
Downgrade HIPBLAS_COMPUTE_32F_FAST_TF32 to HIPBLAS_COMPUTE_32F when XF32 is unsupported on gfx942, and add a smoke gtest that compares the two compute modes on MI300-class hardware. Co-authored-by: Cursor <cursoragent@cursor.com>
sajandhy
changed the title
❌ PR Check — Action Required
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
sajandhy
changed the title
Codecov Report❌ Patch coverage is ❌ Your project status has failed because the head coverage (76.92%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage. Additional details and impacted files@@ Coverage Diff @@
## develop #8868 +/- ##
===========================================
- Coverage 71.34% 71.34% -0.00%
===========================================
Files 2613 2613
Lines 408997 409006 +9
Branches 61114 61115 +1
===========================================
+ Hits 291772 291777 +5
- Misses 95826 95829 +3
- Partials 21399 21400 +1
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
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
2 participants
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.
Downgrade HIPBLAS_COMPUTE_32F_FAST_TF32 to HIPBLAS_COMPUTE_32F when XF32 is unsupported on gfx942, and add a smoke gtest that compares the two compute modes on MI300-class hardware.
Motivation
On gfx942 (MI300-class), HIPBLAS_COMPUTE_32F_FAST_TF32 was silently routing FP32 GEMM through a lower-precision path, producing results that diverged from HIPBLAS_COMPUTE_32F. This fixes ROCM-3139 by ensuring FAST_TF32 requests use full FP32 compute on gfx942 when XF32 is not supported.
Technical Details
Add rocblaslt_internal_supports_xf32_compute() returning false on gfx942.
In hipblasLtMatmulDescCreate(), downgrade HIPBLAS_COMPUTE_32F_FAST_TF32 → HIPBLAS_COMPUTE_32F when XF32 is unsupported.
Gate XF32 heuristic selection in tensile_host.cpp on the same capability check.
Add HipblasltTf32Precision.smoke_FastTf32MatchesFullFp32OnGfx942 gtest: runs the same FP32 NN GEMM with both compute types and requires outputs to match within 1e-4. The smoke_ prefix ensures TheRock PR CI (--gtest_filter=smoke) runs it on gfx942; skipped on other arches.
Link GTest::gmock in clients/CMakeLists.txt (required by existing matrix_transform_gtest.cpp matchers).
Files changed: hipblaslt.cpp, rocblaslt_auxiliary.cpp, rocblaslt-auxiliary.h, tensile_host.cpp, hipblaslt_tf32_precision_gtest.cpp, client CMakeLists.
Test Plan
Build host lib + clients only: invoke build -ca gfx942 -c -d -t
Point device libs at in-tree install: HIPBLASLT_TENSILE_LIBPATH=/lib/hipblaslt/library/gfx942
Run targeted gtest: hipblaslt-test --gtest_filter='HipblasltTf32Precision'
CI: TheRock smoke gtest lane on gfx942 hardware
Test Result
HipblasltTf32Precision.smoke_FastTf32MatchesFullFp32OnGfx942 PASS on gfx942 (MI308X) with patched libhipblaslt.so
Pre-fix stock library reproduced the bug (max diff ~0.37); post-fix max diff 0.0
Submission Checklist