diff options
author | slaren <2141330+slaren@users.noreply.github.com> | 2023-05-01 18:11:07 +0200 |
---|---|---|
committer | GitHub <noreply@github.com> | 2023-05-01 18:11:07 +0200 |
commit | 58b367c2d757c0ea12aec672382462b42204c724 (patch) | |
tree | b2fa89daf71c08788c44e3fb9abf1747ec8ee65d /ggml-cuda.h | |
parent | ea3a0ad6b6b5ca4693b94acd4cb32e2803f66fae (diff) |
cuBLAS: refactor and optimize f16 mat mul performance (#1259)
* cuBLAS: refactor, convert fp16 to fp32 on device
* cuBLAS: use multiple streams, choose smartly between mul_mat_q and mul_mat_f16
* fix build
* cuBLAS: update block_q5_1
Diffstat (limited to 'ggml-cuda.h')
-rw-r--r-- | ggml-cuda.h | 45 |
1 files changed, 5 insertions, 40 deletions
diff --git a/ggml-cuda.h b/ggml-cuda.h index 36782d9..f7d6a8b 100644 --- a/ggml-cuda.h +++ b/ggml-cuda.h @@ -1,54 +1,19 @@ -#include <cublas_v2.h> -#include <cuda_runtime.h> #include "ggml.h" #ifdef __cplusplus extern "C" { #endif -#define CUDA_CHECK(err) \ - do { \ - cudaError_t err_ = (err); \ - if (err_ != cudaSuccess) { \ - fprintf(stderr, "CUDA error %d at %s:%d: %s\n", err_, __FILE__, __LINE__, \ - cudaGetErrorString(err_)); \ - exit(1); \ - } \ - } while (0) - -#define CUBLAS_CHECK(err) \ - do { \ - cublasStatus_t err_ = (err); \ - if (err_ != CUBLAS_STATUS_SUCCESS) { \ - fprintf(stderr, "cuBLAS error %d at %s:%d\n", err_, __FILE__, __LINE__); \ - exit(1); \ - } \ - } while (0) +void ggml_init_cublas(void); -extern cublasHandle_t g_cublasH; -extern cudaStream_t g_cudaStream; -extern cudaStream_t g_cudaStream2; -extern cudaEvent_t g_cudaEvent; +bool ggml_cuda_can_mul_mat(const struct ggml_tensor * src0, const struct ggml_tensor * src1, struct ggml_tensor * dst); +size_t ggml_cuda_mul_mat_get_wsize(const struct ggml_tensor * src0, const struct ggml_tensor * src1, struct ggml_tensor * dst); +void ggml_cuda_mul_mat(const struct ggml_tensor * src0, const struct ggml_tensor * src1, struct ggml_tensor * dst, void * wdata, size_t wsize); -void ggml_init_cublas(void); +// TODO: export these with GGML_API void * ggml_cuda_host_malloc(size_t size); void ggml_cuda_host_free(void * ptr); -void * ggml_cuda_pool_malloc(size_t size, size_t * actual_size); -void ggml_cuda_pool_free(void * ptr, size_t size); - -void dequantize_row_q4_0_cuda(const void * vx, float * y, int k, cudaStream_t stream); -void dequantize_row_q4_1_cuda(const void * vx, float * y, int k, cudaStream_t stream); -void dequantize_row_q4_2_cuda(const void * vx, float * y, int k, cudaStream_t stream); -void dequantize_row_q5_0_cuda(const void * vx, float * y, int k, cudaStream_t stream); -void dequantize_row_q5_1_cuda(const void * vx, float * y, int k, cudaStream_t stream); -void dequantize_row_q8_0_cuda(const void * vx, float * y, int k, cudaStream_t stream); - -cudaError_t ggml_cuda_h2d_tensor_2d(void * dst, const struct ggml_tensor * src, uint64_t i3, uint64_t i2, cudaStream_t stream); - -typedef void (*dequantize_row_q_cuda_t)(const void * x, float * y, int k, cudaStream_t stream); -dequantize_row_q_cuda_t ggml_get_dequantize_row_q_cuda(enum ggml_type type); - #ifdef __cplusplus } #endif |