From 1a0843c4937310d011b5e5a19fcabf04ddeadac3 Mon Sep 17 00:00:00 2001 From: Georgi Gerganov Date: Fri, 27 Oct 2023 13:05:33 +0300 Subject: [PATCH 1/3] cuda : utilize tensor cores with multiple GPU devices --- ggml-cuda.cu | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/ggml-cuda.cu b/ggml-cuda.cu index 1ba951f688d82..922fdb199035a 100644 --- a/ggml-cuda.cu +++ b/ggml-cuda.cu @@ -6361,7 +6361,7 @@ inline void ggml_cuda_op_mul_mat_cublas( const int compute_capability = g_compute_capabilities[id]; - if (compute_capability >= CC_VOLTA && (src0->type == GGML_TYPE_F16 || ggml_is_quantized(src0->type)) && ggml_is_contiguous(src0) && row_diff == src0->ne[1]) { + if (compute_capability >= CC_VOLTA && (src0->type == GGML_TYPE_F16 || ggml_is_quantized(src0->type)) && ggml_is_contiguous(src0)) { // convert src0 and src1 to fp16, multiply as fp16, convert dst to fp32 half * src0_as_f16 = nullptr; size_t src0_as = 0; @@ -6386,7 +6386,7 @@ inline void ggml_cuda_op_mul_mat_cublas( const half * src1_ptr = src1->type == GGML_TYPE_F16 ? (const half *) src1_ddq_i : src1_as_f16; size_t dst_as = 0; - half * dst_f16 = (half *) ggml_cuda_pool_malloc(row_diff*src1_ncols * sizeof(half), &dst_as); + half * dst_f16 = (half *) ggml_cuda_pool_malloc(ldc*src1_ncols * sizeof(half), &dst_as); const half alpha_f16 = 1.0f; const half beta_f16 = 0.0f; @@ -6402,7 +6402,7 @@ inline void ggml_cuda_op_mul_mat_cublas( CUBLAS_GEMM_DEFAULT_TENSOR_OP)); const to_fp32_cuda_t to_fp32_cuda = ggml_get_to_fp32_cuda(GGML_TYPE_F16); - to_fp32_cuda(dst_f16, dst_dd_i, row_diff*src1_ncols, stream); + to_fp32_cuda(dst_f16, dst_dd_i, ldc*src1_ncols, stream); ggml_cuda_pool_free(dst_f16, dst_as); From 706ff4c2e0d4508a2d085a9fcb2ea2d573666040 Mon Sep 17 00:00:00 2001 From: Georgi Gerganov Date: Fri, 27 Oct 2023 22:17:47 +0300 Subject: [PATCH 2/3] cuda : try to fix main device write --- ggml-cuda.cu | 51 ++++++++++++++++++++++++++++++++++----------------- 1 file changed, 34 insertions(+), 17 deletions(-) diff --git a/ggml-cuda.cu b/ggml-cuda.cu index 922fdb199035a..f99126bbd32c2 100644 --- a/ggml-cuda.cu +++ b/ggml-cuda.cu @@ -6355,11 +6355,13 @@ inline void ggml_cuda_op_mul_mat_cublas( int id; CUDA_CHECK(cudaGetDevice(&id)); + const int compute_capability = g_compute_capabilities[id]; + // the main device has a larger memory buffer to hold the results from all GPUs // ldc == nrows of the matrix that cuBLAS writes into - int ldc = dst->backend == GGML_BACKEND_GPU && id == g_main_device ? ne0 : row_diff; + const int ldc = dst->backend == GGML_BACKEND_GPU && id == g_main_device ? ne0 : row_diff; - const int compute_capability = g_compute_capabilities[id]; + const bool is_split = row_diff != src0->ne[1]; if (compute_capability >= CC_VOLTA && (src0->type == GGML_TYPE_F16 || ggml_is_quantized(src0->type)) && ggml_is_contiguous(src0)) { // convert src0 and src1 to fp16, multiply as fp16, convert dst to fp32 @@ -6385,26 +6387,41 @@ inline void ggml_cuda_op_mul_mat_cublas( } const half * src1_ptr = src1->type == GGML_TYPE_F16 ? (const half *) src1_ddq_i : src1_as_f16; - size_t dst_as = 0; - half * dst_f16 = (half *) ggml_cuda_pool_malloc(ldc*src1_ncols * sizeof(half), &dst_as); + if (!is_split) { + const half alpha = 1.0f; + const half beta = 0.0f; - const half alpha_f16 = 1.0f; - const half beta_f16 = 0.0f; + size_t dst_as = 0; + half * dst_f16 = (half *) ggml_cuda_pool_malloc(row_diff*src1_ncols * sizeof(half), &dst_as); - CUBLAS_CHECK(cublasSetStream(g_cublas_handles[id], stream)); - CUBLAS_CHECK( - cublasGemmEx(g_cublas_handles[id], CUBLAS_OP_T, CUBLAS_OP_N, - row_diff, src1_ncols, ne10, - &alpha_f16, src0_ptr, CUDA_R_16F, ne00, + CUBLAS_CHECK(cublasSetStream(g_cublas_handles[id], stream)); + CUBLAS_CHECK( + cublasGemmEx(g_cublas_handles[id], CUBLAS_OP_T, CUBLAS_OP_N, + row_diff, src1_ncols, ne10, + &alpha, src0_ptr, CUDA_R_16F, ne00, src1_ptr, CUDA_R_16F, ne10, - &beta_f16, dst_f16, CUDA_R_16F, ldc, - CUBLAS_COMPUTE_16F, - CUBLAS_GEMM_DEFAULT_TENSOR_OP)); + &beta, dst_f16, CUDA_R_16F, ldc, + CUBLAS_COMPUTE_16F, + CUBLAS_GEMM_DEFAULT_TENSOR_OP)); - const to_fp32_cuda_t to_fp32_cuda = ggml_get_to_fp32_cuda(GGML_TYPE_F16); - to_fp32_cuda(dst_f16, dst_dd_i, ldc*src1_ncols, stream); + const to_fp32_cuda_t to_fp32_cuda = ggml_get_to_fp32_cuda(GGML_TYPE_F16); + to_fp32_cuda(dst_f16, dst_dd_i, row_diff*src1_ncols, stream); - ggml_cuda_pool_free(dst_f16, dst_as); + ggml_cuda_pool_free(dst_f16, dst_as); + } else { + const float alpha = 1.0f; + const float beta = 0.0f; + + CUBLAS_CHECK(cublasSetStream(g_cublas_handles[id], stream)); + CUBLAS_CHECK( + cublasGemmEx(g_cublas_handles[id], CUBLAS_OP_T, CUBLAS_OP_N, + row_diff, src1_ncols, ne10, + &alpha, src0_ptr, CUDA_R_16F, ne00, + src1_ptr, CUDA_R_16F, ne10, + &beta, dst_dd_i, CUDA_R_32F, ldc, + CUBLAS_COMPUTE_32F, + CUBLAS_GEMM_DEFAULT_TENSOR_OP)); + } if (src0_as != 0) { ggml_cuda_pool_free(src0_as_f16, src0_as); From cd3e20fb50d73080dd6d72ad68e401cc42cd53cf Mon Sep 17 00:00:00 2001 From: Georgi Gerganov Date: Fri, 27 Oct 2023 23:11:50 +0300 Subject: [PATCH 3/3] cuda : fix multi-gpu with tensor cores --- ggml-cuda.cu | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/ggml-cuda.cu b/ggml-cuda.cu index f99126bbd32c2..9c7cf35715aa9 100644 --- a/ggml-cuda.cu +++ b/ggml-cuda.cu @@ -461,7 +461,12 @@ static_assert(K_QUANTS_PER_ITERATION == 1 || K_QUANTS_PER_ITERATION == 2, "K_QUA #define GGML_CUDA_PEER_MAX_BATCH_SIZE 128 #endif // GGML_CUDA_PEER_MAX_BATCH_SIZE +#ifdef GGML_CUDA_FORCE_MMQ #define MUL_MAT_SRC1_COL_STRIDE 128 +#else +// with tensor cores, we copy the entire hidden state to the devices in one go +#define MUL_MAT_SRC1_COL_STRIDE 4096 +#endif #define MAX_STREAMS 8 static cudaStream_t g_cudaStreams[GGML_CUDA_MAX_DEVICES][MAX_STREAMS] = { nullptr };