From 1b16b8c90dc3f555db70df3f72af00d677f84ffe Mon Sep 17 00:00:00 2001 From: 0cc4m Date: Sun, 23 Apr 2023 09:59:45 +0200 Subject: [PATCH] Move CLBlast implementation to separate file Add buffer reuse code (adapted from slaren's cuda implementation) --- Makefile | 11 ++- ggml-opencl.cpp | 239 ++++++++++++++++++++++++++++++++++++++++++++++++ ggml-opencl.h | 31 +++++++ ggml.c | 207 +---------------------------------------- 4 files changed, 281 insertions(+), 207 deletions(-) create mode 100644 ggml-opencl.cpp create mode 100644 ggml-opencl.h diff --git a/Makefile b/Makefile index 34824d690..ca945fa99 100644 --- a/Makefile +++ b/Makefile @@ -34,8 +34,8 @@ endif # # keep standard at C11 and C++11 -CFLAGS = -I. -O3 -DNDEBUG -std=c11 -fPIC -CXXFLAGS = -I. -I./examples -O3 -DNDEBUG -std=c++11 -fPIC +CFLAGS = -I. -O3 -DNODEBUG -std=c11 -fPIC +CXXFLAGS = -I. -I./examples -O3 -DNODEBUG -std=c++11 -fPIC LDFLAGS = # warnings @@ -105,8 +105,8 @@ ifdef LLAMA_OPENBLAS LDFLAGS += -lopenblas endif ifdef LLAMA_CUBLAS - CFLAGS += -DGGML_USE_CUBLAS -I/usr/local/cuda/include - LDFLAGS += -lcublas -lculibos -lcudart -lcublasLt -lpthread -ldl -lrt -L/usr/local/cuda/lib64 + CFLAGS += -DGGML_USE_CUBLAS -I/usr/local/cuda/include -I/opt/cuda/include + LDFLAGS += -lcublas -lculibos -lcudart -lcublasLt -lpthread -ldl -lrt -L/usr/local/cuda/lib64 -L/opt/cuda/lib64 OBJS += ggml-cuda.o NVCC = nvcc NVCCFLAGS = --forward-unknown-to-host-compiler -arch=native @@ -116,6 +116,9 @@ endif ifdef LLAMA_CLBLAST CFLAGS += -DGGML_USE_CLBLAST LDFLAGS += -lclblast -lOpenCL + OBJS += ggml-opencl.o +ggml-opencl.o: ggml-opencl.cpp ggml-opencl.h + $(CXX) $(CXXFLAGS) -c $< -o $@ endif ifdef LLAMA_GPROF CFLAGS += -pg diff --git a/ggml-opencl.cpp b/ggml-opencl.cpp new file mode 100644 index 000000000..fa462ff8e --- /dev/null +++ b/ggml-opencl.cpp @@ -0,0 +1,239 @@ +#include "ggml-opencl.h" + +#include +#include +#include + +#include + +struct scoped_spin_lock { + std::atomic_flag& lock; + scoped_spin_lock(std::atomic_flag& lock) : lock(lock) { + while (lock.test_and_set(std::memory_order_acquire)) { + ; // spin + } + } + ~scoped_spin_lock() { + lock.clear(std::memory_order_release); + } + scoped_spin_lock(const scoped_spin_lock&) = delete; + scoped_spin_lock& operator=(const scoped_spin_lock&) = delete; +}; + +struct cl_buffer { + cl_mem mem; + size_t size = 0; +}; + +cl_platform_id platform; +cl_device_id device; +cl_context context; +cl_command_queue queue; +cl_program program; +cl_kernel kernel_q4_0, kernel_q4_1; +bool cl_initialized = false; +size_t cl_size_a = 0, cl_size_b = 0, cl_size_c = 0; + +static cl_buffer g_cl_buffer_pool[MAX_CL_BUFFERS]; +static std::atomic_flag g_cl_pool_lock = ATOMIC_FLAG_INIT; + +cl_mem ggml_cl_pool_malloc(size_t size, size_t * actual_size) { + scoped_spin_lock lock(g_cl_pool_lock); + + for (int i = 0; i < MAX_CL_BUFFERS; ++i) { + cl_buffer& b = g_cl_buffer_pool[i]; + if (b.size >= size && b.size != 0) { + cl_mem mem = b.mem; + *actual_size = b.size; + b.size = 0; + return mem; + } + } + cl_int err; + cl_mem mem = clCreateBuffer(context, 0, size, NULL, &err); + *actual_size = size; + CL_CHECK(err, "clCreateBuffer"); + return mem; +} + +void ggml_cl_pool_free(cl_mem mem, size_t size) { + scoped_spin_lock lock(g_cl_pool_lock); + + for (int i = 0; i < MAX_CL_BUFFERS; ++i) { + cl_buffer& b = g_cl_buffer_pool[i]; + if (b.size == 0) { + b.mem = mem; + b.size = size; + return; + } + } + fprintf(stderr, "WARNING: cl buffer pool full, increase MAX_CL_BUFFERS\n"); + clReleaseMemObject(mem); +} + +cl_program build_program_from_source(cl_context ctx, cl_device_id dev, const char* program_buffer) { + cl_program program; + char *program_log; + size_t program_size, log_size; + int err; + + program_size = strlen(program_buffer); + + program = clCreateProgramWithSource(ctx, 1, + (const char**)&program_buffer, &program_size, &err); + if(err < 0) { + fprintf(stderr, "OpenCL error creating program"); + exit(1); + } + + err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL); + if(err < 0) { + + clGetProgramBuildInfo(program, dev, CL_PROGRAM_BUILD_LOG, + 0, NULL, &log_size); + program_log = (char*) malloc(log_size + 1); + program_log[log_size] = '\0'; + clGetProgramBuildInfo(program, dev, CL_PROGRAM_BUILD_LOG, + log_size + 1, program_log, NULL); + printf("%s\n", program_log); + free(program_log); + exit(1); + } + + return program; +} + +void ggml_cl_init() { + cl_int err = 0; + char * KCPP_CLBLAST_PLATFORM = getenv("KCPP_CLBLAST_PLATFORM"); + char * KCPP_CLBLAST_DEVICES = getenv("KCPP_CLBLAST_DEVICES"); + int plat_num = (KCPP_CLBLAST_PLATFORM == NULL ? 0 : atoi(KCPP_CLBLAST_PLATFORM)); + int dev_num = (KCPP_CLBLAST_DEVICES == NULL ? 0 : atoi(KCPP_CLBLAST_DEVICES)); + printf("\nInitializing CLBlast (First Run)..."); + printf("\nAttempting to use: Platform=%d, Device=%d (If invalid, program will crash)\n",plat_num,dev_num); + cl_uint num_platforms; + clGetPlatformIDs(0, NULL, &num_platforms); + cl_platform_id* platforms = (cl_platform_id*)malloc(num_platforms*sizeof(cl_platform_id)); + clGetPlatformIDs(num_platforms, platforms, NULL); + platform = platforms[plat_num]; + char platform_buffer[1024]; + clGetPlatformInfo(platform, CL_PLATFORM_NAME, sizeof(platform_buffer), &platform_buffer, NULL); + cl_uint num_devices; + clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, NULL, &num_devices); + cl_device_id* devices = (cl_device_id*)malloc(num_devices*sizeof(cl_device_id)); + clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, num_devices, devices, NULL); + device = devices[dev_num]; + char device_buffer[1024]; + clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(device_buffer), &device_buffer, NULL); + printf("Using Platform: %s Device: %s\n", platform_buffer, device_buffer); + context = clCreateContext(NULL, 1, &device, NULL, NULL, &err); + if (err != CL_SUCCESS) { + printf("Error creating OpenCL context: %d\n", err); + fflush(stdout); + } + queue = clCreateCommandQueue(context, device, 0, &err); + + if (err != CL_SUCCESS) { + printf("Error creating OpenCL Command Queue: %d\n", err); + fflush(stdout); + } + + free(platforms); + free(devices); + + program = build_program_from_source(context, device, clblast_dequant); + + // Prepare dequantize kernels + kernel_q4_0 = clCreateKernel(program, "dequantize_row_q4_0", &err); + if(err < 0) { + printf("Error creating OpenCL dequantize q4_0 kernel: %d\n", err); + fflush(stdout); + }; + kernel_q4_1 = clCreateKernel(program, "dequantize_row_q4_1", &err); + if(err < 0) { + printf("Error creating OpenCL dequantize q4_1 kernel: %d\n", err); + fflush(stdout); + }; + cl_initialized = true; +} + +void ggml_cl_sgemm_wrapper(const CLBlastLayout order, const CLBlastTranspose trans_a, const CLBlastTranspose trans_b, const int m, const int n, const int k, const float alpha, const void *host_a, const int lda, const float *host_b, const int ldb, const float beta, float *host_c, const int ldc, const int btype) { + cl_int err = 0; + + cl_event events[4]; + events[0] = NULL; + events[1] = NULL; + events[2] = NULL; + events[3] = NULL; + + bool dequant = (btype == 2 || btype == 3); + cl_kernel kernel = btype == 2 ? kernel_q4_0 : kernel_q4_1; + + size_t global = n * k, local = 16, qb_size; + cl_mem cl_buffer_a, cl_buffer_qb, cl_buffer_b, cl_buffer_c; + + size_t buf_size_a, buf_size_qb, buf_size_b, buf_size_c; + + // Prepare buffers + cl_buffer_a = ggml_cl_pool_malloc(m * k * sizeof(float), &buf_size_a); + if (dequant) { + qb_size = global * (sizeof(float) * (btype == 2 ? 1 : 2) + 16) / 32; + cl_buffer_qb = ggml_cl_pool_malloc(qb_size, &buf_size_qb); + } + cl_buffer_b = ggml_cl_pool_malloc(n*k*sizeof(float), &buf_size_b); + cl_buffer_c = ggml_cl_pool_malloc(m*n*sizeof(float), &buf_size_c); + + if (dequant) { + err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &cl_buffer_qb); + err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &cl_buffer_b); + if(err < 0) { + printf("Error setting OpenCL kernel args: %d\n", err); + fflush(stdout); + } + clEnqueueWriteBuffer(queue, cl_buffer_qb, CL_FALSE, 0, qb_size, host_b, 0, NULL, events + 1); + } else { + clEnqueueWriteBuffer(queue, cl_buffer_b, CL_FALSE, 0, n*k*sizeof(float), host_b, 0, NULL, events + 1); + } + + clEnqueueWriteBuffer(queue, cl_buffer_a, CL_FALSE, 0, m*k*sizeof(float), host_a, 0, NULL, events); + clEnqueueWriteBuffer(queue, cl_buffer_c, CL_FALSE, 0, m*n*sizeof(float), host_c, 0, NULL, events + 2); + if (dequant) { + err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, &local, 1, events + 1, events + 3); + if(err < 0) { + printf("Error enqueueing OpenCL dequantize kernel: %d\n", err); + fflush(stdout); + } + } + clWaitForEvents(dequant ? 4 : 3, events); + clReleaseEvent(events[0]); + clReleaseEvent(events[1]); + clReleaseEvent(events[2]); + if (dequant) { + clReleaseEvent(events[3]); + } + + // Call the SGEMM routine. + CLBlastStatusCode status = CLBlastSgemm(order, + trans_a, trans_b, + m, n, k, + alpha, + cl_buffer_a, 0, lda, + cl_buffer_b, 0, ldb, + beta, + cl_buffer_c, 0, ldc, + &queue, events); + + clEnqueueReadBuffer(queue, cl_buffer_c, CL_TRUE, 0, m*n*sizeof(float), host_c, 1, events, events + 1); + + // Wait for completion + clWaitForEvents(2, events); + clReleaseEvent(events[0]); + clReleaseEvent(events[1]); + + ggml_cl_pool_free(cl_buffer_a, buf_size_a); + if (dequant) { + ggml_cl_pool_free(cl_buffer_qb, buf_size_qb); + } + ggml_cl_pool_free(cl_buffer_b, buf_size_b); + ggml_cl_pool_free(cl_buffer_c, buf_size_c); +} diff --git a/ggml-opencl.h b/ggml-opencl.h new file mode 100644 index 000000000..2486845d8 --- /dev/null +++ b/ggml-opencl.h @@ -0,0 +1,31 @@ +#pragma once + +#define CL_TARGET_OPENCL_VERSION 110 +#include +#define MAX_CL_BUFFERS 16 + +#ifdef __cplusplus +extern "C" { +#endif + +// Buffer reuse code adapted from cuda implementation by slaren +#define CL_CHECK(err, name) \ + do { \ + cl_int err_ = (err); \ + if (err_ != CL_SUCCESS) { \ + fprintf(stderr, "OpenCL %s error %d at %s:%d\n", name, err_, __FILE__, __LINE__); \ + exit(1); \ + } \ + } while (0) + +cl_mem ggml_cl_pool_malloc(size_t size, size_t * actual_size); +void ggml_cl_pool_free(cl_mem mem, size_t size); + +cl_program build_program_from_source(cl_context ctx, cl_device_id dev, const char* program_buffer); +void ggml_cl_init(); + +void ggml_cl_sgemm_wrapper(const CLBlastLayout order, const CLBlastTranspose trans_a, const CLBlastTranspose trans_b, const int m, const int n, const int k, const float alpha, const void *host_a, const int lda, const float *host_b, const int ldb, const float beta, float *host_c, const int ldc, const int btype); + +#ifdef __cplusplus +} +#endif diff --git a/ggml.c b/ggml.c index ad4526298..3daff3fe9 100644 --- a/ggml.c +++ b/ggml.c @@ -150,206 +150,7 @@ inline static void* ggml_aligned_malloc(size_t size) { #elif defined(GGML_USE_CUBLAS) #include "ggml-cuda.h" #elif defined(GGML_USE_CLBLAST) -#define CL_TARGET_OPENCL_VERSION 110 -#include -#include -#include - -cl_platform_id platform; -cl_device_id device; -cl_context context; -cl_command_queue queue; -cl_program program; -cl_kernel kernel_q4_0, kernel_q4_1; -bool cl_initialized = false; -size_t cl_size_a = 0, cl_size_b = 0, cl_size_c = 0; - -cl_program build_program_from_source(cl_context ctx, cl_device_id dev, const char* program_buffer) { - cl_program program; - char *program_log; - size_t program_size, log_size; - int err; - - program_size = strlen(program_buffer); - - program = clCreateProgramWithSource(ctx, 1, - (const char**)&program_buffer, &program_size, &err); - if(err < 0) { - perror("OpenCL error creating program"); - exit(1); - } - - err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL); - if(err < 0) { - - clGetProgramBuildInfo(program, dev, CL_PROGRAM_BUILD_LOG, - 0, NULL, &log_size); - program_log = (char*) malloc(log_size + 1); - program_log[log_size] = '\0'; - clGetProgramBuildInfo(program, dev, CL_PROGRAM_BUILD_LOG, - log_size + 1, program_log, NULL); - printf("%s\n", program_log); - free(program_log); - exit(1); - } - - return program; -} - -static void ggml_cl_init() { - cl_int err = 0; - char * KCPP_CLBLAST_PLATFORM = getenv("KCPP_CLBLAST_PLATFORM"); - char * KCPP_CLBLAST_DEVICES = getenv("KCPP_CLBLAST_DEVICES"); - int plat_num = (KCPP_CLBLAST_PLATFORM == NULL ? 0 : atoi(KCPP_CLBLAST_PLATFORM)); - int dev_num = (KCPP_CLBLAST_DEVICES == NULL ? 0 : atoi(KCPP_CLBLAST_DEVICES)); - printf("\nInitializing CLBlast (First Run)..."); - printf("\nAttempting to use: Platform=%d, Device=%d (If invalid, program will crash)\n",plat_num,dev_num); - cl_uint num_platforms; - clGetPlatformIDs(0, NULL, &num_platforms); - cl_platform_id* platforms = (cl_platform_id*)malloc(num_platforms*sizeof(cl_platform_id)); - clGetPlatformIDs(num_platforms, platforms, NULL); - platform = platforms[plat_num]; - char platform_buffer[1024]; - clGetPlatformInfo(platform, CL_PLATFORM_NAME, sizeof(platform_buffer), &platform_buffer, NULL); - cl_uint num_devices; - clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, NULL, &num_devices); - cl_device_id* devices = (cl_device_id*)malloc(num_devices*sizeof(cl_device_id)); - clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, num_devices, devices, NULL); - device = devices[dev_num]; - char device_buffer[1024]; - clGetDeviceInfo(device, CL_DEVICE_NAME, sizeof(device_buffer), &device_buffer, NULL); - printf("Using Platform: %s Device: %s\n", platform_buffer, device_buffer); - context = clCreateContext(NULL, 1, &device, NULL, NULL, &err); - if (err != CL_SUCCESS) { - printf("Error creating OpenCL context: %d\n", err); - fflush(stdout); - } - queue = clCreateCommandQueue(context, device, 0, &err); - - if (err != CL_SUCCESS) { - printf("Error creating OpenCL Command Queue: %d\n", err); - fflush(stdout); - } - - free(platforms); - free(devices); - - program = build_program_from_source(context, device, clblast_dequant); - - // Prepare dequantize kernels - kernel_q4_0 = clCreateKernel(program, "dequantize_row_q4_0", &err); - if(err < 0) { - printf("Error creating OpenCL dequantize q4_0 kernel: %d\n", err); - fflush(stdout); - }; - kernel_q4_1 = clCreateKernel(program, "dequantize_row_q4_1", &err); - if(err < 0) { - printf("Error creating OpenCL dequantize q4_1 kernel: %d\n", err); - fflush(stdout); - }; - cl_initialized = true; -} - -static void ggml_cl_sgemm_wrapper(const CBLAS_ORDER order, const CBLAS_TRANSPOSE trans_a, const CBLAS_TRANSPOSE trans_b, const int m, const int n, const int k, const float alpha, const void *host_a, const int lda, const float *host_b, const int ldb, const float beta, float *host_c, const int ldc, const int btype) { - cl_int err = 0; - - cl_event events[4]; - events[0] = NULL; - events[1] = NULL; - events[2] = NULL; - events[3] = NULL; - - if (!cl_initialized) { - ggml_cl_init(); - } - - bool dequant = (btype == 2 || btype == 3); - cl_kernel kernel = btype == 2 ? kernel_q4_0 : kernel_q4_1; - - size_t global = n * k, local = 16, qb_size; - cl_mem cl_buffer_a, cl_buffer_qb, cl_buffer_b, cl_buffer_c; - - // Prepare buffers - cl_buffer_a = clCreateBuffer(context, CL_MEM_READ_ONLY, m*k*sizeof(float), NULL, &err); - if (err != CL_SUCCESS) { - printf("Error creating OpenCL Buffer A: %d\n", err); - fflush(stdout); - } - if (dequant) { - qb_size = global * (sizeof(float) * (btype == 2 ? 1 : 2) + 16) / 32; - cl_buffer_qb = clCreateBuffer(context, CL_MEM_READ_ONLY, qb_size, NULL, &err); - if (err != CL_SUCCESS) { - printf("Error creating OpenCL Buffer QB: %d\n", err); - fflush(stdout); - } - } - cl_buffer_b = clCreateBuffer(context, CL_MEM_READ_WRITE, n*k*sizeof(float), NULL, &err); - if (err != CL_SUCCESS) { - printf("Error creating OpenCL Buffer B: %d\n", err); - fflush(stdout); - } - cl_buffer_c = clCreateBuffer(context, CL_MEM_READ_WRITE, m*n*sizeof(float), NULL, &err); - if (err != CL_SUCCESS) { - printf("Error creating OpenCL Buffer C: %d\n", err); - fflush(stdout); - } - - if (dequant) { - err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &cl_buffer_qb); - err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &cl_buffer_b); - if(err < 0) { - printf("Error setting OpenCL kernel args: %d\n", err); - fflush(stdout); - } - clEnqueueWriteBuffer(queue, cl_buffer_qb, CL_FALSE, 0, qb_size, host_b, 0, NULL, events + 1); - } else { - clEnqueueWriteBuffer(queue, cl_buffer_b, CL_FALSE, 0, n*k*sizeof(float), host_b, 0, NULL, events + 1); - } - - clEnqueueWriteBuffer(queue, cl_buffer_a, CL_FALSE, 0, m*k*sizeof(float), host_a, 0, NULL, events); - clEnqueueWriteBuffer(queue, cl_buffer_c, CL_FALSE, 0, m*n*sizeof(float), host_c, 0, NULL, events + 2); - if (dequant) { - err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, &local, 1, events + 1, events + 3); - if(err < 0) { - printf("Error enqueueing OpenCL dequantize kernel: %d\n", err); - fflush(stdout); - } - } - clWaitForEvents(dequant ? 4 : 3, events); - clReleaseEvent(events[0]); - clReleaseEvent(events[1]); - clReleaseEvent(events[2]); - if (dequant) { - clReleaseEvent(events[3]); - } - - // Call the SGEMM routine. - CLBlastStatusCode status = CLBlastSgemm(order, - trans_a, trans_b, - m, n, k, - alpha, - cl_buffer_a, 0, lda, - cl_buffer_b, 0, ldb, - beta, - cl_buffer_c, 0, ldc, - &queue, events); - - clEnqueueReadBuffer(queue, cl_buffer_c, CL_TRUE, 0, m*n*sizeof(float), host_c, 1, events, events + 1); - - // Wait for completion - if (status == CLBlastSuccess) { - clWaitForEvents(2, events); - clReleaseEvent(events[0]); - clReleaseEvent(events[1]); - } - - clReleaseMemObject(cl_buffer_a); - if (dequant) { - clReleaseMemObject(cl_buffer_qb); - } - clReleaseMemObject(cl_buffer_b); - clReleaseMemObject(cl_buffer_c); -} +#include "ggml-opencl.h" #endif #undef MIN @@ -7774,7 +7575,7 @@ static void ggml_compute_forward_mul_mat_f32( CUDA_CHECK(cudaMemcpyAsync(d, d_D, sizeof(float) * d_ne, cudaMemcpyDeviceToHost, g_cudaStream)); #elif defined(GGML_USE_CLBLAST) // zT = y * xT - ggml_cl_sgemm_wrapper(CblasRowMajor, CblasNoTrans, CblasTrans, + ggml_cl_sgemm_wrapper(CLBlastLayoutRowMajor, CLBlastTransposeNo, CLBlastTransposeYes, ne11, ne01, ne10, 1.0f, y, ne10, x, ne10, @@ -8008,7 +7809,7 @@ static void ggml_compute_forward_mul_mat_f16_f32( float * d = (float *) ((char *) dst->data + i02*nb2 + i03*nb3); // zT = y * xT - ggml_cl_sgemm_wrapper(CblasRowMajor, CblasNoTrans, CblasTrans, + ggml_cl_sgemm_wrapper(CLBlastLayoutRowMajor, CLBlastTransposeNo, CLBlastTransposeYes, ne11, ne01, ne10, 1.0f, y, ne10, x, ne10, @@ -8281,7 +8082,7 @@ static void ggml_compute_forward_mul_mat_q_f32( CUDA_CHECK(cudaMemcpyAsync(d, d_D, sizeof(float) * d_ne, cudaMemcpyDeviceToHost, g_cudaStream)); #elif defined(GGML_USE_CLBLAST) // zT = y * xT - ggml_cl_sgemm_wrapper(CblasRowMajor, CblasNoTrans, CblasTrans, + ggml_cl_sgemm_wrapper(CLBlastLayoutRowMajor, CLBlastTransposeNo, CLBlastTransposeYes, ne11, ne01, ne10, 1.0f, y, ne10, x, ne10,