mirror of
https://github.com/ggml-org/llama.cpp.git
synced 2026-09-27 21:46:57 +02:00
opencl: add bin kernel kernel_gemm_noshuffle_q5_k_f32_32b_trans_ila_a8_bin, kernel_gemm_noshuffle_q5_k_q8_1_dp4a_ila_a8_bin (#29401)
* opencl: add A8 Q5_K non-MoE non dp4a + dp4a binary kernel * opencl: fix s transpose - s only transposed for bin kernels --------- Co-authored-by: Li He <[email protected]>
This commit is contained in:
shaofeiqi
Li He
@@ -194,6 +194,7 @@ set(GGML_OPENCL_KERNELS
|
||||
gemv_noshuffle_q6_k_f32_32b_trans
|
||||
gemv_noshuffle_q5_k_f32
|
||||
gemm_noshuffle_q5_k_f32
|
||||
gemv_noshuffle_q5_k_f32_32b_trans
|
||||
mul
|
||||
neg
|
||||
norm
|
||||
|
||||
@@ -1264,6 +1264,9 @@ struct ggml_backend_opencl_context {
|
||||
cl_kernel kernel_gemv_noshuffle_q5_k_f32;
|
||||
cl_kernel kernel_gemv_noshuffle_q5_k_f32_mc3; // multi-column (N=3) verify GEMV (spec/MTP)
|
||||
cl_kernel kernel_gemm_noshuffle_q5_k_f32;
|
||||
cl_kernel kernel_gemm_noshuffle_q5_k_f32_32b_trans_ila_a8_bin;
|
||||
cl_kernel kernel_gemm_noshuffle_q5_k_q8_1_dp4a_ila_a8_bin;
|
||||
cl_kernel kernel_gemv_noshuffle_q5_k_f32_32b_trans;
|
||||
cl_kernel kernel_gemv_noshuffle_q5_0_f32;
|
||||
cl_kernel kernel_gemm_noshuffle_q5_0_f32;
|
||||
cl_kernel kernel_gemm_noshuffle_q5_0_q8_1_dp4a = nullptr; // dp4a (int8) dense q5_0 prefill GEMM
|
||||
@@ -4455,6 +4458,55 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
|
||||
}
|
||||
}
|
||||
|
||||
backend_ctx->kernel_gemv_noshuffle_q5_k_f32_32b_trans = nullptr;
|
||||
backend_ctx->kernel_gemm_noshuffle_q5_k_f32_32b_trans_ila_a8_bin = nullptr;
|
||||
backend_ctx->kernel_gemm_noshuffle_q5_k_q8_1_dp4a_ila_a8_bin = nullptr;
|
||||
if (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E) {
|
||||
{
|
||||
std::string opts = std::string("-cl-std=") + opencl_c_std +
|
||||
" -cl-mad-enable "
|
||||
" -DSIMDGROUP_WIDTH=" +
|
||||
std::to_string(backend_ctx->adreno_wave_size);
|
||||
#ifdef GGML_OPENCL_EMBED_KERNELS
|
||||
const std::string kernel_src {
|
||||
#include "gemv_noshuffle_q5_k_f32_32b_trans.cl.h"
|
||||
};
|
||||
#else
|
||||
const std::string kernel_src = read_file("gemv_noshuffle_q5_k_f32_32b_trans.cl");
|
||||
#endif
|
||||
cl_program prog = build_program_from_source(backend_ctx, kernel_src.c_str(), opts);
|
||||
CL_CHECK((backend_ctx->kernel_gemv_noshuffle_q5_k_f32_32b_trans =
|
||||
clCreateKernel(prog, "gemv_noshuffle_q5_k_f32_32b_trans", &err), err));
|
||||
CL_CHECK(clReleaseProgram(prog));
|
||||
GGML_LOG_CONT(".");
|
||||
}
|
||||
|
||||
if (use_adreno_bin_kernels(backend_ctx)) {
|
||||
size_t bin_size = 0;
|
||||
const char * kernel_bin = (const char *)backend_ctx->get_adreno_bin_kernel("gemm_noshuffle_q5_k_f32_32b_trans_ila_a8", &bin_size);
|
||||
if (kernel_bin && bin_size > 0) {
|
||||
cl_program bin_prog =
|
||||
build_program_from_binary(backend_ctx->context, backend_ctx->device, kernel_bin, "", bin_size);
|
||||
|
||||
CL_CHECK((backend_ctx->kernel_gemm_noshuffle_q5_k_f32_32b_trans_ila_a8_bin =
|
||||
clCreateKernel(bin_prog, "kernel_gemm_noshuffle_q5_k_f32_32b_trans_ila_a8", &err), err));
|
||||
CL_CHECK(clReleaseProgram(bin_prog));
|
||||
GGML_LOG_CONT(".");
|
||||
}
|
||||
|
||||
kernel_bin = (const char *)backend_ctx->get_adreno_bin_kernel("gemm_noshuffle_q5_k_q8_1_dp4a_ila_a8", &bin_size);
|
||||
if (kernel_bin && bin_size > 0) {
|
||||
cl_program bin_prog =
|
||||
build_program_from_binary(backend_ctx->context, backend_ctx->device, kernel_bin, "", bin_size);
|
||||
|
||||
CL_CHECK((backend_ctx->kernel_gemm_noshuffle_q5_k_q8_1_dp4a_ila_a8_bin =
|
||||
clCreateKernel(bin_prog, "kernel_gemm_noshuffle_q5_k_q8_1_dp4a_ila_a8", &err), err));
|
||||
CL_CHECK(clReleaseProgram(bin_prog));
|
||||
GGML_LOG_CONT(".");
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
std::string CL_moe_compile_opts = std::string("-cl-std=") + opencl_c_std +
|
||||
" -cl-mad-enable "
|
||||
" -cl-fast-relaxed-math";
|
||||
@@ -8749,6 +8801,20 @@ inline bool use_q4_k_bin_kernels(const ggml_backend_opencl_context *backend_ctx,
|
||||
#endif
|
||||
}
|
||||
|
||||
inline bool use_q5_k_bin_kernels(const ggml_backend_opencl_context *backend_ctx, const ggml_tensor *tensor) {
|
||||
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
|
||||
if (!backend_ctx->kernel_gemv_noshuffle_q5_k_f32_32b_trans ||
|
||||
!backend_ctx->kernel_gemm_noshuffle_q5_k_f32_32b_trans_ila_a8_bin) {
|
||||
return false;
|
||||
}
|
||||
return (tensor->ne[0] % 256 == 0) && (tensor->ne[1] % 64 == 0);
|
||||
#else
|
||||
GGML_UNUSED(backend_ctx);
|
||||
GGML_UNUSED(tensor);
|
||||
return false;
|
||||
#endif
|
||||
}
|
||||
|
||||
static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_tensor * op) {
|
||||
ggml_backend_opencl_device_context * dev_ctx = (ggml_backend_opencl_device_context *)dev->context;
|
||||
ggml_backend_opencl_context * backend_ctx = dev_ctx->backend_ctx;
|
||||
@@ -11147,8 +11213,29 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
|
||||
|
||||
GGML_ASSERT(K % 32 == 0);
|
||||
|
||||
// Transpose q, d, dm as ushort, qh as uchar
|
||||
transpose_2d_as_16b(backend_ctx, extra->q, extra->q, size_q, K/4, M);
|
||||
if (use_q5_k_bin_kernels(backend_ctx, tensor)) {
|
||||
cl_int err;
|
||||
cl_image_format wimg_fmt;
|
||||
cl_image_desc wimg_desc;
|
||||
|
||||
// transpose q as 32-bit words (M-first); qh/d/dm stay in their existing layout
|
||||
// (both new ILA kernels read qh via the existing [K/8][M] uchar plane directly).
|
||||
GGML_ASSERT(M % 64 == 0);
|
||||
transpose_2d_as_32b(backend_ctx, extra->q, extra->q, size_q, K/8, M);
|
||||
|
||||
wimg_fmt = { CL_R, CL_UNSIGNED_INT32 };
|
||||
memset(&wimg_desc, 0, sizeof(wimg_desc));
|
||||
wimg_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
|
||||
wimg_desc.image_width = (size_t)M * K / 8;
|
||||
wimg_desc.buffer = extra->q;
|
||||
CL_CHECK((extra->q_img = clCreateImage(context, CL_MEM_READ_ONLY, &wimg_fmt, &wimg_desc, NULL, &err), err));
|
||||
|
||||
// Transpose s as uchar
|
||||
transpose_2d_as_8b(backend_ctx, extra->s, extra->s, size_s, K/256*12, M, true, true);
|
||||
} else {
|
||||
// Transpose q as ushort
|
||||
transpose_2d_as_16b(backend_ctx, extra->q, extra->q, size_q, K/4, M);
|
||||
}
|
||||
transpose_2d_as_8b (backend_ctx, extra->qh, extra->qh, size_qh, K/8, M);
|
||||
transpose_2d_as_16b(backend_ctx, extra->d, extra->d, size_d, K/256, M);
|
||||
transpose_2d_as_16b(backend_ctx, extra->dm, extra->dm, size_dm, K/256, M);
|
||||
@@ -12333,21 +12420,28 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
|
||||
|
||||
size_t size_q = extra->size_q;
|
||||
size_t size_qh = extra->size_qh;
|
||||
size_t size_s = extra->size_s;
|
||||
size_t size_d = extra->size_d;
|
||||
size_t size_dm = extra->size_dm;
|
||||
|
||||
static ggml_cl_buffer buf_trans_q;
|
||||
static ggml_cl_buffer buf_trans_qh;
|
||||
static ggml_cl_buffer buf_trans_s;
|
||||
static ggml_cl_buffer buf_trans_d;
|
||||
static ggml_cl_buffer buf_trans_dm;
|
||||
|
||||
buf_trans_q.allocate(backend_ctx->context, size_q);
|
||||
buf_trans_qh.allocate(backend_ctx->context, size_qh);
|
||||
buf_trans_s.allocate(backend_ctx->context, size_s);
|
||||
buf_trans_d.allocate(backend_ctx->context, size_d);
|
||||
buf_trans_dm.allocate(backend_ctx->context, size_dm);
|
||||
|
||||
// Reverse transpose q, qh, d, dm
|
||||
transpose_2d_as_16b(backend_ctx, extra->q, buf_trans_q.buffer, size_q, M, K/4);
|
||||
if (use_q5_k_bin_kernels(backend_ctx, tensor)) {
|
||||
transpose_2d_as_32b(backend_ctx, extra->q, buf_trans_q.buffer, size_q, M, K/8);
|
||||
transpose_2d_as_8b (backend_ctx, extra->s, buf_trans_s.buffer, size_s, M, K/256*12, true, true);
|
||||
} else {
|
||||
transpose_2d_as_16b(backend_ctx, extra->q, buf_trans_q.buffer, size_q, M, K/4);
|
||||
}
|
||||
transpose_2d_as_8b (backend_ctx, extra->qh, buf_trans_qh.buffer, size_qh, M, K/8);
|
||||
transpose_2d_as_16b(backend_ctx, extra->d, buf_trans_d.buffer, size_d, M, K/256);
|
||||
transpose_2d_as_16b(backend_ctx, extra->dm, buf_trans_dm.buffer, size_dm, M, K/256);
|
||||
@@ -12355,7 +12449,7 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
|
||||
cl_kernel kernel = backend_ctx->kernel_restore_block_q5_K_noshuffle;
|
||||
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &buf_trans_q.buffer));
|
||||
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &buf_trans_qh.buffer));
|
||||
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra->s));
|
||||
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &buf_trans_s.buffer));
|
||||
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &buf_trans_d.buffer));
|
||||
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &buf_trans_dm.buffer));
|
||||
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), &data_device));
|
||||
@@ -22443,6 +22537,217 @@ static void ggml_cl_mul_mat_q6_K_f32_adreno(ggml_backend_t backend, const ggml_t
|
||||
#endif
|
||||
}
|
||||
|
||||
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
|
||||
static void ggml_cl_mul_mat_q5_K_f32_adreno_ila(ggml_backend_t backend, const ggml_tensor * src0,
|
||||
const ggml_tensor * src1, ggml_tensor * dst) {
|
||||
GGML_ASSERT(src0);
|
||||
GGML_ASSERT(src0->extra);
|
||||
GGML_ASSERT(src1);
|
||||
GGML_ASSERT(src1->extra);
|
||||
GGML_ASSERT(dst);
|
||||
GGML_ASSERT(dst->extra);
|
||||
|
||||
ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context;
|
||||
|
||||
ggml_tensor_extra_cl_q5_K * extra0_q5_k = (ggml_tensor_extra_cl_q5_K *)src0->extra;
|
||||
ggml_tensor_extra_cl * extra1 = (ggml_tensor_extra_cl *)src1->extra;
|
||||
ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *)dst->extra;
|
||||
|
||||
cl_ulong offset1 = extra1->offset + src1->view_offs;
|
||||
cl_ulong offsetd = extrad->offset + dst->view_offs;
|
||||
|
||||
const int ne00 = src0->ne[0];
|
||||
const int ne01 = src0->ne[1];
|
||||
|
||||
const int ne1 = dst->ne[1];
|
||||
|
||||
GGML_ASSERT(ne00 % ggml_blck_size(src0->type) == 0);
|
||||
|
||||
cl_context context = backend_ctx->context;
|
||||
cl_kernel kernel;
|
||||
|
||||
cl_int err;
|
||||
cl_buffer_region region;
|
||||
cl_image_format img_fmt;
|
||||
cl_image_desc img_desc;
|
||||
|
||||
const int M = ne01;
|
||||
const int N = ne1;
|
||||
const int K = ne00;
|
||||
|
||||
if (ne1 == 1) {
|
||||
cl_mem b_sub_buf = nullptr;
|
||||
cl_mem b_img = nullptr;
|
||||
|
||||
region.origin = offset1;
|
||||
region.size = (size_t)K * N * sizeof(float);
|
||||
CL_CHECK((b_sub_buf = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err));
|
||||
|
||||
img_fmt = { CL_RGBA, CL_FLOAT };
|
||||
memset(&img_desc, 0, sizeof(img_desc));
|
||||
img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
|
||||
img_desc.image_width = (size_t)K * N / 4;
|
||||
img_desc.buffer = b_sub_buf;
|
||||
CL_CHECK((b_img = clCreateImage(context, CL_MEM_READ_ONLY, &img_fmt, &img_desc, NULL, &err), err));
|
||||
|
||||
kernel = backend_ctx->kernel_gemv_noshuffle_q5_k_f32_32b_trans;
|
||||
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q5_k->q_img));
|
||||
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q5_k->qh));
|
||||
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_q5_k->d));
|
||||
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra0_q5_k->dm));
|
||||
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extra0_q5_k->s));
|
||||
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), &b_img));
|
||||
CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), &extrad->data_device));
|
||||
CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_ulong), &offsetd));
|
||||
CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_int), &ne00));
|
||||
CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_int), &ne01));
|
||||
|
||||
size_t local_work_size[3] = { 64, 8, 1 };
|
||||
size_t global_work_size[3] = { (size_t)ne01, 8, 1 };
|
||||
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
|
||||
|
||||
CL_CHECK(clReleaseMemObject(b_img));
|
||||
CL_CHECK(clReleaseMemObject(b_sub_buf));
|
||||
} else {
|
||||
static const char * q5_k_bin_dp4a_env = getenv("GGML_OPENCL_Q5_K_BIN_DP4A");
|
||||
bool q5_k_bin_dp4a_on = q5_k_bin_dp4a_env
|
||||
? (atoi(q5_k_bin_dp4a_env) != 0)
|
||||
: true;
|
||||
// dot prod has to be available
|
||||
q5_k_bin_dp4a_on = backend_ctx->has_integer_dot && q5_k_bin_dp4a_on;
|
||||
|
||||
if (q5_k_bin_dp4a_on && backend_ctx->kernel_gemm_noshuffle_q5_k_q8_1_dp4a_ila_a8_bin) {
|
||||
const int dp4a_N_pad = CEIL_DIV(N, 32) * 32;
|
||||
const size_t n_blocks = (size_t)dp4a_N_pad * (K / 32);
|
||||
|
||||
backend_ctx->prealloc_moe_qa.allocate(context, (size_t)dp4a_N_pad * K * sizeof(cl_char));
|
||||
backend_ctx->prealloc_moe_da.allocate(context, n_blocks * sizeof(cl_half));
|
||||
backend_ctx->prealloc_moe_sa.allocate(context, n_blocks * sizeof(cl_half));
|
||||
|
||||
cl_mem b_sub = nullptr;
|
||||
region.origin = offset1;
|
||||
region.size = (size_t)K * N * sizeof(float);
|
||||
CL_CHECK((b_sub = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err));
|
||||
|
||||
cl_int tb = (cl_int)((size_t)N * (K / 32));
|
||||
cl_kernel qk = backend_ctx->kernel_quant_a_q8_1;
|
||||
CL_CHECK(clSetKernelArg(qk, 0, sizeof(cl_mem), &b_sub));
|
||||
CL_CHECK(clSetKernelArg(qk, 1, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer));
|
||||
CL_CHECK(clSetKernelArg(qk, 2, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer));
|
||||
CL_CHECK(clSetKernelArg(qk, 3, sizeof(cl_mem), &backend_ctx->prealloc_moe_sa.buffer));
|
||||
CL_CHECK(clSetKernelArg(qk, 4, sizeof(cl_int), &tb));
|
||||
size_t q_local[1] = { 64 };
|
||||
size_t q_global[1] = { (size_t)CEIL_DIV(tb, 64) * 64 };
|
||||
backend_ctx->enqueue_ndrange_kernel(qk, 1, q_global, q_local, dst);
|
||||
|
||||
cl_mem d_sub = nullptr;
|
||||
cl_mem d_img = nullptr;
|
||||
region.origin = offsetd;
|
||||
region.size = (size_t)M * N * sizeof(float);
|
||||
CL_CHECK((d_sub = clCreateSubBuffer(extrad->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err));
|
||||
|
||||
img_fmt = { CL_R, CL_FLOAT };
|
||||
memset(&img_desc, 0, sizeof(img_desc));
|
||||
img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
|
||||
img_desc.image_width = (size_t)M * N;
|
||||
img_desc.buffer = d_sub;
|
||||
CL_CHECK((d_img = clCreateImage(context, CL_MEM_WRITE_ONLY, &img_fmt, &img_desc, NULL, &err), err));
|
||||
|
||||
kernel = backend_ctx->kernel_gemm_noshuffle_q5_k_q8_1_dp4a_ila_a8_bin;
|
||||
|
||||
cl_uint k_arg = 0;
|
||||
CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &extra0_q5_k->q_img));
|
||||
CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &extra0_q5_k->qh));
|
||||
CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &extra0_q5_k->d));
|
||||
CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &extra0_q5_k->dm));
|
||||
CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &extra0_q5_k->s));
|
||||
CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer));
|
||||
CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer));
|
||||
CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &backend_ctx->prealloc_moe_sa.buffer));
|
||||
CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &d_img));
|
||||
CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_uint), &ne00));
|
||||
CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_uint), &ne01));
|
||||
CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_int), &N));
|
||||
|
||||
size_t local_work_size_dp4a[3] = { 64, 1, 1 };
|
||||
size_t global_work_size_dp4a[3] = { 64, (size_t)(M / 64), (size_t)(dp4a_N_pad / 32) };
|
||||
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size_dp4a, local_work_size_dp4a, dst);
|
||||
|
||||
CL_CHECK(clReleaseMemObject(b_sub));
|
||||
CL_CHECK(clReleaseMemObject(d_img));
|
||||
CL_CHECK(clReleaseMemObject(d_sub));
|
||||
return;
|
||||
}
|
||||
|
||||
const int gemm_tile_n = 64;
|
||||
int N_pad = CEIL_DIV(N, gemm_tile_n) * gemm_tile_n;
|
||||
|
||||
cl_mem b_sub_buf = nullptr;
|
||||
cl_mem b_padded = nullptr;
|
||||
cl_mem b_buf = nullptr;
|
||||
if (N_pad == N) {
|
||||
region.origin = offset1;
|
||||
region.size = (size_t)K * N * sizeof(float);
|
||||
CL_CHECK((b_sub_buf = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err));
|
||||
b_buf = b_sub_buf;
|
||||
} else {
|
||||
CL_CHECK((b_padded = clCreateBuffer(context, CL_MEM_READ_WRITE, (size_t)K * N_pad * sizeof(float), NULL, &err), err));
|
||||
const float zero = 0.0f;
|
||||
CL_CHECK(clEnqueueFillBuffer(backend_ctx->queue, b_padded, &zero, sizeof(zero), 0, (size_t)K * N_pad * sizeof(float), 0, NULL, NULL));
|
||||
CL_CHECK(clEnqueueCopyBuffer(backend_ctx->queue, extra1->data_device, b_padded, offset1, 0, (size_t)K * N * sizeof(float), 0, NULL, NULL));
|
||||
b_buf = b_padded;
|
||||
}
|
||||
|
||||
img_fmt = { CL_R, CL_FLOAT };
|
||||
memset(&img_desc, 0, sizeof(img_desc));
|
||||
img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
|
||||
img_desc.image_width = (size_t)K * N_pad;
|
||||
img_desc.buffer = b_buf;
|
||||
cl_mem b_img;
|
||||
CL_CHECK((b_img = clCreateImage(context, CL_MEM_READ_ONLY, &img_fmt, &img_desc, NULL, &err), err));
|
||||
|
||||
region.origin = offsetd;
|
||||
region.size = (size_t)M * N * sizeof(float);
|
||||
cl_mem d_sub_buf;
|
||||
CL_CHECK((d_sub_buf = clCreateSubBuffer(extrad->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err));
|
||||
img_fmt = { CL_R, CL_FLOAT };
|
||||
memset(&img_desc, 0, sizeof(img_desc));
|
||||
img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
|
||||
img_desc.image_width = (size_t)M * N;
|
||||
img_desc.buffer = d_sub_buf;
|
||||
cl_mem d_img;
|
||||
CL_CHECK((d_img = clCreateImage(context, CL_MEM_WRITE_ONLY, &img_fmt, &img_desc, NULL, &err), err));
|
||||
|
||||
kernel = backend_ctx->kernel_gemm_noshuffle_q5_k_f32_32b_trans_ila_a8_bin;
|
||||
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q5_k->q_img));
|
||||
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q5_k->qh));
|
||||
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_q5_k->d));
|
||||
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra0_q5_k->dm));
|
||||
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extra0_q5_k->s));
|
||||
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), &b_img));
|
||||
CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_mem), &d_img));
|
||||
CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_uint), &ne00));
|
||||
CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_uint), &ne01));
|
||||
CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), &N));
|
||||
|
||||
size_t local_work_size[3] = { 64, 2, 2 };
|
||||
size_t m_tiles = (size_t)CEIL_DIV(M, 64);
|
||||
size_t global_work_size[3] = { 64, m_tiles, (size_t)CEIL_DIV(N_pad, gemm_tile_n) };
|
||||
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
|
||||
|
||||
CL_CHECK(clReleaseMemObject(b_img));
|
||||
if (b_sub_buf) {
|
||||
CL_CHECK(clReleaseMemObject(b_sub_buf));
|
||||
}
|
||||
if (b_padded) {
|
||||
CL_CHECK(clReleaseMemObject(b_padded));
|
||||
}
|
||||
CL_CHECK(clReleaseMemObject(d_img));
|
||||
CL_CHECK(clReleaseMemObject(d_sub_buf));
|
||||
}
|
||||
}
|
||||
#endif // GGML_OPENCL_USE_ADRENO_KERNELS
|
||||
|
||||
static void ggml_cl_mul_mat_q5_K_f32_adreno(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
|
||||
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
|
||||
GGML_ASSERT(src0);
|
||||
@@ -22491,6 +22796,20 @@ static void ggml_cl_mul_mat_q5_K_f32_adreno(ggml_backend_t backend, const ggml_t
|
||||
static const bool q5k_mc3 = (getenv("GGML_OPENCL_Q5K_MC3") != nullptr);
|
||||
const bool use_q5k_mc3 = q5k_mc3 && (ne1 >= 2 && ne1 <= 4) && (ne01 < 32768);
|
||||
|
||||
const bool use_bin = use_q5_k_bin_kernels(backend_ctx, src0);
|
||||
|
||||
if (use_bin) {
|
||||
if (use_q5k_mc3) {
|
||||
static bool warned = false;
|
||||
if (!warned) {
|
||||
GGML_LOG_WARN("ggml_opencl: GGML_OPENCL_Q5K_MC3 is bypassed by Q5_K binary kernels\n");
|
||||
warned = true;
|
||||
}
|
||||
}
|
||||
ggml_cl_mul_mat_q5_K_f32_adreno_ila(backend, src0, src1, dst);
|
||||
return;
|
||||
}
|
||||
|
||||
if (ne1 == 1 || use_q5k_mc3) {
|
||||
cl_mem q_img = nullptr;
|
||||
cl_mem qh_img = nullptr;
|
||||
|
||||
@@ -0,0 +1,141 @@
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
#pragma OPENCL EXTENSION cl_khr_subgroups : enable
|
||||
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
|
||||
|
||||
#define QK_K 256
|
||||
#define K_SCALE_SIZE 12
|
||||
#define N_SIMDGROUP 8
|
||||
#define SIMDGROUP_WIDTH 64
|
||||
|
||||
inline void get_scale_min_k4(
|
||||
int j,
|
||||
global const uchar * q,
|
||||
uint stride,
|
||||
uchar * d,
|
||||
uchar * m
|
||||
) {
|
||||
if (j < 4) {
|
||||
*d = q[j*stride] & 63;
|
||||
*m = q[(j+4)*stride] & 63;
|
||||
} else {
|
||||
*d = (q[(j+4)*stride] & 0x0F) | ((q[(j-4)*stride] & 0xC0) >> 2);
|
||||
*m = ((q[(j+4)*stride] >> 4) & 0x0F) | ((q[j*stride] & 0xC0) >> 2);
|
||||
}
|
||||
}
|
||||
|
||||
static inline float8 q5_k_to_fp32_packed8(ushort2 q4x8, uint qh_byte, float scale, float minv) {
|
||||
float8 fp32x8;
|
||||
fp32x8.s0 = (float)(( q4x8.s0 & 0x000F) | (((qh_byte >> 0) & 1) << 4)) * scale - minv;
|
||||
fp32x8.s1 = (float)(((q4x8.s0 >> 4) & 0x000F) | (((qh_byte >> 1) & 1) << 4)) * scale - minv;
|
||||
fp32x8.s2 = (float)(((q4x8.s0 >> 8) & 0x000F) | (((qh_byte >> 2) & 1) << 4)) * scale - minv;
|
||||
fp32x8.s3 = (float)(((q4x8.s0 >> 12) & 0x000F) | (((qh_byte >> 3) & 1) << 4)) * scale - minv;
|
||||
fp32x8.s4 = (float)(( q4x8.s1 & 0x000F) | (((qh_byte >> 4) & 1) << 4)) * scale - minv;
|
||||
fp32x8.s5 = (float)(((q4x8.s1 >> 4) & 0x000F) | (((qh_byte >> 5) & 1) << 4)) * scale - minv;
|
||||
fp32x8.s6 = (float)(((q4x8.s1 >> 8) & 0x000F) | (((qh_byte >> 6) & 1) << 4)) * scale - minv;
|
||||
fp32x8.s7 = (float)(((q4x8.s1 >> 12) & 0x000F) | (((qh_byte >> 7) & 1) << 4)) * scale - minv;
|
||||
return fp32x8;
|
||||
}
|
||||
|
||||
__attribute__((qcom_reqd_sub_group_size("half")))
|
||||
__kernel void gemv_noshuffle_q5_k_f32_32b_trans(
|
||||
read_only image1d_buffer_t src0_q,
|
||||
__global uchar * src0_qh,
|
||||
__global half * src0_d,
|
||||
__global half * src0_dm,
|
||||
__global uchar * src0_s,
|
||||
__read_only image1d_buffer_t src1,
|
||||
__global float * dst,
|
||||
ulong offsetd,
|
||||
int ne00,
|
||||
int ne01
|
||||
) {
|
||||
uint i01 = get_global_id(0);
|
||||
uint sgid = get_local_id(1);
|
||||
uint slid = get_sub_group_local_id();
|
||||
|
||||
int num_subblocks = ne00 / 32;
|
||||
|
||||
__private float sum = 0.0f;
|
||||
|
||||
// Loop over sub-blocks of 32 elements, N_SIMDGROUP sub-blocks per iter.
|
||||
for (uint ib = sgid; ib < num_subblocks; ib += N_SIMDGROUP) {
|
||||
uint sb = ib / 8;
|
||||
uint j = ib % 8;
|
||||
|
||||
// Load d and dmin for this super-block.
|
||||
half d_val = src0_d[sb * ne01 + i01];
|
||||
half dm_val = src0_dm[sb * ne01 + i01];
|
||||
|
||||
// Load sub-block scale and min. s is transposed [nb][12][M]; stride ne01 per code.
|
||||
global const uchar * sc = src0_s + sb * K_SCALE_SIZE * ne01 + i01;
|
||||
uchar sv, mn;
|
||||
get_scale_min_k4(j, sc, ne01, &sv, &mn);
|
||||
|
||||
float scale = (float)d_val * (float)sv;
|
||||
float minv = (float)dm_val * (float)mn;
|
||||
|
||||
// Load 4 uints of quants (32 nibbles = 32 elements), column-major stride ne01.
|
||||
uint q_base = ib * ne01 * 4 + i01;
|
||||
|
||||
uint4 regQ;
|
||||
regQ.s0 = read_imageui(src0_q, q_base).x;
|
||||
regQ.s1 = read_imageui(src0_q, q_base + ne01).x;
|
||||
regQ.s2 = read_imageui(src0_q, q_base + ne01 * 2).x;
|
||||
regQ.s3 = read_imageui(src0_q, q_base + ne01 * 3).x;
|
||||
|
||||
uint qh_grp = ib * 4;
|
||||
uint qh_word = (uint)src0_qh[(qh_grp + 0) * ne01 + i01]
|
||||
| ((uint)src0_qh[(qh_grp + 1) * ne01 + i01] << 8)
|
||||
| ((uint)src0_qh[(qh_grp + 2) * ne01 + i01] << 16)
|
||||
| ((uint)src0_qh[(qh_grp + 3) * ne01 + i01] << 24);
|
||||
|
||||
// Load activations: 32 floats = 8 float4s.
|
||||
uint y_offset = ib * 8;
|
||||
|
||||
float4 y_local = (slid < 8) ? read_imagef(src1, (y_offset + slid)) : (float4)0.0f;
|
||||
float4 y0 = sub_group_broadcast(y_local, 0);
|
||||
float4 y1 = sub_group_broadcast(y_local, 1);
|
||||
float4 y2 = sub_group_broadcast(y_local, 2);
|
||||
float4 y3 = sub_group_broadcast(y_local, 3);
|
||||
float4 y4 = sub_group_broadcast(y_local, 4);
|
||||
float4 y5 = sub_group_broadcast(y_local, 5);
|
||||
float4 y6 = sub_group_broadcast(y_local, 6);
|
||||
float4 y7 = sub_group_broadcast(y_local, 7);
|
||||
|
||||
float8 fp32x8 = q5_k_to_fp32_packed8(as_ushort2(regQ.s0), qh_word & 0xFF, scale, minv);
|
||||
float4 acc = y0 * fp32x8.lo;
|
||||
acc += y1 * fp32x8.hi;
|
||||
|
||||
fp32x8 = q5_k_to_fp32_packed8(as_ushort2(regQ.s1), (qh_word >> 8) & 0xFF, scale, minv);
|
||||
acc += y2 * fp32x8.lo;
|
||||
acc += y3 * fp32x8.hi;
|
||||
|
||||
fp32x8 = q5_k_to_fp32_packed8(as_ushort2(regQ.s2), (qh_word >> 16) & 0xFF, scale, minv);
|
||||
acc += y4 * fp32x8.lo;
|
||||
acc += y5 * fp32x8.hi;
|
||||
|
||||
fp32x8 = q5_k_to_fp32_packed8(as_ushort2(regQ.s3), (qh_word >> 24) & 0xFF, scale, minv);
|
||||
acc += y6 * fp32x8.lo;
|
||||
acc += y7 * fp32x8.hi;
|
||||
|
||||
sum += ((acc.s0 + acc.s1) + (acc.s2 + acc.s3));
|
||||
}
|
||||
|
||||
// reduction in local memory over N_SIMDGROUP subgroups
|
||||
__local float reduceLM[SIMDGROUP_WIDTH * (N_SIMDGROUP - 1)];
|
||||
if (sgid > 0) {
|
||||
reduceLM[SIMDGROUP_WIDTH * (sgid - 1) + slid] = sum;
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
if (sgid == 0) {
|
||||
for (uint i = 0; i < N_SIMDGROUP - 1; ++i) {
|
||||
sum += reduceLM[SIMDGROUP_WIDTH * i + slid];
|
||||
}
|
||||
}
|
||||
|
||||
// 1 output per thread in subgroup 0
|
||||
if (sgid == 0) {
|
||||
dst = dst + (offsetd >> 2);
|
||||
dst[i01] = sum;
|
||||
}
|
||||
}
|
||||
Reference in New Issue
Block a user