converted the cl file to be a string literal instead

This commit is contained in:
Concedo 2023-04-16 15:57:30 +08:00
parent 5a4d1b5d15
commit 8bf2e50a11
2 changed files with 40 additions and 1 deletions

View file

@ -13,6 +13,7 @@
#define CL_TARGET_OPENCL_VERSION 110 #define CL_TARGET_OPENCL_VERSION 110
#include <clblast_c.h> #include <clblast_c.h>
#include <ggml_clblast_dequant.cl>
cl_platform_id platform; cl_platform_id platform;
cl_device_id device; cl_device_id device;
@ -69,6 +70,39 @@ cl_program build_program(cl_context ctx, cl_device_id dev, const char* filename)
return program; 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) { 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_int err = 0;
@ -115,7 +149,7 @@ static void ggml_cl_sgemm_wrapper(const enum CBLAS_ORDER order, const enum CBLAS
free(platforms); free(platforms);
free(devices); free(devices);
program = build_program(context, device, "ggml_clblast_dequant.cl"); program = build_program_from_source(context, device, clblast_dequant);
// Prepare dequantize kernels // Prepare dequantize kernels
kernel_q4_0 = clCreateKernel(program, "dequantize_row_q4_0", &err); kernel_q4_0 = clCreateKernel(program, "dequantize_row_q4_0", &err);

View file

@ -1,3 +1,6 @@
#define MULTILINE_QUOTE(...) #__VA_ARGS__
const char * clblast_dequant = MULTILINE_QUOTE(
struct __attribute__ ((packed)) block_q4_0 struct __attribute__ ((packed)) block_q4_0
{ {
float d; 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 + 0] = (vi & 0xf) * d + m;
result[index + 1] = (vi >> 4) * d + m; result[index + 1] = (vi >> 4) * d + m;
} }
);