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

Remove all constants

Browse source
This commit is contained in:
Henri Vasserman 2023-05-14 12:47:41 +03:00 • committed by GitHub
parent b8fb5cdf5c
commit 0453ce3f8b
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
1 changed files with 13 additions and 27 deletions
Download patch file Download diff file Expand all files Collapse all files

40
ggml-opencl.c
View file

@ -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);