mul_mv_q4_1_f32_flat
mul_mv_q4_k_f32
mul_mv_q4_k_f32_flat
+ mul_mv_q5_0_f32
+ mul_mv_q5_0_f32_flat
+ mul_mv_q5_1_f32
+ mul_mv_q5_1_f32_flat
mul_mv_q5_k_f32
mul_mv_q5_k_f32_flat
mul_mv_q6_k_f32
mul_mm_f16_f32_l4_lm
mul_mm_q4_0_f32_l4_lm
mul_mm_q4_1_f32_l4_lm
+ mul_mm_q5_0_f32_l4_lm
+ mul_mm_q5_1_f32_l4_lm
mul_mm_q8_0_f32_l4_lm
mul_mm_iq4_nl_f32_l4_lm
mul_mm_q4_k_f32_l4_lm
cl_kernel kernel_convert_block_q4_0_trans4_ns, kernel_restore_block_q4_0_trans4_ns;
cl_kernel kernel_convert_block_q4_1, kernel_restore_block_q4_1;
cl_kernel kernel_convert_block_q4_1_trans4_ns, kernel_restore_block_q4_1_trans4_ns;
+ cl_kernel kernel_convert_block_q5_0, kernel_restore_block_q5_0;
cl_kernel kernel_convert_block_q5_0_trans4_ns, kernel_restore_block_q5_0_trans4_ns;
+ cl_kernel kernel_convert_block_q5_1, kernel_restore_block_q5_1;
cl_kernel kernel_convert_block_q5_1_trans4_ns, kernel_restore_block_q5_1_trans4_ns;
cl_kernel kernel_convert_block_q4_k_trans4_ns, kernel_restore_block_q4_k_trans4_ns;
cl_kernel kernel_convert_block_q5_k_trans4_ns, kernel_restore_block_q5_k_trans4_ns;
cl_kernel kernel_mul_mat_q4_0_f32_1d_8x_flat, kernel_mul_mat_q4_0_f32_1d_16x_flat;
cl_kernel kernel_mul_mv_q4_1_f32;
cl_kernel kernel_mul_mv_q4_1_f32_flat;
+ cl_kernel kernel_mul_mv_q5_0_f32;
+ cl_kernel kernel_mul_mv_q5_0_f32_flat;
+ cl_kernel kernel_mul_mv_q5_1_f32;
+ cl_kernel kernel_mul_mv_q5_1_f32_flat;
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;
cl_kernel kernel_mul_mm_f16_f32_l4_lm;
cl_kernel kernel_mul_mm_q4_0_f32_l4_lm;
cl_kernel kernel_mul_mm_q4_1_f32_l4_lm;
+ cl_kernel kernel_mul_mm_q5_0_f32_l4_lm;
+ cl_kernel kernel_mul_mm_q5_1_f32_l4_lm;
cl_kernel kernel_mul_mm_q8_0_f32_l4_lm;
cl_kernel kernel_mul_mm_q4_k_f32_l4_lm;
cl_kernel kernel_mul_mm_q5_k_f32_l4_lm;
CL_CHECK((backend_ctx->kernel_restore_block_q4_1 = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_q4_1", &err), err));
CL_CHECK((backend_ctx->kernel_convert_block_q4_1_trans4_ns = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_q4_1_trans4_ns", &err), err));
CL_CHECK((backend_ctx->kernel_restore_block_q4_1_trans4_ns = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_q4_1_trans4_ns", &err), err));
+ CL_CHECK((backend_ctx->kernel_convert_block_q5_0 = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_q5_0", &err), err));
+ CL_CHECK((backend_ctx->kernel_restore_block_q5_0 = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_q5_0", &err), err));
CL_CHECK((backend_ctx->kernel_convert_block_q5_0_trans4_ns = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_q5_0_trans4_ns", &err), err));
CL_CHECK((backend_ctx->kernel_restore_block_q5_0_trans4_ns = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_q5_0_trans4_ns", &err), err));
+ CL_CHECK((backend_ctx->kernel_convert_block_q5_1 = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_q5_1", &err), err));
+ CL_CHECK((backend_ctx->kernel_restore_block_q5_1 = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_q5_1", &err), err));
CL_CHECK((backend_ctx->kernel_convert_block_q5_1_trans4_ns = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_q5_1_trans4_ns", &err), err));
CL_CHECK((backend_ctx->kernel_restore_block_q5_1_trans4_ns = clCreateKernel(backend_ctx->program_cvt, "kernel_restore_block_q5_1_trans4_ns", &err), err));
CL_CHECK((backend_ctx->kernel_convert_block_q4_k_trans4_ns = clCreateKernel(backend_ctx->program_cvt, "kernel_convert_block_q4_k_trans4_ns", &err), err));
GGML_LOG_CONT(".");
}
+ // mul_mv_q5_0_f32
+ {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+ const std::string kernel_src {
+ #include "mul_mv_q5_0_f32.cl.h"
+ };
+#else
+ const std::string kernel_src = read_file("mul_mv_q5_0_f32.cl");
+#endif
+ cl_program prog =
+ build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts);
+
+ CL_CHECK((backend_ctx->kernel_mul_mv_q5_0_f32 = clCreateKernel(prog, "kernel_mul_mv_q5_0_f32", &err), err));
+ CL_CHECK(clReleaseProgram(prog));
+ GGML_LOG_CONT(".");
+ }
+
+ // mul_mv_q5_0_f32_flat
+ {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+ const std::string kernel_src {
+ #include "mul_mv_q5_0_f32_flat.cl.h"
+ };
+#else
+ const std::string kernel_src = read_file("mul_mv_q5_0_f32_flat.cl");
+#endif
+ cl_program prog =
+ build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts);
+
+ CL_CHECK((backend_ctx->kernel_mul_mv_q5_0_f32_flat = clCreateKernel(prog, "kernel_mul_mv_q5_0_f32_flat", &err), err));
+ CL_CHECK(clReleaseProgram(prog));
+ GGML_LOG_CONT(".");
+ }
+
+ // mul_mv_q5_1_f32
+ {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+ const std::string kernel_src {
+ #include "mul_mv_q5_1_f32.cl.h"
+ };
+#else
+ const std::string kernel_src = read_file("mul_mv_q5_1_f32.cl");
+#endif
+ cl_program prog =
+ build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts);
+
+ CL_CHECK((backend_ctx->kernel_mul_mv_q5_1_f32 = clCreateKernel(prog, "kernel_mul_mv_q5_1_f32", &err), err));
+ CL_CHECK(clReleaseProgram(prog));
+ GGML_LOG_CONT(".");
+ }
+
+ // mul_mv_q5_1_f32_flat
+ {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+ const std::string kernel_src {
+ #include "mul_mv_q5_1_f32_flat.cl.h"
+ };
+#else
+ const std::string kernel_src = read_file("mul_mv_q5_1_f32_flat.cl");
+#endif
+ cl_program prog =
+ build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts);
+
+ CL_CHECK((backend_ctx->kernel_mul_mv_q5_1_f32_flat = clCreateKernel(prog, "kernel_mul_mv_q5_1_f32_flat", &err), err));
+ CL_CHECK(clReleaseProgram(prog));
+ GGML_LOG_CONT(".");
+ }
+
// mul_mv_q5_k_f32
{
#ifdef GGML_OPENCL_EMBED_KERNELS
GGML_LOG_CONT(".");
}
+ // mul_mm_q5_0_f32_l4_lm
+ {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+ const std::string kernel_src {
+ #include "mul_mm_q5_0_f32_l4_lm.cl.h"
+ };
+#else
+ const std::string kernel_src = read_file("mul_mm_q5_0_f32_l4_lm.cl");
+#endif
+ cl_program prog =
+ build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts);
+
+ CL_CHECK((backend_ctx->kernel_mul_mm_q5_0_f32_l4_lm = clCreateKernel(prog, "kernel_mul_mm_q5_0_f32_l4_lm", &err), err));
+ GGML_LOG_CONT(".");
+ }
+
+ // mul_mm_q5_1_f32_l4_lm
+ {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+ const std::string kernel_src {
+ #include "mul_mm_q5_1_f32_l4_lm.cl.h"
+ };
+#else
+ const std::string kernel_src = read_file("mul_mm_q5_1_f32_l4_lm.cl");
+#endif
+ cl_program prog =
+ build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts);
+
+ CL_CHECK((backend_ctx->kernel_mul_mm_q5_1_f32_l4_lm = clCreateKernel(prog, "kernel_mul_mm_q5_1_f32_l4_lm", &err), err));
+ GGML_LOG_CONT(".");
+ }
+
// mul_mm_q8_0_f32_l4_lm
{
#ifdef GGML_OPENCL_EMBED_KERNELS
} else if (op->src[0]->type == GGML_TYPE_F32) {
return op->src[1]->type == GGML_TYPE_F32;
} else if (op->src[0]->type == GGML_TYPE_Q4_0 || op->src[0]->type == GGML_TYPE_Q4_1 ||
+ op->src[0]->type == GGML_TYPE_Q5_0 || op->src[0]->type == GGML_TYPE_Q5_1 ||
op->src[0]->type == GGML_TYPE_MXFP4 ||
op->src[0]->type == GGML_TYPE_IQ4_NL ||
op->src[0]->type == GGML_TYPE_Q4_K ||
return;
}
#endif // GGML_OPENCL_USE_ADRENO_KERNELS
- return;
+ cl_kernel kernel = backend_ctx->kernel_convert_block_q5_0;
+ cl_ulong n_blk = ggml_nelements(tensor)/ggml_blck_size(tensor->type);
+ 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->qh));
+ CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra->d));
+ CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_ulong), &n_blk));
+
+ 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(clReleaseMemObject(data_device));
+
+ tensor->extra = extra;
+ return;
}
if (tensor->type == GGML_TYPE_Q5_1) {
ggml_tensor_extra_cl * extra_orig = (ggml_tensor_extra_cl *)tensor->extra;
return;
}
#endif // GGML_OPENCL_USE_ADRENO_KERNELS
+ cl_kernel kernel = backend_ctx->kernel_convert_block_q5_1;
+ cl_ulong n_blk = ggml_nelements(tensor)/ggml_blck_size(tensor->type);
+ 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->qh));
+ CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra->d));
+ CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extra->m));
+ CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &n_blk));
+
+ 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(clReleaseMemObject(data_device));
+
+ tensor->extra = extra;
return;
}
if (tensor->type == GGML_TYPE_MXFP4) {
return;
}
#endif // GGML_OPENCL_USE_ADRENO_KERNELS
- // TODO: normal q5_0
- (void) extra;
+
+ cl_int err;
+ cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
+ ggml_nbytes(tensor), NULL, &err);
+ CL_CHECK(err);
+
+ cl_kernel kernel = backend_ctx->kernel_restore_block_q5_0;
+ CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra->qs));
+ CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra->qh));
+ CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra->d));
+ CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &data_device));
+
+ size_t global_work_size[] = {(size_t)ggml_nelements(tensor)/ggml_blck_size(tensor->type), 1, 1};
+ size_t local_work_size[] = {1, 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(clEnqueueReadBuffer(
+ queue, data_device, CL_TRUE, offset,
+ size, data, 0, NULL, NULL));
+ CL_CHECK(clReleaseMemObject(data_device));
return;
}
if (tensor->type == GGML_TYPE_Q5_1) {
return;
}
#endif // GGML_OPENCL_USE_ADRENO_KERNELS
- // TODO: normal q5_1
- (void) extra;
+ cl_int err;
+ cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
+ ggml_nbytes(tensor), NULL, &err);
+ CL_CHECK(err);
+
+ cl_kernel kernel = backend_ctx->kernel_restore_block_q5_1;
+ CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra->qs));
+ CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra->qh));
+ CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra->d));
+ CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra->m));
+ CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &data_device));
+
+ size_t global_work_size[] = {(size_t)ggml_nelements(tensor)/ggml_blck_size(tensor->type), 1, 1};
+ size_t local_work_size[] = {1, 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(clEnqueueReadBuffer(
+ queue, data_device, CL_TRUE, offset,
+ size, data, 0, NULL, NULL));
+ CL_CHECK(clReleaseMemObject(data_device));
return;
}
if (tensor->type == GGML_TYPE_MXFP4) {
#ifdef GGML_OPENCL_SOA_Q
ggml_tensor_extra_cl_q4_0 * extra0_q4_0 = (ggml_tensor_extra_cl_q4_0 *)src0->extra;
ggml_tensor_extra_cl_q4_1 * extra0_q4_1 = (ggml_tensor_extra_cl_q4_1 *)src0->extra;
+ ggml_tensor_extra_cl_q5_0 * extra0_q5_0 = (ggml_tensor_extra_cl_q5_0 *)src0->extra;
+ ggml_tensor_extra_cl_q5_1 * extra0_q5_1 = (ggml_tensor_extra_cl_q5_1 *)src0->extra;
ggml_tensor_extra_cl_mxfp4 * extra0_mxfp4 = (ggml_tensor_extra_cl_mxfp4 *)src0->extra;
ggml_tensor_extra_cl_q8_0 * extra0_q8_0 = (ggml_tensor_extra_cl_q8_0 *)src0->extra;
ggml_tensor_extra_cl_iq4_nl * extra0_iq4_nl = (ggml_tensor_extra_cl_iq4_nl *)src0->extra;
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
return;
}
+ case GGML_TYPE_Q5_0: {
+ if (ne11 < 32) {
+ break;
+ }
+ if (!ggml_is_contiguous(src0) || !ggml_is_contiguous(src1)) {
+ break;
+ }
+
+ kernel = backend_ctx->kernel_mul_mm_q5_0_f32_l4_lm;
+ nth0 = 128; // calculated as (BM*BN)/(TM*TN)
+
+ 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_q5_0->qs));
+ CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q5_0->qh));
+ CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_q5_0->d));
+ 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), &ne10)); // stride_a
+ CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &ne10)); // stride_b
+ CL_CHECK(clSetKernelArg(kernel, 14, sizeof(int), &ne01)); // stride_d
+ 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));
+
+ // 64 is block tile size BM and BN - change here when BM and BN in the kernel are changed.
+ 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_Q5_1: {
+ if (ne11 < 32) {
+ break;
+ }
+ if (!ggml_is_contiguous(src0) || !ggml_is_contiguous(src1)) {
+ break;
+ }
+
+ kernel = backend_ctx->kernel_mul_mm_q5_1_f32_l4_lm;
+ nth0 = 128; // calculated as (BM*BN)/(TM*TN)
+
+ 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_q5_1->qs));
+ CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q5_1->qh));
+ CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_q5_1->d));
+ CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra0_q5_1->m));
+ CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extra1->data_device));
+ CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &offset1));
+ 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(int), &ne00));
+ CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), &ne01));
+ CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int), &ne02));
+ CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &ne11));
+ CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), &ne12));
+ CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &ne10)); // stride_a
+ CL_CHECK(clSetKernelArg(kernel, 14, sizeof(int), &ne10)); // stride_b
+ CL_CHECK(clSetKernelArg(kernel, 15, sizeof(int), &ne01)); // stride_d
+ CL_CHECK(clSetKernelArg(kernel, 16, sizeof(int), &batch_stride_a));
+ CL_CHECK(clSetKernelArg(kernel, 17, sizeof(int), &batch_stride_b));
+ CL_CHECK(clSetKernelArg(kernel, 18, sizeof(int), &batch_stride_d));
+ CL_CHECK(clSetKernelArg(kernel, 19, sizeof(int), &r2));
+ CL_CHECK(clSetKernelArg(kernel, 20, sizeof(int), &r3));
+
+ // 64 is block tile size BM and BN - change here when BM and BN in the kernel are changed.
+ 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_Q8_0: {
if (ne11 < 32) {
break;
#endif // GGML_OPENCL_SOA_Q
break;
}
+ case GGML_TYPE_Q5_0: {
+#ifdef GGML_OPENCL_SOA_Q
+ 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");
+ }
+
+ kernel = backend_ctx->kernel_mul_mv_q5_0_f32_flat;
+
+ CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q5_0->qs));
+ CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q5_0->qh));
+ CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_q5_0->d));
+ 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), &ne10));
+ CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &ne12));
+ CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), &ne0));
+ CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &ne1));
+ CL_CHECK(clSetKernelArg(kernel, 14, sizeof(int), &r2));
+ CL_CHECK(clSetKernelArg(kernel, 15, sizeof(int), &r3));
+#else
+ 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");
+ }
+
+ kernel = backend_ctx->kernel_mul_mv_q5_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(int), &ne02));
+ CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), &ne10));
+ CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int), &ne12));
+ CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &ne0));
+ CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), &ne1));
+ CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &r2));
+ CL_CHECK(clSetKernelArg(kernel, 14, sizeof(int), &r3));
+#endif // GGML_OPENCL_SOA_Q
+ break;
+ }
+ case GGML_TYPE_Q5_1: {
+#ifdef GGML_OPENCL_SOA_Q
+ 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");
+ }
+
+ kernel = backend_ctx->kernel_mul_mv_q5_1_f32_flat;
+
+ CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q5_1->qs));
+ CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q5_1->qh));
+ CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_q5_1->d));
+ CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra0_q5_1->m));
+ CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &extra1->data_device));
+ CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_ulong), &offset1));
+ 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(int), &ne00));
+ CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), &ne01));
+ CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int), &ne02));
+ CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &ne10));
+ CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), &ne12));
+ CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &ne0));
+ CL_CHECK(clSetKernelArg(kernel, 14, sizeof(int), &ne1));
+ CL_CHECK(clSetKernelArg(kernel, 15, sizeof(int), &r2));
+ CL_CHECK(clSetKernelArg(kernel, 16, sizeof(int), &r3));
+#else
+ 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");
+ }
+
+ kernel = backend_ctx->kernel_mul_mv_q5_1_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(int), &ne02));
+ CL_CHECK(clSetKernelArg(kernel, 9, sizeof(int), &ne10));
+ CL_CHECK(clSetKernelArg(kernel, 10, sizeof(int), &ne12));
+ CL_CHECK(clSetKernelArg(kernel, 11, sizeof(int), &ne0));
+ CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), &ne1));
+ CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &r2));
+ CL_CHECK(clSetKernelArg(kernel, 14, sizeof(int), &r3));
+#endif // GGML_OPENCL_SOA_Q
+ break;
+ }
case GGML_TYPE_Q8_0: {
#ifdef GGML_OPENCL_SOA_Q
kernel = backend_ctx->kernel_mul_mv_q8_0_f32_flat;
if (src0t == GGML_TYPE_Q4_0 || src0t == GGML_TYPE_MXFP4 ||
src0t == GGML_TYPE_Q4_1 ||
+ src0t == GGML_TYPE_Q5_0 ||
+ src0t == GGML_TYPE_Q5_1 ||
src0t == GGML_TYPE_Q8_0 ||
src0t == GGML_TYPE_IQ4_NL ||
src0t == GGML_TYPE_Q2_K) {
((__global ushort8 *)(&(b->qs[0])))[0] = pre_block;
}
+//------------------------------------------------------------------------------
+// kernel_convert_block_q5_0
+// Convert the block_q5_0 format to 3 separate arrays (AOS -> SOA).
+// This kernel does not deshuffle the bits.
+//------------------------------------------------------------------------------
+kernel void kernel_convert_block_q5_0(
+ global struct block_q5_0 * src0,
+ global uchar * dst_qs,
+ global uint * dst_qh,
+ global half * dst_d,
+ ulong n_blk
+) {
+ if (get_global_id(0) >= n_blk) {
+ return;
+ }
+
+ global struct block_q5_0 * b = (global struct block_q5_0 *) src0 + get_global_id(0);
+ global uchar * qs = (global uchar *) dst_qs + (QK5_0/2)*get_global_id(0);
+ global uint * qh = (global uint *) dst_qh + get_global_id(0);
+ global half * d = (global half *) dst_d + get_global_id(0);
+
+ *d = b->d;
+ *qh = *((global uint *)(b->qh));
+
+ for (int i = 0; i < QK5_0/2; ++i) {
+ qs[i] = b->qs[i];
+ }
+}
+
+kernel void kernel_restore_block_q5_0(
+ global uchar * src_qs,
+ global uint * src_qh,
+ global half * src_d,
+ global struct block_q5_0 * dst
+) {
+ global struct block_q5_0 * b = (global struct block_q5_0 *) dst + get_global_id(0);
+ global uchar * qs = (global uchar *) src_qs + (QK5_0/2)*get_global_id(0);
+ global uint * qh = (global uint *) src_qh + get_global_id(0);
+ global half * d = (global half *) src_d + get_global_id(0);
+
+ b->d = *d;
+ *((global uint *)(b->qh)) = *qh;
+ for (int i = 0; i < QK5_0/2; ++i) {
+ b->qs[i] = qs[i];
+ }
+}
+
kernel void kernel_convert_block_q5_0_trans4_ns(
__global struct block_q5_0 * src0,
__global uint * dst_qs,
((__global ushort8 *)(&(b->qs[0])))[0] = pre_block;
}
+//------------------------------------------------------------------------------
+// kernel_convert_block_q5_1
+// Convert the block_q5_1 format to 4 separate arrays (AOS -> SOA).
+// This kernel does not deshuffle the bits.
+//------------------------------------------------------------------------------
+kernel void kernel_convert_block_q5_1(
+ global struct block_q5_1 * src0,
+ global uchar * dst_qs,
+ global uint * dst_qh,
+ global half * dst_d,
+ global half * dst_m,
+ ulong n_blk
+) {
+ if (get_global_id(0) >= n_blk) {
+ return;
+ }
+
+ global struct block_q5_1 * b = (global struct block_q5_1 *) src0 + get_global_id(0);
+ global uchar * qs = (global uchar *) dst_qs + (QK5_1/2)*get_global_id(0);
+ global uint * qh = (global uint *) dst_qh + get_global_id(0);
+ global half * d = (global half *) dst_d + get_global_id(0);
+ global half * m = (global half *) dst_m + get_global_id(0);
+
+ *d = b->d;
+ *m = b->m;
+ *qh = *((global uint *)(b->qh));
+
+ for (int i = 0; i < QK5_1/2; ++i) {
+ qs[i] = b->qs[i];
+ }
+}
+
+kernel void kernel_restore_block_q5_1(
+ global uchar * src_qs,
+ global uint * src_qh,
+ global half * src_d,
+ global half * src_m,
+ global struct block_q5_1 * dst
+) {
+ global struct block_q5_1 * b = (global struct block_q5_1 *) dst + get_global_id(0);
+ global uchar * qs = (global uchar *) src_qs + (QK5_1/2)*get_global_id(0);
+ global uint * qh = (global uint *) src_qh + get_global_id(0);
+ global half * d = (global half *) src_d + get_global_id(0);
+ global half * m = (global half *) src_m + get_global_id(0);
+
+ b->d = *d;
+ b->m = *m;
+ *((global uint *)(b->qh)) = *qh;
+ for (int i = 0; i < QK5_1/2; ++i) {
+ b->qs[i] = qs[i];
+ }
+}
+
kernel void kernel_convert_block_q5_1_trans4_ns(
__global struct block_q5_1 * src0,
__global uint * dst_qs,
--- /dev/null
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#define LOAD_VEC_A 8
+#define LOAD_VEC_B 4
+
+#define BM 64
+#define BN 64
+#define BK 32
+#define TM 4
+#define TN 8
+
+kernel void kernel_mul_mm_q5_0_f32_l4_lm(
+ global uchar4 * src0_qs,
+ global uint * src0_qh,
+ global half * src0_d,
+ 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) / LOAD_VEC_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 / LOAD_VEC_A + loadr_a;
+ int ib = idx / 4;
+ int iqs = idx % 4;
+
+ float d = (float)src0_d[ib];
+ uint qh_val = src0_qh[ib];
+
+ global uchar4 * qs_ptr = src0_qs + ib*4 + iqs;
+ uchar4 q = *qs_ptr;
+
+ uint qh_lo = qh_val >> (iqs * 4);
+ uint qh_hi = qh_val >> (iqs * 4 + 16);
+
+ uchar4 b_lo = (uchar4)((uchar)qh_lo, (uchar)(qh_lo >> 1), (uchar)(qh_lo >> 2), (uchar)(qh_lo >> 3)) & (uchar)1;
+ uchar4 b_hi = (uchar4)((uchar)qh_hi, (uchar)(qh_hi >> 1), (uchar)(qh_hi >> 2), (uchar)(qh_hi >> 3)) & (uchar)1;
+
+ float4 v1 = (convert_float4((q & (uchar)0x0F) | (b_lo << (uchar)4)) - 16.0f) * d;
+ float4 v2 = (convert_float4((q >> (uchar)4) | (b_hi << (uchar)4)) - 16.0f) * d;
+
+ buf_a[(loadr_a * 4 + 0) * BM + loadc_a + l] = v1.s0;
+ buf_a[(loadr_a * 4 + 1) * BM + loadc_a + l] = v1.s1;
+ buf_a[(loadr_a * 4 + 2) * BM + loadc_a + l] = v1.s2;
+ buf_a[(loadr_a * 4 + 3) * BM + loadc_a + l] = v1.s3;
+ buf_a[(loadr_a * 4 + 16) * BM + loadc_a + l] = v2.s0;
+ buf_a[(loadr_a * 4 + 17) * BM + loadc_a + l] = v2.s1;
+ buf_a[(loadr_a * 4 + 18) * BM + loadc_a + l] = v2.s2;
+ buf_a[(loadr_a * 4 + 19) * BM + loadc_a + l] = v2.s3;
+ } else {
+ buf_a[(loadr_a * 4 + 0) * BM + loadc_a + l] = 0.0f;
+ buf_a[(loadr_a * 4 + 1) * BM + loadc_a + l] = 0.0f;
+ buf_a[(loadr_a * 4 + 2) * BM + loadc_a + l] = 0.0f;
+ buf_a[(loadr_a * 4 + 3) * BM + loadc_a + l] = 0.0f;
+ buf_a[(loadr_a * 4 + 16) * BM + loadc_a + l] = 0.0f;
+ buf_a[(loadr_a * 4 + 17) * BM + loadc_a + l] = 0.0f;
+ buf_a[(loadr_a * 4 + 18) * BM + loadc_a + l] = 0.0f;
+ buf_a[(loadr_a * 4 + 19) * 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 / LOAD_VEC_A;
+ 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 int 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];
+ }
+ }
+ }
+}
--- /dev/null
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#define LOAD_VEC_A 8
+#define LOAD_VEC_B 4
+
+#define BM 64
+#define BN 64
+#define BK 32
+#define TM 4
+#define TN 8
+
+kernel void kernel_mul_mm_q5_1_f32_l4_lm(
+ global uchar4 * src0_qs,
+ global uint * src0_qh,
+ global half * src0_d,
+ global half * src0_m,
+ 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) / LOAD_VEC_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 / LOAD_VEC_A + loadr_a;
+ int ib = idx / 4;
+ int iqs = idx % 4;
+
+ float d = (float)src0_d[ib];
+ float m = (float)src0_m[ib];
+ uint qh_val = src0_qh[ib];
+
+ global uchar4 * qs = src0_qs + ib*4 + iqs;
+ uchar4 q = *qs;
+
+ uint qh_lo = qh_val >> (iqs * 4);
+ uint qh_hi = qh_val >> (iqs * 4 + 16);
+
+ uchar4 b_lo = (uchar4)((uchar)qh_lo, (uchar)(qh_lo >> 1), (uchar)(qh_lo >> 2), (uchar)(qh_lo >> 3)) & (uchar)1;
+ uchar4 b_hi = (uchar4)((uchar)qh_hi, (uchar)(qh_hi >> 1), (uchar)(qh_hi >> 2), (uchar)(qh_hi >> 3)) & (uchar)1;
+
+ float4 v1 = convert_float4((q & (uchar)0x0F) | (b_lo << (uchar)4)) * d + m;
+ float4 v2 = convert_float4((q >> (uchar)4) | (b_hi << (uchar)4)) * d + m;
+
+ buf_a[(loadr_a * 4 + 0) * BM + loadc_a + l] = v1.s0;
+ buf_a[(loadr_a * 4 + 1) * BM + loadc_a + l] = v1.s1;
+ buf_a[(loadr_a * 4 + 2) * BM + loadc_a + l] = v1.s2;
+ buf_a[(loadr_a * 4 + 3) * BM + loadc_a + l] = v1.s3;
+ buf_a[(loadr_a * 4 + 16) * BM + loadc_a + l] = v2.s0;
+ buf_a[(loadr_a * 4 + 17) * BM + loadc_a + l] = v2.s1;
+ buf_a[(loadr_a * 4 + 18) * BM + loadc_a + l] = v2.s2;
+ buf_a[(loadr_a * 4 + 19) * BM + loadc_a + l] = v2.s3;
+ } else {
+ buf_a[(loadr_a * 4 + 0) * BM + loadc_a + l] = 0.0f;
+ buf_a[(loadr_a * 4 + 1) * BM + loadc_a + l] = 0.0f;
+ buf_a[(loadr_a * 4 + 2) * BM + loadc_a + l] = 0.0f;
+ buf_a[(loadr_a * 4 + 3) * BM + loadc_a + l] = 0.0f;
+ buf_a[(loadr_a * 4 + 16) * BM + loadc_a + l] = 0.0f;
+ buf_a[(loadr_a * 4 + 17) * BM + loadc_a + l] = 0.0f;
+ buf_a[(loadr_a * 4 + 18) * BM + loadc_a + l] = 0.0f;
+ buf_a[(loadr_a * 4 + 19) * 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 / LOAD_VEC_A;
+ 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 int 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];
+ }
+ }
+ }
+}
--- /dev/null
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#ifdef cl_intel_subgroups
+#pragma OPENCL EXTENSION cl_intel_subgroups : enable
+#else
+#pragma OPENCL EXTENSION cl_khr_subgroups : enable
+#endif
+
+#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 QK5_0 32
+
+struct block_q5_0 {
+ half d;
+ uchar qh[4];
+ uchar qs[QK5_0 / 2];
+};
+
+inline float block_q5_0_dot_y(
+ global const struct block_q5_0 * qb_curr,
+ float sumy,
+ float16 yl,
+ int il,
+ global const float * yb
+) {
+ float d = qb_curr->d;
+
+ float4 acc = (float4)(0.0f, 0.0f, 0.0f, 0.0f);
+
+ global const ushort * qs = ((global const ushort *)((global const uchar *) qb_curr + 6 + il));
+
+ acc.s0 += yl.s0 * (qs[0] & 0x000F);
+ acc.s0 += yl.s1 * (qs[0] & 0x0F00);
+ acc.s0 += yl.s8 * (qs[0] & 0x00F0);
+ acc.s3 += yl.s9 * (qs[0] & 0xF000);
+
+ acc.s0 += yl.s2 * (qs[1] & 0x000F);
+ acc.s1 += yl.s3 * (qs[1] & 0x0F00);
+ acc.s2 += yl.sa * (qs[1] & 0x00F0);
+ acc.s3 += yl.sb * (qs[1] & 0xF000);
+
+ acc.s0 += yl.s4 * (qs[2] & 0x000F);
+ acc.s1 += yl.s5 * (qs[2] & 0x0F00);
+ acc.s2 += yl.sc * (qs[2] & 0x00F0);
+ acc.s3 += yl.sd * (qs[2] & 0xF000);
+
+ acc.s0 += yl.s6 * (qs[3] & 0x000F);
+ acc.s1 += yl.s7 * (qs[3] & 0x0F00);
+ acc.s2 += yl.se * (qs[3] & 0x00F0);
+ acc.s3 += yl.sf * (qs[3] & 0xF000);
+
+ uint qh_val = *((global const uint *)((global const uchar *) qb_curr + 2));
+ uchar qh_lo = (uchar)((qh_val >> il) & 0xFF);
+ uchar qh_hi = (uchar)((qh_val >> (il + 16)) & 0xFF);
+
+ float qh_sum = 0.0f;
+ qh_sum += yb[0] * (float)((qh_lo >> 0) & 1);
+ qh_sum += yb[1] * (float)((qh_lo >> 1) & 1);
+ qh_sum += yb[2] * (float)((qh_lo >> 2) & 1);
+ qh_sum += yb[3] * (float)((qh_lo >> 3) & 1);
+ qh_sum += yb[4] * (float)((qh_lo >> 4) & 1);
+ qh_sum += yb[5] * (float)((qh_lo >> 5) & 1);
+ qh_sum += yb[6] * (float)((qh_lo >> 6) & 1);
+ qh_sum += yb[7] * (float)((qh_lo >> 7) & 1);
+ qh_sum += yb[16] * (float)((qh_hi >> 0) & 1);
+ qh_sum += yb[17] * (float)((qh_hi >> 1) & 1);
+ qh_sum += yb[18] * (float)((qh_hi >> 2) & 1);
+ qh_sum += yb[19] * (float)((qh_hi >> 3) & 1);
+ qh_sum += yb[20] * (float)((qh_hi >> 4) & 1);
+ qh_sum += yb[21] * (float)((qh_hi >> 5) & 1);
+ qh_sum += yb[22] * (float)((qh_hi >> 6) & 1);
+ qh_sum += yb[23] * (float)((qh_hi >> 7) & 1);
+
+ return d * (acc.s0 + acc.s1 + acc.s2 + acc.s3 + 16.0f * qh_sum - 16.0f * sumy);
+}
+
+#undef N_DST
+#undef N_SIMDGROUP
+#undef N_SIMDWIDTH
+
+#ifdef INTEL_GPU
+#define N_DST 4 // each subgroup works on 4 rows
+#define N_SIMDGROUP 1 // number of subgroups in a thread group
+#define N_SIMDWIDTH 16 // assuming subgroup size is 16
+#elif defined (ADRENO_GPU)
+#define N_DST 4
+#define N_SIMDGROUP 1
+#define N_SIMDWIDTH 64
+#endif
+
+inline void mul_vec_q_n_f32(
+ global void * src0,
+ global float * src1,
+ global float * dst,
+ int ne00,
+ int ne01,
+ int ne02,
+ int ne10,
+ int ne12,
+ int ne0,
+ int ne1,
+ int r2,
+ int r3
+) {
+ const ulong nb = ne00/QK5_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 offset0 = first_row * nb + (i12/r2)*(nb*ne01) + (i13/r3)*(nb*ne01*ne02);
+
+ global struct block_q5_0 * x = (global struct block_q5_0 *) src0 + offset0;
+ global float * y = (global float *) src1 + r1*ne10 + im*ne00*ne1;
+
+ float16 yl;
+ float4 sumf = (float4)(0.f, 0.f, 0.f, 0.f);
+
+ int ix = get_sub_group_local_id()/2;
+ int il = 8*(get_sub_group_local_id()%2);
+
+ global float * yb = y + ix * QK5_0 + il;
+
+ for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/2) {
+ float sumy = 0;
+
+ sumy += yb[0];
+ sumy += yb[1];
+ sumy += yb[2];
+ sumy += yb[3];
+ sumy += yb[4];
+ sumy += yb[5];
+ sumy += yb[6];
+ sumy += yb[7];
+
+ sumy += yb[16];
+ sumy += yb[17];
+ sumy += yb[18];
+ sumy += yb[19];
+ sumy += yb[20];
+ sumy += yb[21];
+ sumy += yb[22];
+ sumy += yb[23];
+
+
+ yl.s0 = yb[0];
+ yl.s1 = yb[1]/256.f;
+
+ yl.s2 = yb[2];
+ yl.s3 = yb[3]/256.f;
+
+ yl.s4 = yb[4];
+ yl.s5 = yb[5]/256.f;
+
+ yl.s6 = yb[6];
+ yl.s7 = yb[7]/256.f;
+
+ yl.s8 = yb[16]/16.f;
+ yl.s9 = yb[17]/4096.f;
+
+ yl.sa = yb[18]/16.f;
+ yl.sb = yb[19]/4096.f;
+
+ yl.sc = yb[20]/16.f;
+ yl.sd = yb[21]/4096.f;
+
+ yl.se = yb[22]/16.f;
+ yl.sf = yb[23]/4096.f;
+
+ sumf.s0 += block_q5_0_dot_y(x+ib+0*nb, sumy, yl, il, yb);
+ sumf.s1 += block_q5_0_dot_y(x+ib+1*nb, sumy, yl, il, yb);
+ sumf.s2 += block_q5_0_dot_y(x+ib+2*nb, sumy, yl, il, yb);
+ sumf.s3 += block_q5_0_dot_y(x+ib+3*nb, sumy, yl, il, yb);
+
+ yb += QK5_0 * (N_SIMDWIDTH/2);
+ }
+
+ float4 tot = (float4)(
+ sub_group_reduce_add(sumf.s0), sub_group_reduce_add(sumf.s1),
+ sub_group_reduce_add(sumf.s2), sub_group_reduce_add(sumf.s3)
+ );
+
+ if (get_sub_group_local_id() == 0) {
+ if (first_row + 0 < ne01) {
+ dst[r1*ne0 + im*ne0*ne1 + first_row + 0] = tot.s0;
+ }
+ if (first_row + 1 < ne01) {
+ dst[r1*ne0 + im*ne0*ne1 + first_row + 1] = tot.s1;
+ }
+ if (first_row + 2 < ne01) {
+ dst[r1*ne0 + im*ne0*ne1 + first_row + 2] = tot.s2;
+ }
+ if (first_row + 3 < ne01) {
+ dst[r1*ne0 + im*ne0*ne1 + first_row + 3] = tot.s3;
+ }
+ }
+}
+
+#ifdef INTEL_GPU
+REQD_SUBGROUP_SIZE_16
+#elif defined (ADRENO_GPU)
+REQD_SUBGROUP_SIZE_64
+#endif
+kernel void kernel_mul_mv_q5_0_f32(
+ global void * src0,
+ ulong offset0,
+ global float * src1,
+ ulong offset1,
+ global float * dst,
+ ulong offsetd,
+ int ne00,
+ int ne01,
+ int ne02,
+ int ne10,
+ int ne12,
+ int ne0,
+ int ne1,
+ int r2,
+ int r3
+) {
+ src0 = (global void*)((global char*)src0 + offset0);
+ src1 = (global float*)((global char*)src1 + offset1);
+ dst = (global float*)((global char*)dst + offsetd);
+
+ mul_vec_q_n_f32(src0, src1, dst, ne00, ne01, ne02, ne10, ne12, ne0, ne1, r2, r3);
+}
--- /dev/null
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#ifdef cl_intel_subgroups
+#pragma OPENCL EXTENSION cl_intel_subgroups : enable
+#else
+#pragma OPENCL EXTENSION cl_khr_subgroups : enable
+#endif
+
+#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 QK5_0 32
+
+inline float block_q5_0_dot_y_flat(
+ global const uchar * x,
+ global const uint * qh_ptr,
+ global const half * dh,
+ float sumy,
+ float16 yl,
+ int il,
+ global const float * yb
+) {
+ float d = *dh;
+ global const ushort * qs = ((global const ushort *)(x + il));
+
+ float4 acc = (float4)(0.0f, 0.0f, 0.0f, 0.0f);
+
+ acc.s0 += yl.s0 * (qs[0] & 0x000F);
+ acc.s0 += yl.s1 * (qs[0] & 0x0F00);
+ acc.s0 += yl.s8 * (qs[0] & 0x00F0);
+ acc.s3 += yl.s9 * (qs[0] & 0xF000);
+
+ acc.s0 += yl.s2 * (qs[1] & 0x000F);
+ acc.s1 += yl.s3 * (qs[1] & 0x0F00);
+ acc.s2 += yl.sa * (qs[1] & 0x00F0);
+ acc.s3 += yl.sb * (qs[1] & 0xF000);
+
+ acc.s0 += yl.s4 * (qs[2] & 0x000F);
+ acc.s1 += yl.s5 * (qs[2] & 0x0F00);
+ acc.s2 += yl.sc * (qs[2] & 0x00F0);
+ acc.s3 += yl.sd * (qs[2] & 0xF000);
+
+ acc.s0 += yl.s6 * (qs[3] & 0x000F);
+ acc.s1 += yl.s7 * (qs[3] & 0x0F00);
+ acc.s2 += yl.se * (qs[3] & 0x00F0);
+ acc.s3 += yl.sf * (qs[3] & 0xF000);
+
+ uint qh_val = *qh_ptr;
+ uchar qh_lo = (uchar)((qh_val >> il) & 0xFF);
+ uchar qh_hi = (uchar)((qh_val >> (il + 16)) & 0xFF);
+
+ float qh_sum = 0.0f;
+ qh_sum += yb[0] * (float)((qh_lo >> 0) & 1);
+ qh_sum += yb[1] * (float)((qh_lo >> 1) & 1);
+ qh_sum += yb[2] * (float)((qh_lo >> 2) & 1);
+ qh_sum += yb[3] * (float)((qh_lo >> 3) & 1);
+ qh_sum += yb[4] * (float)((qh_lo >> 4) & 1);
+ qh_sum += yb[5] * (float)((qh_lo >> 5) & 1);
+ qh_sum += yb[6] * (float)((qh_lo >> 6) & 1);
+ qh_sum += yb[7] * (float)((qh_lo >> 7) & 1);
+ qh_sum += yb[16] * (float)((qh_hi >> 0) & 1);
+ qh_sum += yb[17] * (float)((qh_hi >> 1) & 1);
+ qh_sum += yb[18] * (float)((qh_hi >> 2) & 1);
+ qh_sum += yb[19] * (float)((qh_hi >> 3) & 1);
+ qh_sum += yb[20] * (float)((qh_hi >> 4) & 1);
+ qh_sum += yb[21] * (float)((qh_hi >> 5) & 1);
+ qh_sum += yb[22] * (float)((qh_hi >> 6) & 1);
+ qh_sum += yb[23] * (float)((qh_hi >> 7) & 1);
+
+ return d * (acc.s0 + acc.s1 + acc.s2 + acc.s3 + 16.0f * qh_sum - 16.0f * sumy);
+}
+
+#undef N_DST
+#undef N_SIMDGROUP
+#undef N_SIMDWIDTH
+
+#ifdef INTEL_GPU
+#define N_DST 4 // each subgroup works on 4 rows
+#define N_SIMDGROUP 1 // number of subgroups in a thread group
+#define N_SIMDWIDTH 16 // assuming subgroup size is 16
+#elif defined (ADRENO_GPU)
+#define N_DST 4
+#define N_SIMDGROUP 1
+#define N_SIMDWIDTH 64
+#endif
+
+inline void mul_vec_q_n_f32_flat(
+ global void * src0_qs,
+ global void * src0_qh,
+ global void * src0_d,
+ global float * src1,
+ global float * dst,
+ int ne00,
+ int ne01,
+ int ne02,
+ int ne10,
+ int ne12,
+ int ne0,
+ int ne1,
+ int r2,
+ int r3
+) {
+ const ulong nb = ne00/QK5_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 offset0 = first_row * nb + (i12/r2)*(nb*ne01) + (i13/r3)*(nb*ne01*ne02);
+
+ ulong offset0_qs = offset0 * (QK5_0/2);
+
+ global uchar * x = (global uchar *) src0_qs + offset0_qs;
+ global uint * qh = (global uint *) src0_qh + offset0;
+ global half * d = (global half *) src0_d + offset0;
+ global float * y = (global float *) src1 + r1*ne10 + im*ne00*ne1;
+
+ float16 yl;
+ float4 sumf = (float4)(0.f, 0.f, 0.f, 0.f);
+
+ int ix = get_sub_group_local_id()/2;
+ int il = 8*(get_sub_group_local_id()%2);
+
+ global float * yb = y + ix * QK5_0 + il;
+
+ for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/2) {
+ float sumy = 0;
+
+ sumy += yb[0];
+ sumy += yb[1];
+ sumy += yb[2];
+ sumy += yb[3];
+ sumy += yb[4];
+ sumy += yb[5];
+ sumy += yb[6];
+ sumy += yb[7];
+
+ sumy += yb[16];
+ sumy += yb[17];
+ sumy += yb[18];
+ sumy += yb[19];
+ sumy += yb[20];
+ sumy += yb[21];
+ sumy += yb[22];
+ sumy += yb[23];
+
+
+ yl.s0 = yb[0];
+ yl.s1 = yb[1]/256.f;
+
+ yl.s2 = yb[2];
+ yl.s3 = yb[3]/256.f;
+
+ yl.s4 = yb[4];
+ yl.s5 = yb[5]/256.f;
+
+ yl.s6 = yb[6];
+ yl.s7 = yb[7]/256.f;
+
+ yl.s8 = yb[16]/16.f;
+ yl.s9 = yb[17]/4096.f;
+
+ yl.sa = yb[18]/16.f;
+ yl.sb = yb[19]/4096.f;
+
+ yl.sc = yb[20]/16.f;
+ yl.sd = yb[21]/4096.f;
+
+ yl.se = yb[22]/16.f;
+ yl.sf = yb[23]/4096.f;
+
+ sumf.s0 += block_q5_0_dot_y_flat(x + ib*(QK5_0/2) + 0*nb*(QK5_0/2), qh + ib + 0*nb, d + ib + 0*nb, sumy, yl, il, yb);
+ sumf.s1 += block_q5_0_dot_y_flat(x + ib*(QK5_0/2) + 1*nb*(QK5_0/2), qh + ib + 1*nb, d + ib + 1*nb, sumy, yl, il, yb);
+ sumf.s2 += block_q5_0_dot_y_flat(x + ib*(QK5_0/2) + 2*nb*(QK5_0/2), qh + ib + 2*nb, d + ib + 2*nb, sumy, yl, il, yb);
+ sumf.s3 += block_q5_0_dot_y_flat(x + ib*(QK5_0/2) + 3*nb*(QK5_0/2), qh + ib + 3*nb, d + ib + 3*nb, sumy, yl, il, yb);
+
+ yb += QK5_0 * (N_SIMDWIDTH/2);
+ }
+
+ float4 tot = (float4)(
+ sub_group_reduce_add(sumf.s0), sub_group_reduce_add(sumf.s1),
+ sub_group_reduce_add(sumf.s2), sub_group_reduce_add(sumf.s3)
+ );
+
+ if (get_sub_group_local_id() == 0) {
+ if (first_row + 0 < ne01) {
+ dst[r1*ne0 + im*ne0*ne1 + first_row + 0] = tot.s0;
+ }
+ if (first_row + 1 < ne01) {
+ dst[r1*ne0 + im*ne0*ne1 + first_row + 1] = tot.s1;
+ }
+ if (first_row + 2 < ne01) {
+ dst[r1*ne0 + im*ne0*ne1 + first_row + 2] = tot.s2;
+ }
+ if (first_row + 3 < ne01) {
+ dst[r1*ne0 + im*ne0*ne1 + first_row + 3] = tot.s3;
+ }
+ }
+}
+
+#ifdef INTEL_GPU
+REQD_SUBGROUP_SIZE_16
+#elif defined (ADRENO_GPU)
+REQD_SUBGROUP_SIZE_64
+#endif
+kernel void kernel_mul_mv_q5_0_f32_flat(
+ global void * src0_qs,
+ global void * src0_qh,
+ global void * src0_d,
+ global float * src1,
+ ulong offset1,
+ global float * dst,
+ ulong offsetd,
+ int ne00,
+ int ne01,
+ int ne02,
+ int ne10,
+ int ne12,
+ int ne0,
+ int ne1,
+ int r2,
+ int r3
+) {
+ src1 = (global float*)((global char*)src1 + offset1);
+ dst = (global float*)((global char*)dst + offsetd);
+
+ mul_vec_q_n_f32_flat(src0_qs, src0_qh, src0_d, src1, dst, ne00, ne01, ne02, ne10, ne12, ne0, ne1, r2, r3);
+}
--- /dev/null
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#ifdef cl_intel_subgroups
+#pragma OPENCL EXTENSION cl_intel_subgroups : enable
+#else
+#pragma OPENCL EXTENSION cl_khr_subgroups : enable
+#endif
+
+#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 QK5_1 32
+
+struct block_q5_1 {
+ half d;
+ half m;
+ uchar qh[4];
+ uchar qs[QK5_1 / 2];
+};
+
+inline float block_q5_1_dot_y(
+ global const struct block_q5_1 * qb_curr,
+ float sumy,
+ float16 yl,
+ int il,
+ global const float * yb
+) {
+ float d = qb_curr->d;
+ float m = qb_curr->m;
+
+ float4 acc = (float4)(0.0f, 0.0f, 0.0f, 0.0f);
+
+ global const ushort * qs = ((global const ushort *)((global const uchar *) qb_curr + 8 + il));
+
+ acc.s0 += yl.s0 * (qs[0] & 0x000F);
+ acc.s0 += yl.s1 * (qs[0] & 0x0F00);
+ acc.s0 += yl.s8 * (qs[0] & 0x00F0);
+ acc.s3 += yl.s9 * (qs[0] & 0xF000);
+
+ acc.s0 += yl.s2 * (qs[1] & 0x000F);
+ acc.s1 += yl.s3 * (qs[1] & 0x0F00);
+ acc.s2 += yl.sa * (qs[1] & 0x00F0);
+ acc.s3 += yl.sb * (qs[1] & 0xF000);
+
+ acc.s0 += yl.s4 * (qs[2] & 0x000F);
+ acc.s1 += yl.s5 * (qs[2] & 0x0F00);
+ acc.s2 += yl.sc * (qs[2] & 0x00F0);
+ acc.s3 += yl.sd * (qs[2] & 0xF000);
+
+ acc.s0 += yl.s6 * (qs[3] & 0x000F);
+ acc.s1 += yl.s7 * (qs[3] & 0x0F00);
+ acc.s2 += yl.se * (qs[3] & 0x00F0);
+ acc.s3 += yl.sf * (qs[3] & 0xF000);
+
+ uint qh_val = *((global const uint *)((global const uchar *) qb_curr + 4));
+ uchar qh_lo = (uchar)((qh_val >> il) & 0xFF);
+ uchar qh_hi = (uchar)((qh_val >> (il + 16)) & 0xFF);
+
+ float qh_sum = 0.0f;
+ qh_sum += yb[0] * (float)((qh_lo >> 0) & 1);
+ qh_sum += yb[1] * (float)((qh_lo >> 1) & 1);
+ qh_sum += yb[2] * (float)((qh_lo >> 2) & 1);
+ qh_sum += yb[3] * (float)((qh_lo >> 3) & 1);
+ qh_sum += yb[4] * (float)((qh_lo >> 4) & 1);
+ qh_sum += yb[5] * (float)((qh_lo >> 5) & 1);
+ qh_sum += yb[6] * (float)((qh_lo >> 6) & 1);
+ qh_sum += yb[7] * (float)((qh_lo >> 7) & 1);
+ qh_sum += yb[16] * (float)((qh_hi >> 0) & 1);
+ qh_sum += yb[17] * (float)((qh_hi >> 1) & 1);
+ qh_sum += yb[18] * (float)((qh_hi >> 2) & 1);
+ qh_sum += yb[19] * (float)((qh_hi >> 3) & 1);
+ qh_sum += yb[20] * (float)((qh_hi >> 4) & 1);
+ qh_sum += yb[21] * (float)((qh_hi >> 5) & 1);
+ qh_sum += yb[22] * (float)((qh_hi >> 6) & 1);
+ qh_sum += yb[23] * (float)((qh_hi >> 7) & 1);
+
+ return d * (acc.s0 + acc.s1 + acc.s2 + acc.s3 + 16.0f * qh_sum) + sumy * m;
+}
+
+#undef N_DST
+#undef N_SIMDGROUP
+#undef N_SIMDWIDTH
+
+#ifdef INTEL_GPU
+#define N_DST 4 // each subgroup works on 4 rows
+#define N_SIMDGROUP 1 // number of subgroups in a thread group
+#define N_SIMDWIDTH 16 // assuming subgroup size is 16
+#elif defined (ADRENO_GPU)
+#define N_DST 4
+#define N_SIMDGROUP 1
+#define N_SIMDWIDTH 64
+#endif
+
+inline void mul_vec_q_n_f32(
+ global void * src0,
+ global float * src1,
+ global float * dst,
+ int ne00,
+ int ne01,
+ int ne02,
+ int ne10,
+ int ne12,
+ int ne0,
+ int ne1,
+ int r2,
+ int r3
+) {
+ const ulong nb = ne00/QK5_1;
+
+ 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 offset0 = first_row * nb + (i12/r2)*(nb*ne01) + (i13/r3)*(nb*ne01*ne02);
+
+ global struct block_q5_1 * x = (global struct block_q5_1 *) src0 + offset0;
+ global float * y = (global float *) src1 + r1*ne10 + im*ne00*ne1;
+
+ float16 yl;
+ float4 sumf = (float4)(0.f, 0.f, 0.f, 0.f);
+
+ int ix = get_sub_group_local_id()/2;
+ int il = 8*(get_sub_group_local_id()%2);
+
+ global float * yb = y + ix * QK5_1 + il;
+
+ for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/2) {
+ float sumy = 0;
+
+ sumy += yb[0];
+ sumy += yb[1];
+ sumy += yb[2];
+ sumy += yb[3];
+ sumy += yb[4];
+ sumy += yb[5];
+ sumy += yb[6];
+ sumy += yb[7];
+
+ sumy += yb[16];
+ sumy += yb[17];
+ sumy += yb[18];
+ sumy += yb[19];
+ sumy += yb[20];
+ sumy += yb[21];
+ sumy += yb[22];
+ sumy += yb[23];
+
+
+ yl.s0 = yb[0];
+ yl.s1 = yb[1]/256.f;
+
+ yl.s2 = yb[2];
+ yl.s3 = yb[3]/256.f;
+
+ yl.s4 = yb[4];
+ yl.s5 = yb[5]/256.f;
+
+ yl.s6 = yb[6];
+ yl.s7 = yb[7]/256.f;
+
+ yl.s8 = yb[16]/16.f;
+ yl.s9 = yb[17]/4096.f;
+
+ yl.sa = yb[18]/16.f;
+ yl.sb = yb[19]/4096.f;
+
+ yl.sc = yb[20]/16.f;
+ yl.sd = yb[21]/4096.f;
+
+ yl.se = yb[22]/16.f;
+ yl.sf = yb[23]/4096.f;
+
+ sumf.s0 += block_q5_1_dot_y(x+ib+0*nb, sumy, yl, il, yb);
+ sumf.s1 += block_q5_1_dot_y(x+ib+1*nb, sumy, yl, il, yb);
+ sumf.s2 += block_q5_1_dot_y(x+ib+2*nb, sumy, yl, il, yb);
+ sumf.s3 += block_q5_1_dot_y(x+ib+3*nb, sumy, yl, il, yb);
+
+ yb += QK5_1 * (N_SIMDWIDTH/2);
+ }
+
+ float4 tot = (float4)(
+ sub_group_reduce_add(sumf.s0), sub_group_reduce_add(sumf.s1),
+ sub_group_reduce_add(sumf.s2), sub_group_reduce_add(sumf.s3)
+ );
+
+ if (get_sub_group_local_id() == 0) {
+ if (first_row + 0 < ne01) {
+ dst[r1*ne0 + im*ne0*ne1 + first_row + 0] = tot.s0;
+ }
+ if (first_row + 1 < ne01) {
+ dst[r1*ne0 + im*ne0*ne1 + first_row + 1] = tot.s1;
+ }
+ if (first_row + 2 < ne01) {
+ dst[r1*ne0 + im*ne0*ne1 + first_row + 2] = tot.s2;
+ }
+ if (first_row + 3 < ne01) {
+ dst[r1*ne0 + im*ne0*ne1 + first_row + 3] = tot.s3;
+ }
+ }
+}
+
+#ifdef INTEL_GPU
+REQD_SUBGROUP_SIZE_16
+#elif defined (ADRENO_GPU)
+REQD_SUBGROUP_SIZE_64
+#endif
+kernel void kernel_mul_mv_q5_1_f32(
+ global void * src0,
+ ulong offset0,
+ global float * src1,
+ ulong offset1,
+ global float * dst,
+ ulong offsetd,
+ int ne00,
+ int ne01,
+ int ne02,
+ int ne10,
+ int ne12,
+ int ne0,
+ int ne1,
+ int r2,
+ int r3
+) {
+ src0 = (global void*)((global char*)src0 + offset0);
+ src1 = (global float*)((global char*)src1 + offset1);
+ dst = (global float*)((global char*)dst + offsetd);
+
+ mul_vec_q_n_f32(src0, src1, dst, ne00, ne01, ne02, ne10, ne12, ne0, ne1, r2, r3);
+}
--- /dev/null
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+#ifdef cl_intel_subgroups
+#pragma OPENCL EXTENSION cl_intel_subgroups : enable
+#else
+#pragma OPENCL EXTENSION cl_khr_subgroups : enable
+#endif
+
+#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 QK5_1 32
+
+inline float block_q5_1_dot_y_flat(
+ global const uchar * x,
+ global const uint * qh_ptr,
+ global const half * dh,
+ global const half * mh,
+ float sumy,
+ float16 yl,
+ int il,
+ global const float * yb
+) {
+ float d = *dh;
+ float m = *mh;
+ global const ushort * qs = ((global const ushort *)(x + il));
+
+ float4 acc = (float4)(0.0f, 0.0f, 0.0f, 0.0f);
+
+ acc.s0 += yl.s0 * (qs[0] & 0x000F);
+ acc.s0 += yl.s1 * (qs[0] & 0x0F00);
+ acc.s0 += yl.s8 * (qs[0] & 0x00F0);
+ acc.s3 += yl.s9 * (qs[0] & 0xF000);
+
+ acc.s0 += yl.s2 * (qs[1] & 0x000F);
+ acc.s1 += yl.s3 * (qs[1] & 0x0F00);
+ acc.s2 += yl.sa * (qs[1] & 0x00F0);
+ acc.s3 += yl.sb * (qs[1] & 0xF000);
+
+ acc.s0 += yl.s4 * (qs[2] & 0x000F);
+ acc.s1 += yl.s5 * (qs[2] & 0x0F00);
+ acc.s2 += yl.sc * (qs[2] & 0x00F0);
+ acc.s3 += yl.sd * (qs[2] & 0xF000);
+
+ acc.s0 += yl.s6 * (qs[3] & 0x000F);
+ acc.s1 += yl.s7 * (qs[3] & 0x0F00);
+ acc.s2 += yl.se * (qs[3] & 0x00F0);
+ acc.s3 += yl.sf * (qs[3] & 0xF000);
+
+ uint qh_val = *qh_ptr;
+ uchar qh_lo = (uchar)((qh_val >> il) & 0xFF);
+ uchar qh_hi = (uchar)((qh_val >> (il + 16)) & 0xFF);
+
+ float qh_sum = 0.0f;
+ qh_sum += yb[0] * (float)((qh_lo >> 0) & 1);
+ qh_sum += yb[1] * (float)((qh_lo >> 1) & 1);
+ qh_sum += yb[2] * (float)((qh_lo >> 2) & 1);
+ qh_sum += yb[3] * (float)((qh_lo >> 3) & 1);
+ qh_sum += yb[4] * (float)((qh_lo >> 4) & 1);
+ qh_sum += yb[5] * (float)((qh_lo >> 5) & 1);
+ qh_sum += yb[6] * (float)((qh_lo >> 6) & 1);
+ qh_sum += yb[7] * (float)((qh_lo >> 7) & 1);
+ qh_sum += yb[16] * (float)((qh_hi >> 0) & 1);
+ qh_sum += yb[17] * (float)((qh_hi >> 1) & 1);
+ qh_sum += yb[18] * (float)((qh_hi >> 2) & 1);
+ qh_sum += yb[19] * (float)((qh_hi >> 3) & 1);
+ qh_sum += yb[20] * (float)((qh_hi >> 4) & 1);
+ qh_sum += yb[21] * (float)((qh_hi >> 5) & 1);
+ qh_sum += yb[22] * (float)((qh_hi >> 6) & 1);
+ qh_sum += yb[23] * (float)((qh_hi >> 7) & 1);
+
+ return d * (acc.s0 + acc.s1 + acc.s2 + acc.s3 + 16.0f * qh_sum) + sumy * m;
+}
+
+#undef N_DST
+#undef N_SIMDGROUP
+#undef N_SIMDWIDTH
+
+#ifdef INTEL_GPU
+#define N_DST 4 // each subgroup works on 4 rows
+#define N_SIMDGROUP 1 // number of subgroups in a thread group
+#define N_SIMDWIDTH 16 // assuming subgroup size is 16
+#elif defined (ADRENO_GPU)
+#define N_DST 4
+#define N_SIMDGROUP 1
+#define N_SIMDWIDTH 64
+#endif
+
+inline void mul_vec_q_n_f32_flat(
+ global void * src0_qs,
+ global void * src0_qh,
+ global void * src0_d,
+ global void * src0_m,
+ global float * src1,
+ global float * dst,
+ int ne00,
+ int ne01,
+ int ne02,
+ int ne10,
+ int ne12,
+ int ne0,
+ int ne1,
+ int r2,
+ int r3
+) {
+ const ulong nb = ne00/QK5_1;
+
+ 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 offset0 = first_row * nb + (i12/r2)*(nb*ne01) + (i13/r3)*(nb*ne01*ne02);
+
+ ulong offset0_qs = offset0 * (QK5_1/2);
+
+ global uchar * x = (global uchar *) src0_qs + offset0_qs;
+ global uint * qh = (global uint *) src0_qh + offset0;
+ global half * d = (global half *) src0_d + offset0;
+ global half * ms = (global half *) src0_m + offset0;
+ global float * y = (global float *) src1 + r1*ne10 + im*ne00*ne1;
+
+ float16 yl;
+ float4 sumf = (float4)(0.f, 0.f, 0.f, 0.f);
+
+ int ix = get_sub_group_local_id()/2;
+ int il = 8*(get_sub_group_local_id()%2);
+
+ global float * yb = y + ix * QK5_1 + il;
+
+ for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/2) {
+ float sumy = 0;
+
+ sumy += yb[0];
+ sumy += yb[1];
+ sumy += yb[2];
+ sumy += yb[3];
+ sumy += yb[4];
+ sumy += yb[5];
+ sumy += yb[6];
+ sumy += yb[7];
+
+ sumy += yb[16];
+ sumy += yb[17];
+ sumy += yb[18];
+ sumy += yb[19];
+ sumy += yb[20];
+ sumy += yb[21];
+ sumy += yb[22];
+ sumy += yb[23];
+
+
+ yl.s0 = yb[0];
+ yl.s1 = yb[1]/256.f;
+
+ yl.s2 = yb[2];
+ yl.s3 = yb[3]/256.f;
+
+ yl.s4 = yb[4];
+ yl.s5 = yb[5]/256.f;
+
+ yl.s6 = yb[6];
+ yl.s7 = yb[7]/256.f;
+
+ yl.s8 = yb[16]/16.f;
+ yl.s9 = yb[17]/4096.f;
+
+ yl.sa = yb[18]/16.f;
+ yl.sb = yb[19]/4096.f;
+
+ yl.sc = yb[20]/16.f;
+ yl.sd = yb[21]/4096.f;
+
+ yl.se = yb[22]/16.f;
+ yl.sf = yb[23]/4096.f;
+
+ sumf.s0 += block_q5_1_dot_y_flat(x + ib*(QK5_1/2) + 0*nb*(QK5_1/2), qh + ib + 0*nb, d + ib + 0*nb, ms + ib + 0*nb, sumy, yl, il, yb);
+ sumf.s1 += block_q5_1_dot_y_flat(x + ib*(QK5_1/2) + 1*nb*(QK5_1/2), qh + ib + 1*nb, d + ib + 1*nb, ms + ib + 1*nb, sumy, yl, il, yb);
+ sumf.s2 += block_q5_1_dot_y_flat(x + ib*(QK5_1/2) + 2*nb*(QK5_1/2), qh + ib + 2*nb, d + ib + 2*nb, ms + ib + 2*nb, sumy, yl, il, yb);
+ sumf.s3 += block_q5_1_dot_y_flat(x + ib*(QK5_1/2) + 3*nb*(QK5_1/2), qh + ib + 3*nb, d + ib + 3*nb, ms + ib + 3*nb, sumy, yl, il, yb);
+
+ yb += QK5_1 * (N_SIMDWIDTH/2);
+ }
+
+ float4 tot = (float4)(
+ sub_group_reduce_add(sumf.s0), sub_group_reduce_add(sumf.s1),
+ sub_group_reduce_add(sumf.s2), sub_group_reduce_add(sumf.s3)
+ );
+
+ if (get_sub_group_local_id() == 0) {
+ if (first_row + 0 < ne01) {
+ dst[r1*ne0 + im*ne0*ne1 + first_row + 0] = tot.s0;
+ }
+ if (first_row + 1 < ne01) {
+ dst[r1*ne0 + im*ne0*ne1 + first_row + 1] = tot.s1;
+ }
+ if (first_row + 2 < ne01) {
+ dst[r1*ne0 + im*ne0*ne1 + first_row + 2] = tot.s2;
+ }
+ if (first_row + 3 < ne01) {
+ dst[r1*ne0 + im*ne0*ne1 + first_row + 3] = tot.s3;
+ }
+ }
+}
+
+#ifdef INTEL_GPU
+REQD_SUBGROUP_SIZE_16
+#elif defined (ADRENO_GPU)
+REQD_SUBGROUP_SIZE_64
+#endif
+kernel void kernel_mul_mv_q5_1_f32_flat(
+ global void * src0_qs,
+ global void * src0_qh,
+ global void * src0_d,
+ global void * src0_m,
+ global float * src1,
+ ulong offset1,
+ global float * dst,
+ ulong offsetd,
+ int ne00,
+ int ne01,
+ int ne02,
+ int ne10,
+ int ne12,
+ int ne0,
+ int ne1,
+ int r2,
+ int r3
+) {
+ src1 = (global float*)((global char*)src1 + offset1);
+ dst = (global float*)((global char*)dst + offsetd);
+
+ mul_vec_q_n_f32_flat(src0_qs, src0_qh, src0_d, src0_m, src1, dst, ne00, ne01, ne02, ne10, ne12, ne0, ne1, r2, r3);
+}
overview / pkg / ggml / sources / llama.cpp / commitdiff