vbatts/llama.cpp
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

remove constants

Browse source
This commit is contained in:
Henri Vasserman 2023-05-13 22:04:17 +03:00
parent 08737ef720
commit bb5c3e2c70
No known key found for this signature in database
GPG key ID: 2995FC0F58B1A986
1 changed files with 15 additions and 15 deletions
Download patch file Download diff file Expand all files Collapse all files

30
ggml-opencl.c
View file

@ -16,48 +16,48 @@ typedef uchar uint8_t;
typedef int int32_t; typedef int int32_t;
typedef uint uint32_t; typedef uint uint32_t;
constant uint QK4_0 = 32; //constant uint QK4_0 = 32;
struct block_q4_0 struct block_q4_0
{ {
float d; 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 struct block_q4_1
{ {
float d; float d;
float m; 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 struct __attribute__ ((packed)) block_q5_0
{ {
half d; half d;
uint32_t qh; 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 struct block_q5_1
{ {
half d; half d;
half m; half m;
uint32_t qh; 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 struct block_q8_0
{ {
float d; 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) { __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 i = get_global_id(0) / qk;
const uint j = get_local_id(0); 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) { __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 i = get_global_id(0) / qk;
const uint j = get_local_id(0); 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) { __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 i = get_global_id(0) / qk;
const uint j = get_local_id(0); 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) { __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 i = get_global_id(0) / qk;
const uint j = get_local_id(0); 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) { __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 i = get_global_id(0) / qk;
const uint j = get_local_id(0); const uint j = get_local_id(0);