diff --git a/ggml-opencl.c b/ggml-opencl.c index 31ab13b25..a959f3aef 100644 --- a/ggml-opencl.c +++ b/ggml-opencl.c @@ -16,48 +16,48 @@ typedef uchar uint8_t; typedef int int32_t; typedef uint uint32_t; -constant uint QK4_0 = 32; +//constant uint QK4_0 = 32; struct block_q4_0 { float d; - uint8_t qs[QK4_0 / 2]; + uint8_t qs[16]; // QK4_0 / 2 }; -constant uint QK4_1 = 32; +//constant uint QK4_1 = 32; struct block_q4_1 { float d; float m; - uint8_t qs[QK4_1 / 2]; + uint8_t qs[16]; // QK4_1 / 2 }; -constant uint QK5_0 = 32; +//constant uint QK5_0 = 32; struct __attribute__ ((packed)) block_q5_0 { half d; uint32_t qh; - uint8_t qs[QK5_0 / 2]; + uint8_t qs[16]; // QK5_0 / 2 }; -constant uint QK5_1 = 32; +//constant uint QK5_1 = 32; struct block_q5_1 { half d; half m; uint32_t qh; - uint8_t qs[QK5_1 / 2]; + uint8_t qs[16]; // QK5_1 / 2 }; -constant uint QK8_0 = 32; +//constant uint QK8_0 = 32; struct block_q8_0 { float d; - uint8_t qs[QK8_0]; + 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 = QK4_0; + constant uint qk = 32; // QK4_0; const uint i = get_global_id(0) / qk; const uint j = get_local_id(0); @@ -72,7 +72,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 = QK4_1; + constant uint qk = 32; // QK4_1; const uint i = get_global_id(0) / qk; const uint j = get_local_id(0); @@ -88,7 +88,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 = QK5_0; + constant uint qk = 32; // QK5_0; const uint i = get_global_id(0) / qk; const uint j = get_local_id(0); @@ -108,7 +108,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 = QK5_1; + constant uint qk = 32; // QK5_1; const uint i = get_global_id(0) / qk; const uint j = get_local_id(0); @@ -129,7 +129,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 = QK8_0; + constant uint qk = 32; // QK8_0; const uint i = get_global_id(0) / qk; const uint j = get_local_id(0);