Enable hipBLASLt GEMM for gfx115x#5082
Merged
Merged
Conversation
Add support for gfx115x so gfx1150/1151/1152/1153 (Strix Halo/Point, RDNA3.5) are supported in hipBLASLt. hipblaslt_supported_impl() gated hipBLASLt on gfx110x and gfx120x but not gfx115x, so on these parts the lowering pass fell back to rocBLAS for every GEMM. rocBLAS lacks BF16/INT8 Tensile solutions for these parts, so quantized dot ops returned rocblas_status_not_implemented at runtime while FP16/FP32 worked. hipBLASLt ships tuned BF16/INT8 kernels for the gfx115 family, so allow it to use the hipBLASLt path. Signed-off-by: Johny Shaik Mohammad <[email protected]>
Contributor
|
Thank you for your contribution! Since this is an external pull request, a maintainer must review PR and add the "ok-to-test" label if it is approved for testing. |
Contributor
There was a problem hiding this comment.
Pull request overview
Enables the hipBLASLt GEMM path on RDNA3.5 (gfx115x) GPUs by extending the device-family gate used by the GPU backend, preventing fallback to rocBLAS for GEMM on those parts.
Changes:
- Extend
hipblaslt_supported_impl()to treatgfx115*as hipBLASLt-supported (alongside gfx110x/gfx120x and select gfx94x/gfx95x).
Comment on lines
142
to
+145
| return (gfx_name == "gfx90a" or (starts_with(gfx_name, "gfx94") and gfx_name >= "gfx942") or | ||
| (starts_with(gfx_name, "gfx95") and gfx_name >= "gfx950") or | ||
| starts_with(gfx_name, "gfx110") or starts_with(gfx_name, "gfx120")); | ||
| starts_with(gfx_name, "gfx110") or starts_with(gfx_name, "gfx115") or | ||
| starts_with(gfx_name, "gfx120")); |
kahmed10
approved these changes
Jul 21, 2026
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.
Summary
Adds
gfx115tohipblaslt_supported_impl()so gfx1150/1151/1152/1153 (Strix Halo/Point, RDNA3.5) use the hipBLASLt GEMM path.Previously
hipblaslt_supported_impl()gated hipBLASLt on gfx110x and gfx120x but not gfx115x, so on these parts the lowering pass fell back to rocBLAS for every GEMM. rocBLAS lacks BF16/INT8 Tensile solutions for these parts, so quantized dot ops returnedrocblas_status_not_implementedat runtime while FP16/FP32 worked. hipBLASLt ships tuned BF16/INT8 kernels for the gfx115 family, so this allows it to use the hipBLASLt path.Change
src/targets/gpu/device_name.cpp: addstarts_with(gfx_name, "gfx115")to the hipBLASLt-supported predicate.Test plan
MIGRAPHX_USE_HIPBLASLTenabled and run quantized (BF16/INT8) dot ops on a gfx115x device; confirm they no longer fall back to rocBLAS / returnrocblas_status_not_implemented.Made with Cursor