wmantly/llama.cpp-xrip
1
0 Fork 0
Code Issues Pull Requests Actions 32 Packages Projects Releases Wiki Activity

opencl: add initial q2_0 support (#31340)

Browse Source 
* opencl: add q2_0 AOS mv

Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com>
Assisted-by: Codex

* opencl: add q2_0 SOA mv and tiled mm

Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com>
Assisted-by: Codex

* opencl: add q2_0 noshuffle gemm/gemv for adreno

Assisted-by: Codex

* opencl: add q2_0 dp4a for adreno

* opencl: optimize q2_0 adreno decode and narrow batches

---------

Co-authored-by: Li He <lih@qti.qualcomm.com>
This commit is contained in:
shaofeiqiandLi He authored and GitHub committed 2026-10-29 21:30:05 -07:00
1 parent 3c809f4438
commit e0f142d5dd
9 files changed
+1322 -2
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

ggml/src/ggml-opencl/CMakeLists.txt
+6
View File
@@ -114,6 +114,9 @@ set(GGML_OPENCL_KERNELS
gemv_noshuffle_iq2_xs_f32
gemm_noshuffle_iq2_s_q8_1_dp4a
gemv_noshuffle_iq2_s_f32
gemm_noshuffle_q2_0_f32
gemm_noshuffle_q2_0_q8_1_dp4a
gemv_noshuffle_q2_0_f32
gemm_noshuffle_q3_k_q8_1_dp4a
gemv_noshuffle_q3_k_f32
mul_mv_q2_k_f32
@@ -130,6 +133,9 @@ set(GGML_OPENCL_KERNELS
mul_mm_iq2_xs_f32_l4_lm
mul_mv_iq2_s_f32_flat
mul_mm_iq2_s_f32_l4_lm
mul_mv_q2_0_f32
mul_mv_q2_0_f32_flat
mul_mm_q2_0_f32_l4_lm
mul_mv_q3_k_f32_flat
mul_mm_q3_k_f32_l4_lm
mul_mv_q2_k_f32_flat
ggml/src/ggml-opencl/ggml-opencl.cpp
+575 -2
View File
@@ -1019,6 +1019,15 @@ struct ggml_backend_opencl_context {
ggml_cl_buffer iq2_xxs_grid_buffer;
cl_mem iq2_xxs_grid_image = nullptr;
cl_kernel kernel_mul_mv_iq2_xxs_f32;
cl_kernel kernel_convert_block_q2_0, kernel_restore_block_q2_0;
cl_kernel kernel_mul_mv_q2_0_f32_flat;
cl_kernel kernel_mul_mm_q2_0_f32_l4_lm;
cl_kernel kernel_gemv_noshuffle_q2_0_f32 = nullptr;
cl_kernel kernel_gemm_noshuffle_q2_0_f32 = nullptr;
cl_kernel kernel_gemm_noshuffle_q2_0_q8_1_dp4a = nullptr;
cl_kernel kernel_gemm_noshuffle_q2_0_q8_1_dp4a_narrow = nullptr;
int q2_0_mv_nsg = 8;
cl_kernel kernel_mul_mv_q2_0_f32;
cl_kernel kernel_mul_mv_q4_K_f32;
cl_kernel kernel_mul_mv_q4_K_f32_flat;
cl_kernel kernel_mul_mv_q5_K_f32;
@@ -2089,6 +2098,8 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
CL_CHECK((backend_ctx->kernel_restore_block_iq2_xs = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_iq2_xs", &err), err));
CL_CHECK((backend_ctx->kernel_convert_block_iq2_s = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_iq2_s", &err), err));
CL_CHECK((backend_ctx->kernel_restore_block_iq2_s = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_iq2_s", &err), err));
CL_CHECK((backend_ctx->kernel_convert_block_q2_0 = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_q2_0", &err), err));
CL_CHECK((backend_ctx->kernel_restore_block_q2_0 = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_q2_0", &err), err));
CL_CHECK((backend_ctx->kernel_convert_block_q3_K = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_q3_K", &err), err));
CL_CHECK((backend_ctx->kernel_restore_block_q3_K = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_q3_K", &err), err));
CL_CHECK((backend_ctx->kernel_convert_block_q2_K = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_q2_K", &err), err));
@@ -2686,6 +2697,51 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
GGML_LOG_CONT(".");
}
// mul_mv_q2_0_f32_flat
{
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "mul_mv_q2_0_f32_flat.cl.h"
};
#else
const std::string kernel_src = read_file("mul_mv_q2_0_f32_flat.cl");
#endif
cl_program prog = build_program_from_source(backend_ctx, kernel_src.c_str(), compile_opts);
CL_CHECK((backend_ctx->kernel_mul_mv_q2_0_f32_flat = clCreateKernel(prog, "kernel_mul_mv_q2_0_f32_flat", &err), err));
CL_CHECK(clReleaseProgram(prog));
GGML_LOG_CONT(".");
}
// mul_mm_q2_0_f32_l4_lm
{
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "mul_mm_q2_0_f32_l4_lm.cl.h"
};
#else
const std::string kernel_src = read_file("mul_mm_q2_0_f32_l4_lm.cl");
#endif
cl_program prog = build_program_from_source(backend_ctx, kernel_src.c_str(), compile_opts);
CL_CHECK((backend_ctx->kernel_mul_mm_q2_0_f32_l4_lm = clCreateKernel(prog, "kernel_mul_mm_q2_0_f32_l4_lm", &err), err));
CL_CHECK(clReleaseProgram(prog));
GGML_LOG_CONT(".");
}
// mul_mv_q2_0_f32
{
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "mul_mv_q2_0_f32.cl.h"
};
#else
const std::string kernel_src = read_file("mul_mv_q2_0_f32.cl");
#endif
cl_program prog = build_program_from_source(backend_ctx, kernel_src.c_str(), compile_opts);
CL_CHECK((backend_ctx->kernel_mul_mv_q2_0_f32 = clCreateKernel(prog, "kernel_mul_mv_q2_0_f32", &err), err));
CL_CHECK(clReleaseProgram(prog));
GGML_LOG_CONT(".");
}
// mul_mv_q3_k_f32
{
#ifdef GGML_OPENCL_EMBED_KERNELS
@@ -4529,6 +4585,73 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
GGML_LOG_CONT(".");
}
// gemm_noshuffle_q2_0_q8_1_dp4a
{
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "gemm_noshuffle_q2_0_q8_1_dp4a.cl.h"
};
#else
const std::string kernel_src = read_file("gemm_noshuffle_q2_0_q8_1_dp4a.cl");
#endif
cl_program prog = build_program_from_source(backend_ctx, kernel_src.c_str(), compile_opts);
CL_CHECK((backend_ctx->kernel_gemm_noshuffle_q2_0_q8_1_dp4a = clCreateKernel(prog, "kernel_gemm_noshuffle_q2_0_q8_1_dp4a", &err), err));
CL_CHECK(clReleaseProgram(prog));
prog = build_program_from_source(backend_ctx, kernel_src.c_str(), compile_opts + " -DTILESIZE_N=8");
CL_CHECK((backend_ctx->kernel_gemm_noshuffle_q2_0_q8_1_dp4a_narrow = clCreateKernel(prog, "kernel_gemm_noshuffle_q2_0_q8_1_dp4a", &err), err));
CL_CHECK(clReleaseProgram(prog));
GGML_LOG_CONT(".");
}
// gemm_noshuffle_q2_0_f32
{
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "gemm_noshuffle_q2_0_f32.cl.h"
};
#else
const std::string kernel_src = read_file("gemm_noshuffle_q2_0_f32.cl");
#endif
cl_program prog = build_program_from_source(backend_ctx, kernel_src.c_str(), compile_opts);
CL_CHECK((backend_ctx->kernel_gemm_noshuffle_q2_0_f32 = clCreateKernel(prog, "kernel_gemm_noshuffle_q2_0_f32", &err), err));
CL_CHECK(clReleaseProgram(prog));
GGML_LOG_CONT(".");
}
// gemv_noshuffle_q2_0_f32
{
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src {
#include "gemv_noshuffle_q2_0_f32.cl.h"
};
#else
const std::string kernel_src = read_file("gemv_noshuffle_q2_0_f32.cl");
#endif
if (const char * env = getenv("GGML_OPENCL_Q2_0_MV_NSG")) {
const int nsg = atoi(env);
if (nsg == 1 || nsg == 2 || nsg == 4 || nsg == 8) {
backend_ctx->q2_0_mv_nsg = nsg;
}
}
for (;;) {
const std::string opts = compile_opts + " -DQ2_0_MV_NSG=" + std::to_string(backend_ctx->q2_0_mv_nsg);
cl_program prog = build_program_from_source(backend_ctx, kernel_src.c_str(), opts);
cl_kernel kernel = clCreateKernel(prog, "kernel_gemv_noshuffle_q2_0_f32", &err);
CL_CHECK(err);
CL_CHECK(clReleaseProgram(prog));
size_t max_workgroup;
CL_CHECK(clGetKernelWorkGroupInfo(kernel, backend_ctx->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(max_workgroup), &max_workgroup, nullptr));
if (64*(size_t)backend_ctx->q2_0_mv_nsg <= max_workgroup) {
backend_ctx->kernel_gemv_noshuffle_q2_0_f32 = kernel;
break;
}
CL_CHECK(clReleaseKernel(kernel));
GGML_ASSERT(backend_ctx->q2_0_mv_nsg > 1);
backend_ctx->q2_0_mv_nsg /= 2;
}
GGML_LOG_CONT(".");
}
// gemm_noshuffle_iq2_s_q8_1_dp4a
{
#ifdef GGML_OPENCL_EMBED_KERNELS
@@ -8358,10 +8481,11 @@ static void transpose_2d_as_16b(
ggml_backend_opencl_context * backend_ctx,
cl_mem src, cl_mem dst, size_t size,
cl_int stride, cl_int rows,
bool blocking = true
bool blocking = true,
bool auto_local = false
) {
transpose_2d(backend_ctx, backend_ctx->kernel_transpose_16_buf,
src, dst, size, stride, rows, blocking);
src, dst, size, stride, rows, blocking, auto_local);
}
static void transpose_2d_as_32b(
@@ -8983,6 +9107,23 @@ struct ggml_tensor_extra_cl_iq2_s {
}
};
struct ggml_tensor_extra_cl_q2_0 {
bool noshuffle = false;
cl_mem qs = nullptr;
cl_mem d = nullptr;
size_t size_qs = 0, size_d = 0;
~ggml_tensor_extra_cl_q2_0() { reset(); }
void reset() {
noshuffle = false;
if (qs != nullptr) { CL_CHECK(clReleaseMemObject(qs)); qs = nullptr; }
if (d != nullptr) { CL_CHECK(clReleaseMemObject(d)); d = nullptr; }
size_qs = size_d = 0;
}
};
struct ggml_tensor_extra_cl_q6_K {
// Lower 4 bits of quantized weights.
cl_mem ql = nullptr;
@@ -10101,6 +10242,24 @@ static bool ggml_cl_q3_k_use_noshuffle(const ggml_backend_opencl_context * backe
#endif
}
static bool ggml_cl_q2_0_use_noshuffle(const ggml_backend_opencl_context * backend_ctx, const ggml_tensor * tensor) {
#if defined(GGML_OPENCL_SOA_Q) && defined(GGML_OPENCL_USE_ADRENO_KERNELS)
if (const char * env = getenv("GGML_OPENCL_Q2_0_NOSHUFFLE")) {
if (env[0] == '0') {
return false;
}
}
return tensor->type == GGML_TYPE_Q2_0 &&
backend_ctx->gpu_family == ADRENO &&
ggml_is_contiguous(tensor) && !tensor->view_src &&
use_adreno_kernels(backend_ctx, tensor);
#else
GGML_UNUSED(backend_ctx);
GGML_UNUSED(tensor);
return false;
#endif
}
static bool ggml_cl_iq2_s_use_noshuffle(const ggml_backend_opencl_context * backend_ctx, const ggml_tensor * tensor) {
#if defined(GGML_OPENCL_SOA_Q) && defined(GGML_OPENCL_USE_ADRENO_KERNELS)
if (const char * env = getenv("GGML_OPENCL_IQ2_S_NOSHUFFLE")) {
@@ -10698,6 +10857,12 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
return false;
}
return op->src[1]->type == GGML_TYPE_F32 && ggml_is_contiguous_rows(op->src[1]);
} else if (op->src[0]->type == GGML_TYPE_Q2_0) {
const ggml_tensor * weight = op->src[0]->view_src;
if (weight && ggml_cl_q2_0_use_noshuffle(backend_ctx, weight)) {
return false;
}
return op->src[1]->type == GGML_TYPE_F32 && ggml_is_contiguous_rows(op->src[1]);
} else if (op->src[0]->type == GGML_TYPE_Q3_K) {
const ggml_tensor * weight = op->src[0]->view_src;
if (weight && ggml_cl_q3_k_use_noshuffle(backend_ctx, weight)) {
@@ -11102,6 +11267,9 @@ struct ggml_backend_opencl_buffer_context {
for (ggml_tensor_extra_cl_iq2_xxs * e : temp_tensor_extras_iq2_xxs) {
delete e;
}
for (ggml_tensor_extra_cl_q2_0 * e : temp_tensor_extras_q2_0) {
delete e;
}
for (ggml_tensor_extra_cl_q6_K * e : temp_tensor_extras_q6_K_in_use) {
delete e;
}
@@ -11120,6 +11288,9 @@ struct ggml_backend_opencl_buffer_context {
for (ggml_tensor_extra_cl_iq2_xxs * e : temp_tensor_extras_iq2_xxs_in_use) {
delete e;
}
for (ggml_tensor_extra_cl_q2_0 * e : temp_tensor_extras_q2_0_in_use) {
delete e;
}
for (ggml_tensor_extra_cl_q5_K * e : temp_tensor_extras_q5_K) {
delete e;
}
@@ -11368,6 +11539,21 @@ struct ggml_backend_opencl_buffer_context {
return extra;
}
ggml_tensor_extra_cl_q2_0 * ggml_opencl_alloc_temp_tensor_extra_q2_0() {
ggml_tensor_extra_cl_q2_0 * extra;
if (temp_tensor_extras_q2_0.empty()) {
extra = new ggml_tensor_extra_cl_q2_0();
} else {
extra = temp_tensor_extras_q2_0.back();
temp_tensor_extras_q2_0.pop_back();
}
temp_tensor_extras_q2_0_in_use.push_back(extra);
extra->reset();
return extra;
}
ggml_tensor_extra_cl_q6_K * ggml_opencl_alloc_temp_tensor_extra_q6_K() {
ggml_tensor_extra_cl_q6_K * extra;
if (temp_tensor_extras_q6_K.empty()) {
@@ -11469,6 +11655,11 @@ struct ggml_backend_opencl_buffer_context {
}
temp_tensor_extras_iq2_xxs_in_use.clear();
for (ggml_tensor_extra_cl_q2_0 * e : temp_tensor_extras_q2_0_in_use) {
temp_tensor_extras_q2_0.push_back(e);
}
temp_tensor_extras_q2_0_in_use.clear();
q8_0_soa_tensors.clear();
q4_0_soa_tensors.clear();
}
@@ -11510,6 +11701,8 @@ struct ggml_backend_opencl_buffer_context {
std::vector<ggml_tensor_extra_cl_iq2_s *> temp_tensor_extras_iq2_s_in_use;
std::vector<ggml_tensor_extra_cl_iq2_xs *> temp_tensor_extras_iq2_xs_in_use;
std::vector<ggml_tensor_extra_cl_iq2_xxs *> temp_tensor_extras_iq2_xxs_in_use;
std::vector<ggml_tensor_extra_cl_q2_0 *> temp_tensor_extras_q2_0;
std::vector<ggml_tensor_extra_cl_q2_0 *> temp_tensor_extras_q2_0_in_use;
std::vector<ggml_tensor_extra_cl_q6_K *> temp_tensor_extras_q6_K;
std::vector<ggml_tensor_extra_cl_q6_K *> temp_tensor_extras_q6_K_in_use;
@@ -12162,6 +12355,72 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
return;
}
if (tensor->type == GGML_TYPE_Q2_0) {
GGML_ASSERT(offset == 0 && size == ggml_nbytes(tensor));
ggml_tensor_extra_cl * extra_orig = (ggml_tensor_extra_cl *)tensor->extra;
GGML_ASSERT(extra_orig && "Tensors in OpenCL backend should have been allocated and initialized");
ggml_backend_opencl_buffer_context * ctx = (ggml_backend_opencl_buffer_context *) buffer->context;
ggml_tensor_extra_cl_q2_0 * extra = ctx->ggml_opencl_alloc_temp_tensor_extra_q2_0();
extra->noshuffle = ggml_cl_q2_0_use_noshuffle(backend_ctx, tensor);
const size_t blck = (size_t)ggml_blck_size(tensor->type);
const size_t n_blk = ggml_nelements(tensor) / blck;
const size_t size_qs = n_blk * 16;
const size_t size_d = n_blk * 2;
GGML_ASSERT(size_qs + size_d == ggml_nbytes(tensor) && "Incorrect tensor size");
cl_int err;
cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
ggml_nbytes(tensor), NULL, &err);
CL_CHECK(err);
CL_CHECK(clEnqueueWriteBuffer(
queue, data_device, CL_TRUE, 0,
ggml_nbytes(tensor), data, 0, NULL, NULL));
cl_buffer_region region;
region.origin = align_to(extra_orig->offset + tensor->view_offs + offset, backend_ctx->alignment);
region.size = size_qs;
CL_CHECK((extra->qs = clCreateSubBuffer(extra_orig->data_device, CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, &region, &err), err));
region.origin = align_to(region.origin + size_qs, backend_ctx->alignment);
region.size = size_d;
CL_CHECK((extra->d = clCreateSubBuffer(extra_orig->data_device, CL_MEM_READ_WRITE, CL_BUFFER_CREATE_TYPE_REGION, &region, &err), err));
cl_kernel kernel = backend_ctx->kernel_convert_block_q2_0;
cl_ulong nb_arg = (cl_ulong)n_blk;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &data_device));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra->qs));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra->d));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &nb_arg));
size_t global_work_size[] = { (size_t)CEIL_DIV(n_blk, 64)*64, 1, 1 };
size_t local_work_size[] = { 64, 1, 1 };
cl_event evt;
CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 3, NULL, global_work_size, local_work_size, 0, NULL, &evt));
CL_CHECK(clWaitForEvents(1, &evt));
CL_CHECK(clReleaseEvent(evt));
CL_CHECK(clReleaseMemObject(data_device));
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
if (extra->noshuffle) {
const int m = tensor->ne[1];
const int k = tensor->ne[0];
transpose_2d_as_8b(backend_ctx, extra->qs, extra->qs, size_qs, k/4, m, true, true);
transpose_2d_as_16b(backend_ctx, extra->d, extra->d, size_d, k/64, m, true, true);
}
#endif
extra->size_qs = size_qs;
extra->size_d = size_d;
tensor->extra = extra;
return;
}
// We separate the quantized bits and scale from block_q4_0 by using an
// additional kernel, where each thread handles a block. We first read the
// original weights into a temporary buffer, then create two separate
@@ -14083,6 +14342,53 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
CL_CHECK(clReleaseMemObject(data_device));
return;
}
if (tensor->type == GGML_TYPE_Q2_0) {
if (tensor->view_src) {
offset += tensor->view_offs;
tensor = tensor->view_src;
}
ggml_tensor_extra_cl_q2_0 * extra = (ggml_tensor_extra_cl_q2_0 *)tensor->extra;
const size_t blck = (size_t)ggml_blck_size(tensor->type);
ggml_cl_buffer buf_unpacked;
buf_unpacked.allocate(backend_ctx->context, ggml_nbytes(tensor));
const size_t n_blk = (size_t)ggml_nelements(tensor) / blck;
cl_ulong nb_arg = (cl_ulong)n_blk;
cl_mem qs = extra->qs, d = extra->d;
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
ggml_cl_buffer buf_qs, buf_d;
if (extra->noshuffle) {
buf_qs.allocate(context, extra->size_qs);
qs = buf_qs.buffer;
buf_d.allocate(context, extra->size_d);
d = buf_d.buffer;
const int m = tensor->ne[1];
const int k = tensor->ne[0];
transpose_2d_as_8b(backend_ctx, extra->qs, qs, extra->size_qs, m, k/4, true, true);
transpose_2d_as_16b(backend_ctx, extra->d, d, extra->size_d, m, k/64, true, true);
}
#endif
cl_kernel kernel = backend_ctx->kernel_restore_block_q2_0;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &qs));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &d));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &buf_unpacked.buffer));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &nb_arg));
size_t gws[] = { (n_blk + 63) / 64 * 64 };
size_t lws[] = { 64 };
cl_event evt;
CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 1, NULL, gws, lws, 0, NULL, &evt));
CL_CHECK(clWaitForEvents(1, &evt));
CL_CHECK(clReleaseEvent(evt));
CL_CHECK(clEnqueueReadBuffer(queue, buf_unpacked.buffer, CL_TRUE, offset, size, data, 0, NULL, NULL));
return;
}
// In end-to-end runs, get_tensor is usually used to get back the logits,
// where we can simply do clEnqueueReadBuffer since they are f32.
// However, in test-backend-ops, the GPU graph is copied to the CPU backend,
@@ -22052,6 +22358,160 @@ static void ggml_cl_mul_mat_q1_0_f32_adreno(ggml_backend_t backend, const ggml_t
}
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
static bool ggml_cl_mul_mat_q2_0_dp4a(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *)backend->context;
ggml_tensor_extra_cl_q2_0 * extra0 = (ggml_tensor_extra_cl_q2_0 *)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 m = src0->ne[1];
const int n = src1->ne[1];
const int k = src0->ne[0];
const int tile_n = n <= 8 ? 8 : 32;
cl_kernel gemm = n <= 8 ? backend_ctx->kernel_gemm_noshuffle_q2_0_q8_1_dp4a_narrow
: backend_ctx->kernel_gemm_noshuffle_q2_0_q8_1_dp4a;
const char * env = getenv("GGML_OPENCL_Q2_0_DENSE_DP4A");
if (!gemm || !backend_ctx->has_integer_dot || (env && env[0] == '0') ||
n <= 1 || !ggml_is_contiguous(src1) || offset1 % backend_ctx->alignment != 0) {
return false;
}
cl_int err;
cl_buffer_region region = { (size_t)offset1, (size_t)k*n*sizeof(float) };
cl_mem activation = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, &region, &err);
CL_CHECK(err);
const size_t n_blocks = (size_t)n*(k/32);
backend_ctx->prealloc_moe_qa.allocate(backend_ctx->context, (size_t)n*k);
backend_ctx->prealloc_moe_da.allocate(backend_ctx->context, n_blocks*sizeof(cl_half));
backend_ctx->prealloc_moe_sa.allocate(backend_ctx->context, n_blocks*sizeof(cl_half));
cl_int nb = n_blocks;
cl_kernel quant = backend_ctx->kernel_quant_a_q8_1;
CL_CHECK(clSetKernelArg(quant, 0, sizeof(cl_mem), &activation));
CL_CHECK(clSetKernelArg(quant, 1, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer));
CL_CHECK(clSetKernelArg(quant, 2, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer));
CL_CHECK(clSetKernelArg(quant, 3, sizeof(cl_mem), &backend_ctx->prealloc_moe_sa.buffer));
CL_CHECK(clSetKernelArg(quant, 4, sizeof(cl_int), &nb));
size_t q_global[] = { CEIL_DIV(n_blocks, 64)*64 };
size_t q_local[] = { 64 };
backend_ctx->enqueue_ndrange_kernel(quant, 1, q_global, q_local, dst);
CL_CHECK(clSetKernelArg(gemm, 0, sizeof(cl_mem), &extra0->qs));
CL_CHECK(clSetKernelArg(gemm, 1, sizeof(cl_mem), &extra0->d));
CL_CHECK(clSetKernelArg(gemm, 2, sizeof(cl_mem), &backend_ctx->prealloc_moe_qa.buffer));
CL_CHECK(clSetKernelArg(gemm, 3, sizeof(cl_mem), &backend_ctx->prealloc_moe_da.buffer));
CL_CHECK(clSetKernelArg(gemm, 4, sizeof(cl_mem), &extrad->data_device));
CL_CHECK(clSetKernelArg(gemm, 5, sizeof(cl_ulong), &offsetd));
CL_CHECK(clSetKernelArg(gemm, 6, sizeof(int), &m));
CL_CHECK(clSetKernelArg(gemm, 7, sizeof(int), &n));
CL_CHECK(clSetKernelArg(gemm, 8, sizeof(int), &k));
size_t global[] = { 64, (size_t)CEIL_DIV(m, 64), (size_t)CEIL_DIV(n, tile_n) };
size_t local[] = { 64, 1, 1 };
backend_ctx->enqueue_ndrange_kernel(gemm, 3, global, local, dst);
CL_CHECK(clReleaseMemObject(activation));
return true;
}
static void ggml_cl_mul_mat_q2_0_f32_adreno(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *)backend->context;
ggml_tensor_extra_cl_q2_0 * extra0 = (ggml_tensor_extra_cl_q2_0 *)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 m = src0->ne[1];
const int n = src1->ne[1];
const int k = src0->ne[0];
if (n == 1 || !ggml_is_contiguous(src1) || offset1 % backend_ctx->alignment != 0) {
cl_kernel gemv = backend_ctx->kernel_gemv_noshuffle_q2_0_f32;
const int stride_d = dst->ne[0];
const int stride_b = src1->nb[1]/sizeof(float);
CL_CHECK(clSetKernelArg(gemv, 0, sizeof(cl_mem), &extra0->qs));
CL_CHECK(clSetKernelArg(gemv, 1, sizeof(cl_mem), &extra0->d));
CL_CHECK(clSetKernelArg(gemv, 2, sizeof(cl_mem), &extra1->data_device));
CL_CHECK(clSetKernelArg(gemv, 3, sizeof(cl_ulong), &offset1));
CL_CHECK(clSetKernelArg(gemv, 4, sizeof(cl_mem), &extrad->data_device));
CL_CHECK(clSetKernelArg(gemv, 5, sizeof(cl_ulong), &offsetd));
CL_CHECK(clSetKernelArg(gemv, 6, sizeof(int), &k));
CL_CHECK(clSetKernelArg(gemv, 7, sizeof(int), &m));
CL_CHECK(clSetKernelArg(gemv, 8, sizeof(int), &stride_b));
CL_CHECK(clSetKernelArg(gemv, 9, sizeof(int), &stride_d));
const size_t nsg = backend_ctx->q2_0_mv_nsg;
size_t global[] = { (size_t)CEIL_DIV(m, 128)*64, (size_t)n*nsg, 1 };
size_t local[] = { 64, nsg, 1 };
backend_ctx->enqueue_ndrange_kernel(gemv, 3, global, local, dst);
return;
}
if (ggml_cl_mul_mat_q2_0_dp4a(backend, src0, src1, dst)) {
return;
}
cl_context context = backend_ctx->context;
cl_int err;
cl_buffer_region region = { (size_t)offset1, (size_t)k*n*sizeof(float) };
cl_mem b_sub_buf = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, &region, &err);
CL_CHECK(err);
cl_image_format img_fmt = { CL_RGBA, CL_FLOAT };
cl_image_desc 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_mem b_img = clCreateImage(context, CL_MEM_READ_ONLY, &img_fmt, &img_desc, nullptr, &err);
CL_CHECK(err);
const int padded_n = CEIL_DIV(n, 8)*8;
region.origin = 0;
region.size = (size_t)k*padded_n*sizeof(cl_half);
backend_ctx->prealloc_act_trans.allocate(context, region.size);
cl_mem b_sub_buf_trans = clCreateSubBuffer(backend_ctx->prealloc_act_trans.buffer, 0, CL_BUFFER_CREATE_TYPE_REGION, &region, &err);
CL_CHECK(err);
img_fmt = { CL_RGBA, CL_HALF_FLOAT };
img_desc.image_width = (size_t)k*padded_n/4;
img_desc.buffer = b_sub_buf_trans;
cl_mem b_img_trans = clCreateImage(context, 0, &img_fmt, &img_desc, nullptr, &err);
CL_CHECK(err);
const int height_b = std::max(n/4, 1);
const int width_b = k/4;
const int padded_height_b = padded_n/4;
cl_kernel kernel = backend_ctx->kernel_transpose_32_16;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &b_img));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &b_img_trans));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), &height_b));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), &width_b));
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), &padded_height_b));
size_t t_global[] = { (size_t)width_b, (size_t)padded_height_b };
size_t t_local[] = { 1, 16 };
backend_ctx->enqueue_ndrange_kernel(kernel, 2, t_global, t_local, dst);
kernel = backend_ctx->kernel_gemm_noshuffle_q2_0_f32;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0->qs));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0->d));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &b_img_trans));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extrad->data_device));
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_ulong), &offsetd));
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), &m));
CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), &padded_n));
CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), &k));
CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), &n));
size_t global[] = { (size_t)CEIL_DIV(n, 8), (size_t)CEIL_DIV(m, 512)*128, 1 };
size_t local[] = { 1, 128, 1 };
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global, local, dst);
CL_CHECK(clReleaseMemObject(b_img_trans));
CL_CHECK(clReleaseMemObject(b_sub_buf_trans));
CL_CHECK(clReleaseMemObject(b_img));
CL_CHECK(clReleaseMemObject(b_sub_buf));
}
static void ggml_cl_mul_mat_q4_0_f32_adreno_ila(ggml_backend_t backend, const ggml_tensor * src0,
const ggml_tensor * src1, ggml_tensor * dst) {
GGML_ASSERT(src0);
@@ -27245,6 +27705,7 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
// Follow view_src to reach the real SoA extra.
const ggml_tensor * soa0_src = src0->view_src != nullptr ? src0->view_src : src0;
ggml_tensor_extra_cl_q1_0 * extra0_q1_0 = (ggml_tensor_extra_cl_q1_0 *)src0->extra;
ggml_tensor_extra_cl_q2_0 * extra0_q2_0 = (ggml_tensor_extra_cl_q2_0 *)soa0_src->extra;
ggml_tensor_extra_cl_q4_0 * extra0_q4_0 = (ggml_tensor_extra_cl_q4_0 *)soa0_src->extra;
ggml_tensor_extra_cl_q4_1 * extra0_q4_1 = (ggml_tensor_extra_cl_q4_1 *)soa0_src->extra;
ggml_tensor_extra_cl_q5_0 * extra0_q5_0 = (ggml_tensor_extra_cl_q5_0 *)soa0_src->extra;
@@ -27615,6 +28076,12 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
return;
}
// q2_0 x fp32
if (src0t == GGML_TYPE_Q2_0 && src1t == GGML_TYPE_F32 && extra0_q2_0->noshuffle) {
ggml_cl_mul_mat_q2_0_f32_adreno(backend, src0, src1, dst);
return;
}
// q4_0 x fp32
if(src0t == GGML_TYPE_Q4_0 && src1t == GGML_TYPE_F32) {
ggml_cl_mul_mat_q4_0_f32_adreno(backend, src0, src1, dst);
@@ -27932,6 +28399,49 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
return;
}
case GGML_TYPE_Q2_0: {
if (ne11 < 32) {
break;
}
if (!ggml_is_contiguous(src0) || !ggml_is_contiguous(src1)) {
break;
}
kernel = backend_ctx->kernel_mul_mm_q2_0_f32_l4_lm;
nth0 = (backend_ctx->gpu_family == INTEL) ? 64 : 128;
offset0 = src0->view_offs;
int batch_stride_a = ne00*ne01;
int batch_stride_b = ne10*ne11;
int batch_stride_d = ne0*ne1;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q2_0->qs));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q2_0->d));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_ulong), &offset0));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra1->data_device));
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_ulong), &offset1));
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), &extrad->data_device));
CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_ulong), &offsetd));
CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), &ne00));
CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), &ne01));
CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), &ne02));
CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int), &ne11));
CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &ne12));
CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), &ne00));
CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &ne10));
CL_CHECK(clSetKernelArg(kernel, 14, sizeof(int), &ne0));
CL_CHECK(clSetKernelArg(kernel, 15, sizeof(int), &batch_stride_a));
CL_CHECK(clSetKernelArg(kernel, 16, sizeof(int), &batch_stride_b));
CL_CHECK(clSetKernelArg(kernel, 17, sizeof(int), &batch_stride_d));
CL_CHECK(clSetKernelArg(kernel, 18, sizeof(int), &r2));
CL_CHECK(clSetKernelArg(kernel, 19, sizeof(int), &r3));
size_t global_work_size[] = {(size_t)(CEIL_DIV(ne01, 64)*nth0), (size_t)CEIL_DIV(ne11, 64), (size_t)ne12*ne13};
size_t local_work_size[] = {(size_t)nth0, 1, 1};
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
return;
}
case GGML_TYPE_Q4_0: {
if (ne11 < 32) {
break;
@@ -28914,6 +29424,68 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
#endif // GGML_OPENCL_SOA_Q
break;
}
case GGML_TYPE_Q2_0: {
if (backend_ctx->gpu_family == INTEL) {
nth0 = 16;
nth1 = 1;
ndst = 4;
} else if (backend_ctx->gpu_family == ADRENO) {
nth0 = 64;
nth1 = 1;
ndst = 4;
} else {
GGML_ASSERT(false && "TODO: Unknown GPU");
}
#ifdef GGML_OPENCL_SOA_Q
kernel = backend_ctx->kernel_mul_mv_q2_0_f32_flat;
offset0 = src0->view_offs;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q2_0->qs));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q2_0->d));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_ulong), &offset0));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra1->data_device));
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_ulong), &offset1));
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), &extrad->data_device));
CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_ulong), &offsetd));
CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), &ne00));
CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), &ne01));
CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_ulong), &nb01));
CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_ulong), &nb02));
CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_ulong), &nb03));
CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), &ne12));
CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_ulong), &nb11));
CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_ulong), &nb12));
CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_ulong), &nb13));
CL_CHECK(clSetKernelArg(kernel, 16, sizeof(int), &ne0));
CL_CHECK(clSetKernelArg(kernel, 17, sizeof(int), &ne1));
CL_CHECK(clSetKernelArg(kernel, 18, sizeof(int), &r2));
CL_CHECK(clSetKernelArg(kernel, 19, sizeof(int), &r3));
#else
kernel = backend_ctx->kernel_mul_mv_q2_0_f32;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0->data_device));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offset0));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra1->data_device));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offset1));
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extrad->data_device));
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &offsetd));
CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), &ne00));
CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), &ne01));
CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_ulong), &nb01));
CL_CHECK(clSetKernelArg(kernel, 9, sizeof(cl_ulong), &nb02));
CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_ulong), &nb03));
CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &ne12));
CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_ulong), &nb11));
CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_ulong), &nb12));
CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_ulong), &nb13));
CL_CHECK(clSetKernelArg(kernel, 15, sizeof(int), &ne0));
CL_CHECK(clSetKernelArg(kernel, 16, sizeof(int), &ne1));
CL_CHECK(clSetKernelArg(kernel, 17, sizeof(int), &r2));
CL_CHECK(clSetKernelArg(kernel, 18, sizeof(int), &r3));
#endif // GGML_OPENCL_SOA_Q
break;
}
case GGML_TYPE_Q4_0:
// This should have been satisfied.
GGML_ASSERT(ne11 == ne1);
@@ -29951,6 +30523,7 @@ static void ggml_cl_mul_mat(ggml_backend_t backend, const ggml_tensor * src0, co
src0t == GGML_TYPE_Q5_1 ||
src0t == GGML_TYPE_Q8_0 ||
src0t == GGML_TYPE_Q1_0 ||
src0t == GGML_TYPE_Q2_0 ||
src0t == GGML_TYPE_IQ4_NL ||
src0t == GGML_TYPE_Q2_K || src0t == GGML_TYPE_Q3_K || src0t == GGML_TYPE_IQ2_S || src0t == GGML_TYPE_IQ2_XS || src0t == GGML_TYPE_IQ2_XXS) {
// Each SIMD group produces N_DST values in the result. Assuming each
ggml/src/ggml-opencl/kernels/cvt.cl
+33
View File
@@ -3055,3 +3055,36 @@ kernel void kernel_restore_block_iq2_xxs(
q[3] = (ushort)(sas >> 16);
}
}
struct block_q2_0 {
half d;
uchar qs[16];
};
kernel void kernel_convert_block_q2_0(
global struct block_q2_0 * src0,
global uchar * dst_qs,
global half * dst_d,
ulong n_blocks
) {
const ulong i = get_global_id(0);
if (i >= n_blocks) { return; }
dst_d[i] = src0[i].d;
for (int j = 0; j < 16; ++j) {
dst_qs[16*i + j] = src0[i].qs[j];
}
}
kernel void kernel_restore_block_q2_0(
global uchar * src_qs,
global half * src_d,
global struct block_q2_0 * dst,
ulong n_blocks
) {
const ulong i = get_global_id(0);
if (i >= n_blocks) { return; }
dst[i].d = src_d[i];
for (int j = 0; j < 16; ++j) {
dst[i].qs[j] = src_qs[16*i + j];
}
}
ggml/src/ggml-opencl/kernels/gemm_noshuffle_q2_0_f32.cl
+80
View File
@@ -0,0 +1,80 @@
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
#ifdef cl_qcom_reqd_sub_group_size
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
__attribute__((qcom_reqd_sub_group_size("full")))
#endif
kernel void kernel_gemm_noshuffle_q2_0_f32(
global const uchar * src0_qs,
global const half * src0_d,
read_only image1d_buffer_t src1,
global float * dst,
ulong offsetd,
int m,
int n,
int k,
int n_no_padding
) {
dst = (global float *)((global char *)dst + offsetd);
const int gy = get_global_id(0);
const int row = get_global_id(1) * 4;
const int n_4 = n / 4;
if (row >= m) {
return;
}
float8 c0 = 0, c1 = 0, c2 = 0, c3 = 0;
for (int i = 0; i < k; i += 16) {
const float4 scale = convert_float4(vload4(0, src0_d + row + (i / 64) * m));
for (int l = 0; l < 16; l += 4) {
const int ki = i + l;
const uchar4 q = vload4(0, src0_qs + row + (ki / 4) * m);
#pragma unroll
for (int j = 0; j < 4; ++j) {
float8 b;
b.s0123 = convert_float4(read_imageh(src1, gy * 2 + (ki + j) * n_4));
b.s4567 = convert_float4(read_imageh(src1, gy * 2 + 1 + (ki + j) * n_4));
const int4 v = convert_int4((q >> (uchar4)(2*j)) & (uchar4)(3)) - (int4)(1);
const float4 w = convert_float4(v) * scale;
c0 += b * w.s0;
c1 += b * w.s1;
c2 += b * w.s2;
c3 += b * w.s3;
}
}
}
int idx = gy * 8 * m + row;
if (idx + 3 < m * n_no_padding) {
vstore4((float4)(c0.s0, c1.s0, c2.s0, c3.s0), 0, dst + idx);
idx += m;
}
if (idx + 3 < m * n_no_padding) {
vstore4((float4)(c0.s1, c1.s1, c2.s1, c3.s1), 0, dst + idx);
idx += m;
}
if (idx + 3 < m * n_no_padding) {
vstore4((float4)(c0.s2, c1.s2, c2.s2, c3.s2), 0, dst + idx);
idx += m;
}
if (idx + 3 < m * n_no_padding) {
vstore4((float4)(c0.s3, c1.s3, c2.s3, c3.s3), 0, dst + idx);
idx += m;
}
if (idx + 3 < m * n_no_padding) {
vstore4((float4)(c0.s4, c1.s4, c2.s4, c3.s4), 0, dst + idx);
idx += m;
}
if (idx + 3 < m * n_no_padding) {
vstore4((float4)(c0.s5, c1.s5, c2.s5, c3.s5), 0, dst + idx);
idx += m;
}
if (idx + 3 < m * n_no_padding) {
vstore4((float4)(c0.s6, c1.s6, c2.s6, c3.s6), 0, dst + idx);
idx += m;
}
if (idx + 3 < m * n_no_padding) {
vstore4((float4)(c0.s7, c1.s7, c2.s7, c3.s7), 0, dst + idx);
}
}
ggml/src/ggml-opencl/kernels/gemm_noshuffle_q2_0_q8_1_dp4a.cl
+127
View File
@@ -0,0 +1,127 @@
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
#pragma OPENCL EXTENSION cl_khr_subgroups : enable
#ifdef cl_khr_integer_dot_product
#pragma OPENCL EXTENSION cl_khr_integer_dot_product : enable
#endif
#ifndef TILESIZE_N
#define TILESIZE_N 32
#endif
inline uint q2_0_pack(uchar q) {
return (((uint)(q & 3) - 1u) & 255u)
| (((((uint)q >> 2) & 3u) - 1u) & 255u) << 8
| (((((uint)q >> 4) & 3u) - 1u) & 255u) << 16
| (((((uint)q >> 6) & 3u) - 1u) & 255u) << 24;
}
// Pass staged operands by value; local dp4a operands miscompile on X2.
inline int dot8_q8a_v(uint8 qw, uint4 a0, uint4 a1) {
int r = 0;
r = dot_acc_sat_4x8packed_ss_int(qw.s0, a0.x, r);
r = dot_acc_sat_4x8packed_ss_int(qw.s1, a0.y, r);
r = dot_acc_sat_4x8packed_ss_int(qw.s2, a0.z, r);
r = dot_acc_sat_4x8packed_ss_int(qw.s3, a0.w, r);
r = dot_acc_sat_4x8packed_ss_int(qw.s4, a1.x, r);
r = dot_acc_sat_4x8packed_ss_int(qw.s5, a1.y, r);
r = dot_acc_sat_4x8packed_ss_int(qw.s6, a1.z, r);
r = dot_acc_sat_4x8packed_ss_int(qw.s7, a1.w, r);
return r;
}
__attribute__((qcom_wave_pair_mode(1)))
kernel void kernel_gemm_noshuffle_q2_0_q8_1_dp4a(
__global const uchar * src0_qs,
__global const half * src0_d,
__global const uint * src1_qa,
__global const half * src1_da,
__global float * dst,
ulong offsetd,
int m,
int n_no_padding,
int k
) {
dst = (global float *)((global char *)dst + offsetd);
const uint lid = get_local_id(0);
const uint block_id_m = get_global_id(1);
const uint block_id_n = get_global_id(2);
const uint row = block_id_m * 64 + lid;
const uint col_base = block_id_n * TILESIZE_N;
const bool row_valid = row < (uint)m;
const uint rrow = row_valid ? row : 0;
const uint k_u = (uint)k >> 2;
const uint k_b = (uint)k >> 5;
__local uint4 sh_qa4[TILESIZE_N][2];
__local half sh_d[TILESIZE_N];
#define NGROUPS (TILESIZE_N / 4)
float4 acc[NGROUPS];
#pragma unroll
for (int g = 0; g < NGROUPS; ++g) acc[g] = (float4)(0.0f);
for (uint step = 0; step < (uint)k; step += 32) {
const uint sub = step >> 5;
const uint ib = sub >> 1;
const float d_w = (float)src0_d[rrow + ib*(uint)m];
const uint qsb = rrow + sub*8u*(uint)m;
uint8 qw;
qw.s0 = q2_0_pack(src0_qs[qsb + 0u*(uint)m]);
qw.s1 = q2_0_pack(src0_qs[qsb + 1u*(uint)m]);
qw.s2 = q2_0_pack(src0_qs[qsb + 2u*(uint)m]);
qw.s3 = q2_0_pack(src0_qs[qsb + 3u*(uint)m]);
qw.s4 = q2_0_pack(src0_qs[qsb + 4u*(uint)m]);
qw.s5 = q2_0_pack(src0_qs[qsb + 5u*(uint)m]);
qw.s6 = q2_0_pack(src0_qs[qsb + 6u*(uint)m]);
qw.s7 = q2_0_pack(src0_qs[qsb + 7u*(uint)m]);
for (uint idx = lid; idx < TILESIZE_N * 2; idx += 64) {
const uint t = idx >> 1;
const uint v = idx & 1;
const uint c = col_base + t;
sh_qa4[t][v] = (c < (uint)n_no_padding)
? vload4(0, src1_qa + c * k_u + (step >> 2) + (v << 2))
: (uint4)(0u);
}
if (lid < TILESIZE_N) {
const uint c = col_base + lid;
sh_d[lid] = (c < (uint)n_no_padding) ? src1_da[c * k_b + sub] : (half)0;
}
barrier(CLK_LOCAL_MEM_FENCE);
#define LD4(arr, b) ((float4)((float)arr[(b)+0], (float)arr[(b)+1], (float)arr[(b)+2], (float)arr[(b)+3]))
#pragma unroll
for (int g = 0; g < NGROUPS; ++g) {
const int b = g * 4;
float4 rf;
#define TDOT4(T) dot8_q8a_v(qw, (uint4)(sh_qa4[T][0]), (uint4)(sh_qa4[T][1]))
rf.s0 = (float)TDOT4(b+0); rf.s1 = (float)TDOT4(b+1);
rf.s2 = (float)TDOT4(b+2); rf.s3 = (float)TDOT4(b+3);
acc[g] += d_w * LD4(sh_d, b) * rf;
}
#undef TDOT4
#undef LD4
barrier(CLK_LOCAL_MEM_FENCE);
}
if (!row_valid) {
return;
}
#pragma unroll
for (int g = 0; g < NGROUPS; ++g) {
const uint b = (uint)(g * 4);
const float4 a = acc[g];
const uint c0 = col_base + b;
if (c0 + 0 < (uint)n_no_padding) dst[(c0 + 0) * (uint)m + row] = a.s0;
if (c0 + 1 < (uint)n_no_padding) dst[(c0 + 1) * (uint)m + row] = a.s1;
if (c0 + 2 < (uint)n_no_padding) dst[(c0 + 2) * (uint)m + row] = a.s2;
if (c0 + 3 < (uint)n_no_padding) dst[(c0 + 3) * (uint)m + row] = a.s3;
}
#undef NGROUPS
}
ggml/src/ggml-opencl/kernels/gemv_noshuffle_q2_0_f32.cl
+64
View File
@@ -0,0 +1,64 @@
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
#ifndef Q2_0_MV_NSG
#define Q2_0_MV_NSG 8
#endif
kernel void kernel_gemv_noshuffle_q2_0_f32(
global const uchar * src0_qs,
global const half * src0_d,
global const float * src1,
ulong offset1,
global float * dst,
ulong offsetd,
int ne00,
int ne01,
int ne10,
int ne0
) {
src1 = (global const float *)((global const char *)src1 + offset1);
dst = (global float *)((global char *)dst + offsetd);
const uint m = (uint)ne01;
const uint mh = m >> 1;
const uint lid = get_local_id(0);
const uint sgi = get_local_id(1);
const uint col = get_group_id(1);
const uint j = get_group_id(0)*64u + lid;
const uint row = j << 1;
global const float * y = src1 + (ulong)col * (uint)ne10;
float2 sum = 0.0f;
if (j < mh) {
global const ushort * qs = (global const ushort *)src0_qs;
for (uint ib = sgi; ib < (uint)ne00/64; ib += Q2_0_MV_NSG) {
const float2 d = convert_float2(vload2(j + ib*mh, src0_d));
float2 acc = 0.0f;
for (uint u = 0; u < 16; ++u) {
const uint group = ib*16 + u;
const uint q = qs[j + group*mh];
const float4 yv = vload4(group, y);
const int4 v0 = convert_int4(((uint4)(q) >> (uint4)(0, 2, 4, 6)) & (uint4)(3)) - (int4)(1);
const int4 v1 = convert_int4(((uint4)(q) >> (uint4)(8, 10, 12, 14)) & (uint4)(3)) - (int4)(1);
acc.s0 += dot(yv, convert_float4(v0));
acc.s1 += dot(yv, convert_float4(v1));
}
sum += d*acc;
}
}
#if Q2_0_MV_NSG > 1
local float2 part[Q2_0_MV_NSG][64];
part[sgi][lid] = sum;
barrier(CLK_LOCAL_MEM_FENCE);
if (sgi != 0) {
return;
}
for (uint s = 1; s < Q2_0_MV_NSG; ++s) {
sum += part[s][lid];
}
#endif
if (j < mh) {
vstore2(sum, 0, dst + (ulong)col*(uint)ne0 + row);
}
}
ggml/src/ggml-opencl/kernels/mul_mm_q2_0_f32_l4_lm.cl
+171
View File
@@ -0,0 +1,171 @@
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
#ifdef cl_intel_required_subgroup_size
#define INTEL_GPU 1
#endif
#define QK2_0 64
struct block_q2_0 {
half d;
uchar qs[16];
};
#define LOAD_VEC_A 4
#define LOAD_VEC_B 4
#define BM 64
#define BN 64
#define BK 32
#ifdef INTEL_GPU
#define TM 8
#define TN 8
#else
#define TM 4
#define TN 8
#endif
kernel void kernel_mul_mm_q2_0_f32_l4_lm(
global uchar * src0_qs,
global half * src0_d,
ulong offset0,
global float4 * src1,
ulong offset1,
global float * dst,
ulong offsetd,
int ne00,
int ne01,
int ne02,
int ne11,
int ne12,
int stride_a,
int stride_b,
int stride_d,
int batch_stride_a,
int batch_stride_b,
int batch_stride_d,
int r2,
int r3
) {
src1 = (global float4*)((global char*)src1 + offset1);
dst = (global float *)((global char*)dst + offsetd);
local float buf_a[BM * BK];
local float buf_b[BN * BK];
const int batch_idx = get_global_id(2);
const int i13 = batch_idx / ne12;
const int i12 = batch_idx % ne12;
const int i03 = i13 / r3;
const int i02 = i12 / r2;
const int batch_idx_a = i03 * ne02 + i02;
const int ir = get_group_id(0);
const int ic = get_group_id(1);
const int tid = get_local_id(0);
const int th_r = tid % (BM / TM);
const int th_c = tid / (BM / TM);
const int loadr_a = get_local_id(0) % (BK / LOAD_VEC_A);
const int loadc_a = get_local_id(0) / (BK / LOAD_VEC_A);
const int loadr_b = get_local_id(0) % (BK / LOAD_VEC_B);
const int loadc_b = get_local_id(0) / (BK / LOAD_VEC_B);
const int loadstride_a = get_local_size(0) * LOAD_VEC_A / BK;
const int loadstride_b = get_local_size(0) * LOAD_VEC_B / BK;
int pos_a = batch_idx_a * batch_stride_a + ir * BM * stride_a;
int pos_b = (batch_idx * batch_stride_b + ic * BN * stride_b) / LOAD_VEC_B;
float sums[TM * TN];
float cache_a[TM];
float cache_b[TN];
for (int i = 0; i < TM * TN; i++) {
sums[i] = 0.0f;
}
for (int block = 0; block < ne00; block += BK) {
for (int l = 0; l < BM; l += loadstride_a) {
if (ir*BM + loadc_a + l < ne01) {
int idx = pos_a + (loadc_a + l) * stride_a + loadr_a * LOAD_VEC_A;
int ib = idx / QK2_0;
ib += offset0/sizeof(struct block_q2_0);
int e = idx % QK2_0;
const uint q = src0_qs[16*ib + e/4];
const int4 v = convert_int4(((uint4)(q) >> (uint4)(0, 2, 4, 6)) & (uint4)(3)) - (int4)(1);
const float4 v1 = (float)src0_d[ib] * convert_float4(v);
buf_a[(loadr_a * LOAD_VEC_A + 0) * BM + loadc_a + l] = v1.s0;
buf_a[(loadr_a * LOAD_VEC_A + 1) * BM + loadc_a + l] = v1.s1;
buf_a[(loadr_a * LOAD_VEC_A + 2) * BM + loadc_a + l] = v1.s2;
buf_a[(loadr_a * LOAD_VEC_A + 3) * BM + loadc_a + l] = v1.s3;
} else {
buf_a[(loadr_a * LOAD_VEC_A + 0) * BM + loadc_a + l] = 0.0f;
buf_a[(loadr_a * LOAD_VEC_A + 1) * BM + loadc_a + l] = 0.0f;
buf_a[(loadr_a * LOAD_VEC_A + 2) * BM + loadc_a + l] = 0.0f;
buf_a[(loadr_a * LOAD_VEC_A + 3) * BM + loadc_a + l] = 0.0f;
}
}
for (int l = 0; l < BN; l += loadstride_b) {
if (ic*BN + loadc_b + l < ne11) {
int idx = pos_b + (loadc_b + l) * stride_b / LOAD_VEC_B + loadr_b;
buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = src1[idx].s0;
buf_b[(loadr_b * LOAD_VEC_B + 1) * BN + loadc_b + l] = src1[idx].s1;
buf_b[(loadr_b * LOAD_VEC_B + 2) * BN + loadc_b + l] = src1[idx].s2;
buf_b[(loadr_b * LOAD_VEC_B + 3) * BN + loadc_b + l] = src1[idx].s3;
} else {
buf_b[(loadr_b * LOAD_VEC_B + 0) * BN + loadc_b + l] = 0.0f;
buf_b[(loadr_b * LOAD_VEC_B + 1) * BN + loadc_b + l] = 0.0f;
buf_b[(loadr_b * LOAD_VEC_B + 2) * BN + loadc_b + l] = 0.0f;
buf_b[(loadr_b * LOAD_VEC_B + 3) * BN + loadc_b + l] = 0.0f;
}
}
barrier(CLK_LOCAL_MEM_FENCE);
pos_a += BK;
pos_b += BK / LOAD_VEC_B;
for (int i = 0; i < BK; i++) {
for (int j = 0; j < TM; j++) {
cache_a[j] = buf_a[(i) * BM + th_r * TM + j];
}
for (int j = 0; j < TN; j++) {
cache_b[j] = buf_b[(i) * BN + th_c * TN + j];
}
for (int cc = 0; cc < TN; cc++) {
for (int cr = 0; cr < TM; cr++) {
const int sums_idx = cc*TM + cr;
sums[sums_idx] = mad(cache_a[cr], cache_b[cc], sums[sums_idx]);
}
}
}
barrier(CLK_LOCAL_MEM_FENCE);
}
const int dr = ir * BM + th_r * TM;
const int dc = ic * BN + th_c * TN;
const ulong offsets = batch_idx * batch_stride_d;
for (int cc = 0; cc < TN; cc++) {
for (int cr = 0; cr < TM; cr++) {
if (dr + cr < ne01 && dc + cc < ne11) {
dst[offsets + (dc + cc) * stride_d + dr + cr] = sums[cc * TM + cr];
}
}
}
}
ggml/src/ggml-opencl/kernels/mul_mv_q2_0_f32.cl
+134
View File
@@ -0,0 +1,134 @@
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
#ifdef cl_intel_required_subgroup_size
#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable
#define INTEL_GPU 1
#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16)))
#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32)))
#elif defined(cl_qcom_reqd_sub_group_size)
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
#define ADRENO_GPU 1
#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half")))
#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
#endif
#define QK2_0 64
// 00=-1, 01=0, 10=+1, 11=+2
typedef struct {
half d;
uchar qs[QK2_0/4];
} block_q2_0;
#undef N_DST
#undef N_SIMDGROUP
#undef N_SIMDWIDTH
#ifdef INTEL_GPU
#define N_DST 4 // number of rows each SIMD group works on
#define N_SIMDGROUP 1 // number of SIMD groups in a thread group
#define N_SIMDWIDTH 16 // SIMD group size
#elif defined (ADRENO_GPU)
#define N_DST 4
#define N_SIMDGROUP 1
#define N_SIMDWIDTH 64
#endif
#undef BLOCK_STRIDE
// 4 threads cover one block
#define BLOCK_STRIDE (N_SIMDWIDTH/4)
#ifdef INTEL_GPU
REQD_SUBGROUP_SIZE_16
#elif defined (ADRENO_GPU)
REQD_SUBGROUP_SIZE_64
#endif
kernel void kernel_mul_mv_q2_0_f32(
global char * src0,
ulong offset0,
global char * src1,
ulong offset1,
global char * dst,
ulong offsetd,
int ne00,
int ne01,
ulong nb01,
ulong nb02,
ulong nb03,
int ne12,
ulong nb11,
ulong nb12,
ulong nb13,
int ne0,
int ne1,
int r2,
int r3
) {
src0 = src0 + offset0;
src1 = src1 + offset1;
dst = dst + offsetd;
int ix = get_sub_group_local_id()/4; // block index
int it = get_sub_group_local_id()%4; // part of the block
int nb = ne00/QK2_0;
int r0 = get_group_id(0);
int r1 = get_group_id(1);
int im = get_group_id(2);
int first_row = (r0 * N_SIMDGROUP + get_sub_group_id()) * N_DST;
int i12 = im%ne12;
int i13 = im/ne12;
ulong offset_src0 = first_row*nb01 + (i12/r2)*nb02 + (i13/r3)*nb03;
ulong offset_src1 = r1*nb11 + (i12 )*nb12 + (i13 )*nb13;
global block_q2_0 * x = (global block_q2_0 *) (src0 + offset_src0);
global float * y = (global float *) (src1 + offset_src1);
float yl[16];
float sumf[N_DST] = {0.f};
float all_sum;
global float * y4 = y + ix * QK2_0 + 16 * it;
for (int ib = ix; ib < nb; ib += BLOCK_STRIDE) {
for (int i = 0; i < 16; ++i) {
yl[i] = y4[i];
}
global char * xrow = (global char *)(x + ib);
// only ne01 output rows exist; reading past them can pull in an inf scale
for (int row = 0; row < N_DST && first_row + row < ne01; row++) {
global block_q2_0 * xb = (global block_q2_0 *)(xrow + row*nb01);
global uchar * q = xb->qs + 4*it;
float acc = 0.f;
for (int i = 0; i < 4; ++i) {
uchar b = q[i];
acc += yl[4*i+0] * (float)(((b >> 0) & 3) - 1);
acc += yl[4*i+1] * (float)(((b >> 2) & 3) - 1);
acc += yl[4*i+2] * (float)(((b >> 4) & 3) - 1);
acc += yl[4*i+3] * (float)(((b >> 6) & 3) - 1);
}
sumf[row] += (float)xb->d * acc;
}
y4 += BLOCK_STRIDE * QK2_0;
}
global float * dst_f32 = (global float *) dst + im*ne0*ne1 + r1*ne0;
for (int row = 0; row < N_DST; ++row) {
all_sum = sub_group_reduce_add(sumf[row]);
if (first_row + row < ne01) {
if (get_sub_group_local_id() == 0) {
dst_f32[first_row + row] = all_sum;
}
}
}
}
ggml/src/ggml-opencl/kernels/mul_mv_q2_0_f32_flat.cl
+132
View File
@@ -0,0 +1,132 @@
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
#ifdef cl_intel_required_subgroup_size
#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable
#define INTEL_GPU 1
#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16)))
#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32)))
#elif defined(cl_qcom_reqd_sub_group_size)
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
#define ADRENO_GPU 1
#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half")))
#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
#endif
#define QK2_0 64
// 00=-1, 01=0, 10=+1, 11=+2
typedef struct {
half d;
uchar qs[QK2_0/4];
} block_q2_0;
#undef N_DST
#undef N_SIMDGROUP
#undef N_SIMDWIDTH
#ifdef INTEL_GPU
#define N_DST 4 // number of rows each SIMD group works on
#define N_SIMDGROUP 1 // number of SIMD groups in a thread group
#define N_SIMDWIDTH 16 // SIMD group size
#elif defined (ADRENO_GPU)
#define N_DST 4
#define N_SIMDGROUP 1
#define N_SIMDWIDTH 64
#endif
#undef BLOCK_STRIDE
// 4 threads cover one block
#define BLOCK_STRIDE (N_SIMDWIDTH/4)
#ifdef INTEL_GPU
REQD_SUBGROUP_SIZE_16
#elif defined (ADRENO_GPU)
REQD_SUBGROUP_SIZE_64
#endif
kernel void kernel_mul_mv_q2_0_f32_flat(
global uchar * src0_qs,
global half * src0_d,
ulong offset0,
global char * src1,
ulong offset1,
global char * dst,
ulong offsetd,
int ne00,
int ne01,
ulong nb01,
ulong nb02,
ulong nb03,
int ne12,
ulong nb11,
ulong nb12,
ulong nb13,
int ne0,
int ne1,
int r2,
int r3
) {
src1 = src1 + offset1;
dst = dst + offsetd;
int ix = get_sub_group_local_id()/4; // block index
int it = get_sub_group_local_id()%4; // part of the block
int nb = ne00/QK2_0;
int r0 = get_group_id(0);
int r1 = get_group_id(1);
int im = get_group_id(2);
int first_row = (r0 * N_SIMDGROUP + get_sub_group_id()) * N_DST;
int i12 = im%ne12;
int i13 = im/ne12;
ulong offset_src0 = (offset0 + first_row*nb01 + (i12/r2)*nb02 + (i13/r3)*nb03)/sizeof(block_q2_0);
ulong row_stride = nb01/sizeof(block_q2_0);
ulong offset_src1 = r1*nb11 + (i12 )*nb12 + (i13 )*nb13;
global float * y = (global float *) (src1 + offset_src1);
float yl[16];
float sumf[N_DST] = {0.f};
float all_sum;
global float * y4 = y + ix * QK2_0 + 16 * it;
for (int ib = ix; ib < nb; ib += BLOCK_STRIDE) {
for (int i = 0; i < 16; ++i) {
yl[i] = y4[i];
}
// only ne01 output rows exist; reading past them can pull in an inf scale
for (int row = 0; row < N_DST && first_row + row < ne01; row++) {
const ulong b = offset_src0 + ib + row*row_stride;
global uchar * q = src0_qs + 16*b + 4*it;
float acc = 0.f;
for (int i = 0; i < 4; ++i) {
uchar b = q[i];
acc += yl[4*i+0] * (float)(((b >> 0) & 3) - 1);
acc += yl[4*i+1] * (float)(((b >> 2) & 3) - 1);
acc += yl[4*i+2] * (float)(((b >> 4) & 3) - 1);
acc += yl[4*i+3] * (float)(((b >> 6) & 3) - 1);
}
sumf[row] += (float)src0_d[b] * acc;
}
y4 += BLOCK_STRIDE * QK2_0;
}
global float * dst_f32 = (global float *) dst + im*ne0*ne1 + r1*ne0;
for (int row = 0; row < N_DST; ++row) {
all_sum = sub_group_reduce_add(sumf[row]);
if (first_row + row < ne01) {
if (get_sub_group_local_id() == 0) {
dst_f32[first_row + row] = all_sum;
}
}
}
}