From 0453ce3f8b955c91a931c0c2a79f7403b0584b7f Mon Sep 17 00:00:00 2001
From: Henri Vasserman <henv@hot.ee>
Date: Sun, 14 May 2023 12:47:41 +0300
Subject: [PATCH] Remove all constants

---
 ggml-opencl.c | 40 +++++++++++++---------------------------
 1 file changed, 13 insertions(+), 27 deletions(-)

diff --git a/ggml-opencl.c b/ggml-opencl.c
index 32c1fac95..a9ab4620f 100644
--- a/ggml-opencl.c
+++ b/ggml-opencl.c
@@ -10,56 +10,49 @@
 #include "ggml.h"
 
 #define MULTILINE_QUOTE(...) #__VA_ARGS__
-const char * clblast_dequant = MULTILINE_QUOTE(
+static const char * program_source = MULTILINE_QUOTE(
 
 typedef uchar uint8_t;
 typedef int int32_t;
 typedef uint uint32_t;
 
-//constant uint QK4_0 = 32;
 struct block_q4_0
 {
     float d;
-    uint8_t qs[16]; // QK4_0 / 2
+    uint8_t qs[16]; /* QK4_0 / 2 */
 };
 
-//constant uint QK4_1 = 32;
 struct block_q4_1
 {
     float d;
     float m;
-    uint8_t qs[16]; // QK4_1 / 2
+    uint8_t qs[16]; /* QK4_1 / 2 */
 };
 
-//constant uint QK5_0 = 32;
 struct __attribute__ ((packed)) block_q5_0
 {
     half d;
     uint32_t qh;
-    uint8_t qs[16]; // QK5_0 / 2
+    uint8_t qs[16]; /* QK5_0 / 2 */
 };
 
-//constant uint QK5_1 = 32;
 struct block_q5_1
 {
     half d;
     half m;
     uint32_t qh;
-    uint8_t qs[16]; // QK5_1 / 2
+    uint8_t qs[16]; /* QK5_1 / 2 */
 };
 
-//constant uint QK8_0 = 32;
 struct block_q8_0
 {
     float d;
-    uint8_t qs[16]; // QK8_0 / 2
+    uint8_t qs[16]; /* QK8_0 / 2 */
 };
 
 
 __kernel void dequantize_row_q4_0(__global struct block_q4_0* x, __global float* y) {
-    constant uint qk = 32; // QK4_0;
-
-    const uint i = get_global_id(0) / qk;
+    const uint i = get_global_id(0) / 32; /* QK4_0 */
     const uint j = get_local_id(0);
 
     const float d = x[i].d;
@@ -72,9 +65,7 @@ __kernel void dequantize_row_q4_0(__global struct block_q4_0* x, __global float*
 }
 
 __kernel void dequantize_row_q4_1(__global struct block_q4_1* x, __global float* y) {
-    constant uint qk = 32; // QK4_1;
-
-    const uint i = get_global_id(0) / qk;
+    const uint i = get_global_id(0) / 32; /* QK4_1 */
     const uint j = get_local_id(0);
 
     const float d = x[i].d;
@@ -88,9 +79,7 @@ __kernel void dequantize_row_q4_1(__global struct block_q4_1* x, __global float*
 }
 
 __kernel void dequantize_row_q5_0(__global struct block_q5_0* x, __global float* y) {
-    constant uint qk = 32; // QK5_0;
-
-    const uint i = get_global_id(0) / qk;
+    const uint i = get_global_id(0) / 32; /* QK5_0 */
     const uint j = get_local_id(0);
 
     const float d = vload_half(0, (__global half*) &x[i].d);
@@ -108,9 +97,7 @@ __kernel void dequantize_row_q5_0(__global struct block_q5_0* x, __global float*
 }
 
 __kernel void dequantize_row_q5_1(__global struct block_q5_1* x, __global float* y) {
-    constant uint qk = 32; // QK5_1;
-
-    const uint i = get_global_id(0) / qk;
+    const uint i = get_global_id(0) / 32; /* QK5_1 */
     const uint j = get_local_id(0);
 
     const float d = vload_half(0, (__global half*) &x[i].d);
@@ -129,8 +116,7 @@ __kernel void dequantize_row_q5_1(__global struct block_q5_1* x, __global float*
 }
 
 __kernel void dequantize_row_q8_0(__global struct block_q8_0* x, __global float* y) {
-    constant uint qk = 32; // QK8_0;
-    const uint i = get_global_id(0) / qk;
+    const uint i = get_global_id(0) / 32; /* QK8_0 */
     const uint j = get_local_id(0);
 
     const float d = x[i].d;
@@ -300,13 +286,13 @@ void ggml_cl_init(void) {
 
     out_of_order = CL_TRUE;
     queue = clCreateCommandQueue(context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &err);
-    if (err == CL_INVALID_PROPERTY) {
+    if (err == CL_INVALID_PROPERTY || err == CL_INVALID_VALUE) {
         out_of_order = CL_FALSE;
         queue = clCreateCommandQueue(context, device, 0, &err);
     }
     CL_CHECK(err, "clCreateCommandQueue");
 
-    program = build_program_from_source(context, device, clblast_dequant);
+    program = build_program_from_source(context, device, program_source);
 
     // Prepare dequantize kernels
     kernel_q4_0 = clCreateKernel(program, "dequantize_row_q4_0", &err);