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

Fix q8_0 dequant kernel

Browse source
This commit is contained in:
0cc4m 2023-04-29 10:37:58 +02:00
parent 1560c10f24
commit d6be497ef6
1 changed files with 2 additions and 5 deletions
Download patch file Download diff file Expand all files Collapse all files

7
ggml-opencl-dequant.cl
View file

@ -116,17 +116,14 @@ __kernel void dequantize_row_q5_1(__global struct block_q5_1* blocks, __global f
struct block_q8_0
{
float d;
uchar qs[32];
char qs[32];
};
__kernel void dequantize_row_q8_0(__global struct block_q8_0* blocks, __global float* result) {
const uint i = get_global_id(0) / 32;
const uint l = get_local_id(0);
const float d = blocks[i].d;
const uint index = i*32 + l;
result[index] = blocks[i].qs[l] * d;
result[i*32 + l] = blocks[i].qs[l] * blocks[i].d;
}
);