converted the cl file to be a string literal instead

LostRuins · Apr 16, 2023 · 8bf2e50 · 8bf2e50
1 parent 5a4d1b5
commit 8bf2e50
Show file tree

Hide file tree

Showing 2 changed files with 40 additions and 1 deletion.
diff --git a/ggml_blas_adapter.c b/ggml_blas_adapter.c
@@ -13,6 +13,7 @@
 
 #define CL_TARGET_OPENCL_VERSION 110
 #include <clblast_c.h>
+#include <ggml_clblast_dequant.cl>
 
 cl_platform_id platform;
 cl_device_id device;
@@ -69,6 +70,39 @@ cl_program build_program(cl_context ctx, cl_device_id dev, const char* filename)
    return program;
 }
 
+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_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;
 
@@ -115,7 +149,7 @@ 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");
+        program = build_program_from_source(context, device, clblast_dequant);
 
         // Prepare dequantize kernels
         kernel_q4_0 = clCreateKernel(program, "dequantize_row_q4_0", &err);

diff --git a/ggml_clblast_dequant.cl b/ggml_clblast_dequant.cl
@@ -1,3 +1,6 @@
+#define MULTILINE_QUOTE(...) #__VA_ARGS__
+const char * clblast_dequant = MULTILINE_QUOTE(
+
 struct __attribute__ ((packed)) block_q4_0
 {
     float d;
@@ -39,3 +42,5 @@ __kernel void dequantize_row_q4_1(__global struct block_q4_1* blocks, __global f
     result[index + 0] = (vi & 0xf) * d + m;
     result[index + 1] = (vi >> 4) * d + m;
 }
+
+);