-
Notifications
You must be signed in to change notification settings - Fork 34
hipblaslt Fallback hotfix #652
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Changes from 1 commit
File filter
Filter by extension
Conversations
Jump to
Diff view
Diff view
There are no files selected for viewing
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -1985,7 +1985,7 @@ void cublas_gemm(const Tensor *inputA, const Tensor *inputB, Tensor *outputD, | |
|
|
||
| #ifdef USE_HIPKITTENS_GEMM | ||
|
|
||
| bool use_hipkittens = false; | ||
| bool use_hipkittens = false, hipkittens_gemm_complete = false; | ||
| if (is_mxfp8) { | ||
| bool is_gfx950 = (cuda::sm_arch() == 95); | ||
| bool force_hipblaslt = false; | ||
|
|
@@ -2001,18 +2001,19 @@ void cublas_gemm(const Tensor *inputA, const Tensor *inputB, Tensor *outputD, | |
| if (use_hipkittens) { | ||
| auto param = CanonicalizeGemmInput(*inputA, transa, *inputB, transb, m, n, k); | ||
|
|
||
| kittens_mxfp8_gemm(param.A, param.B, outputD->data.dptr, | ||
| param.A_scale_inv, param.B_scale_inv, | ||
| m, n, k, is_transa, is_transb, | ||
| static_cast<int>(param.Atype), | ||
| static_cast<int>(param.Btype), | ||
| inputBias->data.dptr, | ||
| static_cast<int>(inputBias->data.dtype), | ||
| outputPreGelu->data.dptr, | ||
| static_cast<int>(outputD->data.dtype), | ||
| static_cast<int>(outputPreGelu->data.dtype), | ||
| workspace, workspaceSize, s); | ||
| } else { | ||
| hipkittens_gemm_complete = kittens_mxfp8_gemm(param.A, param.B, outputD->data.dptr, | ||
|
wenchenvincent marked this conversation as resolved.
Outdated
|
||
| param.A_scale_inv, param.B_scale_inv, | ||
| m, n, k, is_transa, is_transb, | ||
| static_cast<int>(param.Atype), | ||
| static_cast<int>(param.Btype), | ||
| inputBias->data.dptr, | ||
| static_cast<int>(inputBias->data.dtype), | ||
| outputPreGelu->data.dptr, | ||
| static_cast<int>(outputD->data.dtype), | ||
| static_cast<int>(outputPreGelu->data.dtype), | ||
| workspace, workspaceSize, s); | ||
| } | ||
| if (!use_hipkittens || !hipkittens_gemm_complete) { | ||
|
Collaborator
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. nit: no need to check use_hipkittens - hipkittens_gemm_complete can only be True if use_hipkittens is True. Only use_hipkittens can be sufficient (nohipkittens_gemm_complete) , if assign it at line 2004
Contributor
Author
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. I have simplified the check as you mentioned. |
||
| #endif | ||
| // FIXME(https://amd-hub.atlassian.net/browse/ROCM-26110): Remove this workaround once hipBLASLt supports NN/NT | ||
| // layouts for MXFP8 on gfx1250. | ||
|
|
||
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
This is unused variable
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Sorry, missed that. Should be fixed now.