diff --git a/expose.cpp b/expose.cpp index 893f427a0..d8b62b3e1 100644 --- a/expose.cpp +++ b/expose.cpp @@ -23,6 +23,8 @@ extern "C" { + std::string platformenv, deviceenv; + //return val: 0=fail, 1=(original ggml, alpaca), 2=(ggmf), 3=(ggjt) static FileFormat file_format = FileFormat::BADFORMAT; @@ -33,15 +35,15 @@ extern "C" //first digit is whether configured, second is platform, third is devices int parseinfo = inputs.clblast_info; - + std::string usingclblast = "KCPP_CLBLAST_CONFIGURED="+std::to_string(parseinfo>0?1:0); putenv((char*)usingclblast.c_str()); - parseinfo = parseinfo%100; //keep last 2 digits + parseinfo = parseinfo%100; //keep last 2 digits int platform = parseinfo/10; int devices = parseinfo%10; - std::string platformenv = "KCPP_CLBLAST_PLATFORM="+std::to_string(platform); - std::string deviceenv = "KCPP_CLBLAST_DEVICES="+std::to_string(devices); + platformenv = "KCPP_CLBLAST_PLATFORM="+std::to_string(platform); + deviceenv = "KCPP_CLBLAST_DEVICES="+std::to_string(devices); putenv((char*)platformenv.c_str()); putenv((char*)deviceenv.c_str()); @@ -61,7 +63,7 @@ extern "C" printf("\n---\nRetrying as GPT-J model: (ver %d)\nAttempting to Load...\n---\n", file_format); lr = gpttype_load_model(inputs, file_format); } - + if (lr == ModelLoadResult::FAIL || lr == ModelLoadResult::RETRY_LOAD) { return false; @@ -92,14 +94,14 @@ extern "C" } else { - printf("\n---\nIdentified as LLAMA model: (ver %d)\nAttempting to Load...\n---\n", file_format); + printf("\n---\nIdentified as LLAMA model: (ver %d)\nAttempting to Load...\n---\n", file_format); return llama_load_model(inputs, file_format); } } generation_outputs generate(const generation_inputs inputs, generation_outputs &output) { - if (file_format == FileFormat::GPTJ_1 || file_format == FileFormat::GPTJ_2 || file_format==FileFormat::GPTJ_3 + if (file_format == FileFormat::GPTJ_1 || file_format == FileFormat::GPTJ_2 || file_format==FileFormat::GPTJ_3 || file_format==FileFormat::GPT2_1 || file_format==FileFormat::GPT2_2 ) { return gpttype_generate(inputs, output); @@ -107,6 +109,6 @@ extern "C" else { return llama_generate(inputs, output); - } + } } } diff --git a/ggml.c b/ggml.c index c72084093..7b00b6ad3 100644 --- a/ggml.c +++ b/ggml.c @@ -6635,7 +6635,8 @@ static void ggml_compute_forward_mul_mat_f32( ne11, ne01, ne10, 1.0f, y, ne10, x, ne00, - 0.0f, d, ne01); + 0.0f, d, ne01, + params->type); } } @@ -6807,7 +6808,8 @@ static void ggml_compute_forward_mul_mat_f16_f32( ne11, ne01, ne10, 1.0f, y, ne10, x, ne00, - 0.0f, d, ne01); + 0.0f, d, ne01, + params->type); } } @@ -7002,6 +7004,7 @@ static void ggml_compute_forward_mul_mat_q_f32( for (int64_t i03 = 0; i03 < ne03; i03++) { for (int64_t i02 = 0; i02 < ne02; i02++) { +#ifndef GGML_USE_CLBLAST { size_t id = 0; for (int64_t i01 = 0; i01 < ne01; ++i01) { @@ -7011,6 +7014,9 @@ static void ggml_compute_forward_mul_mat_q_f32( } const float * x = wdata; +#else + const void* x = src0->data + i03*nb03 + i02*nb02; +#endif const float * y = (float *) ((char *) src1->data + i02*nb12 + i03*nb13); float * d = (float *) ((char *) dst->data + i02*nb2 + i03*nb3); @@ -7020,7 +7026,8 @@ static void ggml_compute_forward_mul_mat_q_f32( ne11, ne01, ne10, 1.0f, y, ne10, x, ne00, - 0.0f, d, ne01); + 0.0f, d, ne01, + type); } } diff --git a/ggml_blas_adapter.c b/ggml_blas_adapter.c index d41b001c9..a276faf58 100644 --- a/ggml_blas_adapter.c +++ b/ggml_blas_adapter.c @@ -18,15 +18,65 @@ 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; -static void ggml_cl_sgemm_wrapper(const enum CBLAS_ORDER order, const enum CBLAS_TRANSPOSE trans_a, const enum CBLAS_TRANSPOSE trans_b, const int m, const int n, const int k, const float alpha, const float *host_a, const int lda, const float *host_b, const int ldb, const float beta, float *host_c, const int ldc) { +// Function taken from https://github.com/rsnemmen/OpenCL-examples/blob/master/add_numbers/add_numbers.c +cl_program build_program(cl_context ctx, cl_device_id dev, const char* filename) { + + cl_program program; + FILE *program_handle; + char *program_buffer, *program_log; + size_t program_size, log_size; + int err; + + program_handle = fopen(filename, "r"); + if(program_handle == NULL) { + perror("OpenCL kernel file not found"); + exit(1); + } + fseek(program_handle, 0, SEEK_END); + program_size = ftell(program_handle); + rewind(program_handle); + program_buffer = (char*)malloc(program_size + 1); + program_buffer[program_size] = '\0'; + fread(program_buffer, sizeof(char), program_size, program_handle); + fclose(program_handle); + + program = clCreateProgramWithSource(ctx, 1, + (const char**)&program_buffer, &program_size, &err); + if(err < 0) { + perror("OpenCL error creating program"); + exit(1); + } + free(program_buffer); + + 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_sgemm_wrapper(const enum CBLAS_ORDER order, const enum CBLAS_TRANSPOSE trans_a, const enum 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[3]; + cl_event events[4]; events[0] = NULL; events[1] = NULL; events[2] = NULL; + events[3] = NULL; if (!cl_initialized) { char * KCPP_CLBLAST_PLATFORM = getenv("KCPP_CLBLAST_PLATFORM"); @@ -65,33 +115,82 @@ static void ggml_cl_sgemm_wrapper(const enum CBLAS_ORDER order, const enum CBLAS free(platforms); free(devices); + program = build_program(context, device, "ggml_clblast_dequant.cl"); + + // 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; } + 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_mem cl_buffer_a = clCreateBuffer(context, CL_MEM_READ_ONLY, m*k*sizeof(float), NULL, &err); + 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); } - cl_mem cl_buffer_b = clCreateBuffer(context, CL_MEM_READ_ONLY, n*k*sizeof(float), NULL, &err); + 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_mem cl_buffer_c = clCreateBuffer(context, CL_MEM_READ_WRITE, m*n*sizeof(float), NULL, &err); + 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_b, CL_FALSE, 0, n*k*sizeof(float), host_b, 0, NULL, events + 1); clEnqueueWriteBuffer(queue, cl_buffer_c, CL_FALSE, 0, m*n*sizeof(float), host_c, 0, NULL, events + 2); - clWaitForEvents(3, events); + 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, @@ -114,20 +213,22 @@ static void ggml_cl_sgemm_wrapper(const enum CBLAS_ORDER order, const enum CBLAS } clReleaseMemObject(cl_buffer_a); + if (dequant) { + clReleaseMemObject(cl_buffer_qb); + } clReleaseMemObject(cl_buffer_b); clReleaseMemObject(cl_buffer_c); } - #endif #endif #if defined(GGML_USE_ACCELERATE) || defined(GGML_USE_OPENBLAS) #if GGML_USE_CLBLAST -#define do_blas_sgemm(Order, TransA, TransB,M, N, K,alpha, A, lda, B, ldb, beta, C, ldc) ({\ -ggml_cl_sgemm_wrapper(Order, TransA, TransB, M, N, K, alpha, A, lda, B, ldb, beta, C, ldc);\ +#define do_blas_sgemm(Order, TransA, TransB,M, N, K,alpha, A, lda, B, ldb, beta, C, ldc, btype) ({\ +ggml_cl_sgemm_wrapper(Order, TransA, TransB, M, N, K, alpha, A, lda, B, ldb, beta, C, ldc, btype);\ }) #else -#define do_blas_sgemm(Order, TransA, TransB,M, N, K,alpha, A, lda, B, ldb, beta, C, ldc) ({\ +#define do_blas_sgemm(Order, TransA, TransB,M, N, K,alpha, A, lda, B, ldb, beta, C, ldc, btype) ({\ cblas_sgemm(Order, TransA, TransB, M, N, K, alpha, A, lda, B, ldb, beta, C, ldc);\ }) #endif diff --git a/ggml_clblast_dequant.cl b/ggml_clblast_dequant.cl new file mode 100644 index 000000000..87bdfe9bf --- /dev/null +++ b/ggml_clblast_dequant.cl @@ -0,0 +1,41 @@ +struct __attribute__ ((packed)) block_q4_0 +{ + float d; + uchar qs[16]; +}; + +__kernel void dequantize_row_q4_0(__global struct block_q4_0* blocks, __global float* result) { + uint i, l; + i = get_global_id(0) / 32; + l = get_local_id(0); + + float d = blocks[i].d; + + uchar vi = blocks[i].qs[l]; + + uint index = i*32 + l*2; + result[index + 0] = ((vi & 0xf) - 8)*d; + result[index + 1] = ((vi >> 4) - 8)*d; +} + +struct __attribute__ ((packed)) block_q4_1 +{ + float d; + float m; + uchar qs[16]; +}; + +__kernel void dequantize_row_q4_1(__global struct block_q4_1* blocks, __global float* result) { + uint i, l; + i = get_global_id(0) / 32; + l = get_local_id(0); + + float d = blocks[i].d; + float m = blocks[i].m; + + uchar vi = blocks[i].qs[l]; + + uint index = i*32 + l*2; + result[index + 0] = (vi & 0xf) * d + m; + result[index + 1] = (vi >> 4) * d + m; +}