Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
29 changes: 17 additions & 12 deletions transformer_engine/common/gemm/rocm_gemm.cu
Original file line number Diff line number Diff line change
Expand Up @@ -1999,18 +1999,23 @@ 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, gemm_stream);
} else {
use_hipkittens = 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, gemm_stream);
}
if (!use_hipkittens) {
if (is_mxfp8) {
NVTE_CHECK(inputBias->data.dptr == nullptr,
"hipBLASLt MXFP8 GEMM does not support bias");
}
#endif
// FIXME(https://amd-hub.atlassian.net/browse/ROCM-26110): Remove this workaround once hipBLASLt supports NN/NT
// layouts for MXFP8 on gfx1250.
Expand Down
2 changes: 1 addition & 1 deletion transformer_engine/jax/cpp_extensions/gemm.py
Original file line number Diff line number Diff line change
Expand Up @@ -653,7 +653,7 @@ def _dims_are_consecutive(dims):
n = reduce(operator.mul, rhs_non_contracting_shape)
k = lhs_contracting_size
layout = ("T" if lhs_is_transposed else "N") + ("T" if rhs_is_transposed else "N")
workspace_size = _hipkittens_workspace_bytes(m, n, k, layout)
workspace_size = max(_hipkittens_workspace_bytes(m, n, k, layout), get_cublas_workspace_size_bytes())
else:
# Declare cuBLAS workspace
workspace_size = get_cublas_workspace_size_bytes()
Expand Down
Loading