Merge remote-tracking branch 'occam/clblast-gpu-dequant' into concedo

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);
+        platformenv = "KCPP_CLBLAST_PLATFORM="+std::to_string(platform);
-        std::string deviceenv = "KCPP_CLBLAST_DEVICES="+std::to_string(devices);
+        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);
-        }       
+        }
     }
 }

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);
             }
         }
 

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) ({\
+#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);\
+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

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;
+}