vbatts/llama.cpp
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

Move CLBlast implementation to separate file

Browse source 
Add buffer reuse code (adapted from slaren's cuda implementation)
This commit is contained in:
0cc4m 2023-04-23 09:59:45 +02:00
parent 6f66870726
commit 1b16b8c90d
4 changed files with 281 additions and 207 deletions
Download patch file Download diff file Expand all files Collapse all files

11
Makefile
View file

@ -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

239
ggml-opencl.cpp Normal file
View file

@ -0,0 +1,239 @@
#include "ggml-opencl.h"
#include <atomic>
#include <cstdio>
#include <cstring>
#include <ggml_clblast_dequant.cl>
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);
}

31
ggml-opencl.h Normal file
View file

@ -0,0 +1,31 @@
#pragma once
#define CL_TARGET_OPENCL_VERSION 110
#include <clblast_c.h>
#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

207
ggml.c
View file

@ -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 <clblast_c.h>
#include <ggml_clblast_dequant.cl>
#include <cblas.h>
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,