From 78b1d8351f64e421e93a139f5ed2288f24ee46c6 Mon Sep 17 00:00:00 2001 From: 0cc4m Date: Fri, 19 May 2023 21:18:57 +0200 Subject: [PATCH 1/3] Add OpenCL compile options --- ggml-opencl.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/ggml-opencl.cpp b/ggml-opencl.cpp index 604019c3e..ff72de5a8 100644 --- a/ggml-opencl.cpp +++ b/ggml-opencl.cpp @@ -287,7 +287,9 @@ static cl_program build_program_from_source(cl_context ctx, cl_device_id dev, co exit(1); } - err = clBuildProgram(p, 0, NULL, NULL, NULL, NULL); + const char* compile_opts = "-cl-mad-enable -cl-unsafe-math-optimizations -cl-finite-math-only -cl-fast-relaxed-math"; + + err = clBuildProgram(p, 0, NULL, compile_opts, NULL, NULL); if(err < 0) { clGetProgramBuildInfo(p, dev, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size); From 285f8f990b412435405abc5e03140866eb00f658 Mon Sep 17 00:00:00 2001 From: 0cc4m Date: Sat, 20 May 2023 07:26:38 +0200 Subject: [PATCH 2/3] Explicitely set CLBlast GEMM type --- ggml-opencl.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/ggml-opencl.cpp b/ggml-opencl.cpp index ff72de5a8..4a0c03c2b 100644 --- a/ggml-opencl.cpp +++ b/ggml-opencl.cpp @@ -552,7 +552,7 @@ static void ggml_cl_mul_mat_f32(const ggml_tensor * src0, const ggml_tensor * sr // compute cl_event ev_sgemm; - clblast::StatusCode status = clblast::Gemm(clblast::Layout::kColMajor, + clblast::StatusCode status = clblast::Gemm(clblast::Layout::kColMajor, clblast::Transpose::kYes, clblast::Transpose::kNo, ne01, ne11, ne10, alpha, @@ -650,7 +650,7 @@ static void ggml_cl_mul_mat_f16(const ggml_tensor * src0, const ggml_tensor * sr // compute cl_event ev_sgemm; - clblast::StatusCode status = clblast::Gemm(clblast::Layout::kColMajor, + clblast::StatusCode status = clblast::Gemm(clblast::Layout::kColMajor, clblast::Transpose::kYes, clblast::Transpose::kNo, ne01, ne11, ne10, alpha, @@ -757,7 +757,7 @@ static void ggml_cl_mul_mat_q_f32(const ggml_tensor * src0, const ggml_tensor * CL_CHECK(clFinish(queue), "clFinish"); // compute - clblast::StatusCode status = clblast::Gemm(clblast::Layout::kColMajor, + clblast::StatusCode status = clblast::Gemm(clblast::Layout::kColMajor, clblast::Transpose::kYes, clblast::Transpose::kNo, ne01, ne11, ne10, alpha, From 02914698f0c7083ecc69f344a356bc54ec405e61 Mon Sep 17 00:00:00 2001 From: 0cc4m Date: Sat, 20 May 2023 07:45:56 +0200 Subject: [PATCH 3/3] Update Q4_0, Q4_1 and Q8_0 to use half instead of float --- ggml-opencl.cpp | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/ggml-opencl.cpp b/ggml-opencl.cpp index 4a0c03c2b..af9c7651f 100644 --- a/ggml-opencl.cpp +++ b/ggml-opencl.cpp @@ -27,7 +27,7 @@ constant uint QK4_0 = 32; constant uint QR4_0 = 2; struct block_q4_0 { - float d; + half d; uint8_t qs[QK4_0 / 2]; }; @@ -35,8 +35,8 @@ constant uint QK4_1 = 32; constant uint QR4_1 = 2; struct block_q4_1 { - float d; - float m; + half d; + half m; uint8_t qs[QK4_1 / 2]; }; @@ -63,7 +63,7 @@ constant uint QK8_0 = 32; constant uint QR8_0 = 1; struct block_q8_0 { - float d; + half d; uint8_t qs[QK8_0]; }; @@ -75,7 +75,7 @@ __kernel void convert_fp16_to_fp32(__global half* x, __global float* y) { } void dequantize_q4_0(__global const struct block_q4_0* x, const int ib, const int iqs, float* v0, float* v1) { - const float d = x[ib].d; + const float d = vload_half(0, &x[ib].d); const uint8_t vui = x[ib].qs[iqs]; @@ -86,8 +86,8 @@ void dequantize_q4_0(__global const struct block_q4_0* x, const int ib, const in *v1 = (vi1 - 8)*d; } void dequantize_q4_1(__global const struct block_q4_1* x, const int ib, const int iqs, float* v0, float* v1) { - const float d = x[ib].d; - const float m = x[ib].m; + const float d = vload_half(0, &x[ib].d); + const float m = vload_half(0, &x[ib].m); const uint8_t vui = x[ib].qs[iqs]; @@ -127,7 +127,7 @@ void dequantize_q5_1(__global const struct block_q5_1* x, const int ib, const in *v1 = x1*d + m; } void dequantize_q8_0(__global const struct block_q8_0* x, const int ib, const int iqs, float* v0, float* v1) { - const float d = x[ib].d; + const float d = vload_half(0, &x[ib].d); const int8_t vi0 = x[ib].qs[iqs + 0]; const int8_t vi1 = x[ib].qs[iqs + 1];