Skip to content

Add hipblasLt implementation for batched gemm to improve performance for CDNA3 only #16457

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.

Sign up for GitHub

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

Open
wants to merge 2 commits into
base: master
Choose a base branch
from
+950 −0
Open

Add hipblasLt implementation for batched gemm to improve performance for CDNA3 only #16457

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
c2429d4
Make a fix to hipblaslt
peizhang56 Sep 26, 2025
71100d1
Fix code review comments
peizhang56 Sep 29, 2025
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
30 changes: 30 additions & 0 deletions ggml/src/ggml-cuda/ggml-cuda.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -54,6 +54,10 @@
#include "ggml-cuda/pad_reflect_1d.cuh"
#include "ggml.h"

#if defined(GGML_USE_HIP)
#include "ggml-cuda/vendors/hipblaslt_batched_gemm.cuh"
#endif // defined(GGML_USE_HIP)

#include <algorithm>
#include <array>
#include <atomic>
Expand Down Expand Up @@ -535,6 +539,12 @@ ggml_backend_cuda_context::~ggml_backend_cuda_context() {
CUDA_CHECK(cudaEventDestroy(copy_event));
}
for (int i = 0; i < GGML_CUDA_MAX_DEVICES; ++i) {
#if defined(GGML_USE_HIP)
const int cc = ggml_cuda_info().devices[i].cc;
if(ggml::vendors::getHipblasltBatchedGemmEnvVal() != 0 && GGML_CUDA_CC_IS_CDNA3(cc)){
ggml::vendors::HipblasltBatchedGemmSingleton::getInstance().hipblasltDestroy(i);
}
#endif // defined(GGML_USE_HIP)
for (int j = 0; j < GGML_CUDA_MAX_STREAMS; ++j) {
if (streams[i][j] != nullptr) {
CUDA_CHECK(cudaStreamDestroy(streams[i][j]));
Expand Down Expand Up @@ -1839,6 +1849,15 @@ static void ggml_cuda_mul_mat_batched_cublas_impl(ggml_backend_cuda_context & ct
cudaStream_t main_stream = ctx.stream();
CUBLAS_CHECK(cublasSetStream(ctx.cublas_handle(), main_stream));

#if defined(GGML_USE_HIP)
int cur_dev_id = ggml_cuda_get_device();
const int cur_dev_cc = ggml_cuda_info().devices[cur_dev_id].cc;
if(ggml::vendors::getHipblasltBatchedGemmEnvVal() != 0 && GGML_CUDA_CC_IS_CDNA3(cur_dev_cc)){
auto& instance = ggml::vendors::HipblasltBatchedGemmSingleton::getInstance();
instance.hipblasLtSetStream(ctx.device, main_stream);
}
#endif // defined(GGML_USE_HIP)

float * dst_ddf = (float *) dst->data;
const size_t ts_src1 = ggml_type_size(src1->type);
GGML_ASSERT(nb10 == ts_src1);
Expand Down Expand Up @@ -1962,6 +1981,16 @@ static void ggml_cuda_mul_mat_batched_cublas_impl(ggml_backend_cuda_context & ct

CUDA_CHECK(cudaGetLastError());

#if defined(GGML_USE_HIP)
CUBLAS_CHECK(ggml::vendors::hipblasGemmBatchedEx(ctx.cublas_handle(), CUBLAS_OP_T, CUBLAS_OP_N,
ne01, ne11, ne10,
alpha, (const void **) (ptrs_src.get() + 0*ne23), cu_data_type_a, nb01/nb00,
(const void **) (ptrs_src.get() + 1*ne23), cu_data_type_b, s11,
beta, ( void **) (ptrs_dst.get() + 0*ne23), cu_data_type, ne0,
ne23,
cu_compute_type,
CUBLAS_GEMM_DEFAULT_TENSOR_OP));
#else
CUBLAS_CHECK(
cublasGemmBatchedEx(ctx.cublas_handle(), CUBLAS_OP_T, CUBLAS_OP_N,
ne01, ne11, ne10,
Expand All @@ -1971,6 +2000,7 @@ static void ggml_cuda_mul_mat_batched_cublas_impl(ggml_backend_cuda_context & ct
ne23,
cu_compute_type,
CUBLAS_GEMM_DEFAULT_TENSOR_OP));
#endif // defined(GGML_USE_HIP)
}

// Convert output back to F32 if needed
Expand Down
Loading