diag
div
gelu
- gemv_noshuffle_general
- gemv_noshuffle
get_rows
glu
group_norm
im2col_f32
im2col_f16
mean
- mul_mat_Ab_Bi_8x4
mul_mv_f16_f16
mul_mv_f16_f32_1row
mul_mv_f16_f32_l4
mul_mm_q4_k_f32_l4_lm
mul_mm_q5_k_f32_l4_lm
mul_mm_q6_k_f32_l4_lm
- mul_mm_q8_0_f32_8x4
+ gemv_noshuffle_q4_0_f32
+ gemv_noshuffle_q4_0_f32_spec
+ gemm_noshuffle_q4_0_f32
gemv_noshuffle_q4_1_f32
gemm_noshuffle_q4_1_f32
gemv_noshuffle_iq4_nl_f32
gemm_noshuffle_iq4_nl_f32
- gemv_noshuffle_general_q8_0_f32
+ gemv_noshuffle_q8_0_f32
+ gemm_noshuffle_q8_0_f32
gemv_noshuffle_q4_k_f32
gemm_noshuffle_q4_k_f32
gemv_noshuffle_q6_k_f32
cl_kernel kernel_transpose_16_4x1;
// Gemm and Gemv related programs, kernels, etc
- cl_program program_CL_gemm;
- cl_program program_CL_gemv_general;
- cl_program program_CL_gemv_4096_1_11008;
- cl_program program_CL_gemv_4096_1_4096;
- cl_program program_CL_gemv_11008_1_4096;
- cl_program program_CL_gemv_32000_1_4096;
- cl_kernel CL_mul_mat_Ab_Bi_8x4;
- cl_kernel CL_mul_mat_vec_q4_0_f32_1d_4x_flat_general;
- cl_kernel CL_mul_mat_vec_q4_0_f32_1d_4x_flat_4096_1_11008;
- cl_kernel CL_mul_mat_vec_q4_0_f32_1d_4x_flat_4096_1_4096;
- cl_kernel CL_mul_mat_vec_q4_0_f32_1d_4x_flat_11008_1_4096;
- cl_kernel CL_mul_mat_vec_q4_0_f32_1d_4x_flat_32000_1_4096;
+ cl_kernel kernel_gemm_noshuffle_q4_0_f32;
+ cl_kernel kernel_gemv_noshuffle_q4_0_f32;
+ cl_kernel kernel_gemv_noshuffle_q4_0_f32_4096_1_11008;
+ cl_kernel kernel_gemv_noshuffle_q4_0_f32_4096_1_4096;
+ cl_kernel kernel_gemv_noshuffle_q4_0_f32_11008_1_4096;
+ cl_kernel kernel_gemv_noshuffle_q4_0_f32_32000_1_4096;
cl_kernel kernel_gemv_noshuffle_q4_1_f32;
cl_kernel kernel_gemm_noshuffle_q4_1_f32;
- cl_kernel kernel_mul_mm_q8_0_f32_8x4;
- cl_kernel CL_mul_mat_vec_q8_0_f32;
+ cl_kernel kernel_gemm_noshuffle_q8_0_f32;
+ cl_kernel kernel_gemv_noshuffle_q8_0_f32;
cl_kernel kernel_gemv_noshuffle_q4_k_f32;
cl_kernel kernel_gemm_noshuffle_q4_k_f32;
cl_kernel kernel_gemv_noshuffle_q6_K_f32;
" -DSIMDGROUP_WIDTH=" +
std::to_string(backend_ctx->adreno_wave_size);
if (backend_ctx->has_vector_subgroup_broadcast) {
- CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAT ";
+ CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAST ";
}
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src_CL_gemv_general {
- #include "gemv_noshuffle_general.cl.h"
+ #include "gemv_noshuffle_q4_0_f32.cl.h"
};
#else
- const std::string kernel_src_CL_gemv_general = read_file("gemv_noshuffle_general.cl");
+ const std::string kernel_src_CL_gemv_general = read_file("gemv_noshuffle_q4_0_f32.cl");
#endif
- backend_ctx->program_CL_gemv_general = build_program_from_source(
+ cl_program prog = build_program_from_source(
backend_ctx->context, backend_ctx->device, kernel_src_CL_gemv_general.c_str(), CL_gemv_compile_opts);
- CL_CHECK((backend_ctx->CL_mul_mat_vec_q4_0_f32_1d_4x_flat_general = clCreateKernel(backend_ctx->program_CL_gemv_general, "kernel_gemv_noshuffle", &err), err));
+ CL_CHECK((backend_ctx->kernel_gemv_noshuffle_q4_0_f32 = clCreateKernel(prog, "kernel_gemv_noshuffle_q4_0_f32", &err), err));
+ CL_CHECK(clReleaseProgram(prog));
GGML_LOG_CONT(".");
}
" -DSIMDGROUP_WIDTH=" +
std::to_string(backend_ctx->adreno_wave_size);
if (backend_ctx->has_vector_subgroup_broadcast) {
- CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAT ";
+ CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAST ";
}
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src_CL_gemv {
- #include "gemv_noshuffle.cl.h"
+ #include "gemv_noshuffle_q4_0_f32_spec.cl.h"
};
#else
- const std::string kernel_src_CL_gemv = read_file("gemv_noshuffle.cl");
+ const std::string kernel_src_CL_gemv = read_file("gemv_noshuffle_q4_0_f32_spec.cl");
#endif
- backend_ctx->program_CL_gemv_4096_1_4096 = build_program_from_source(
+ cl_program prog = build_program_from_source(
backend_ctx->context, backend_ctx->device, kernel_src_CL_gemv.c_str(), CL_gemv_compile_opts);
- CL_CHECK((backend_ctx->CL_mul_mat_vec_q4_0_f32_1d_4x_flat_4096_1_4096 = clCreateKernel(backend_ctx->program_CL_gemv_4096_1_4096, "kernel_gemv_noshuffle", &err), err));
+ CL_CHECK((backend_ctx->kernel_gemv_noshuffle_q4_0_f32_4096_1_4096 = clCreateKernel(prog, "kernel_gemv_noshuffle_q4_0_f32", &err), err));
+ CL_CHECK(clReleaseProgram(prog));
GGML_LOG_CONT(".");
// Gemv 2048, 16384
" -DSIMDGROUP_WIDTH=" +
std::to_string(backend_ctx->adreno_wave_size);
if (backend_ctx->has_vector_subgroup_broadcast) {
- CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAT ";
+ CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAST ";
}
- backend_ctx->program_CL_gemv_4096_1_11008 = build_program_from_source(
+ prog = build_program_from_source(
backend_ctx->context, backend_ctx->device, kernel_src_CL_gemv.c_str(), CL_gemv_compile_opts);
- CL_CHECK((backend_ctx->CL_mul_mat_vec_q4_0_f32_1d_4x_flat_4096_1_11008 = clCreateKernel(backend_ctx->program_CL_gemv_4096_1_11008, "kernel_gemv_noshuffle", &err), err));
+ CL_CHECK((backend_ctx->kernel_gemv_noshuffle_q4_0_f32_4096_1_11008 = clCreateKernel(prog, "kernel_gemv_noshuffle_q4_0_f32", &err), err));
+ CL_CHECK(clReleaseProgram(prog));
GGML_LOG_CONT(".");
// Gemv 5504, 44032
" -DSIMDGROUP_WIDTH=" +
std::to_string(backend_ctx->adreno_wave_size);
if (backend_ctx->has_vector_subgroup_broadcast) {
- CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAT ";
+ CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAST ";
}
- backend_ctx->program_CL_gemv_11008_1_4096 = build_program_from_source(
+ prog = build_program_from_source(
backend_ctx->context, backend_ctx->device, kernel_src_CL_gemv.c_str(), CL_gemv_compile_opts);
- CL_CHECK((backend_ctx->CL_mul_mat_vec_q4_0_f32_1d_4x_flat_11008_1_4096 = clCreateKernel(backend_ctx->program_CL_gemv_11008_1_4096, "kernel_gemv_noshuffle", &err), err));
+ CL_CHECK((backend_ctx->kernel_gemv_noshuffle_q4_0_f32_11008_1_4096 = clCreateKernel(prog, "kernel_gemv_noshuffle_q4_0_f32", &err), err));
+ CL_CHECK(clReleaseProgram(prog));
GGML_LOG_CONT(".");
// Gemv 16000, 128000
std::to_string(backend_ctx->adreno_wave_size);
if (backend_ctx->has_vector_subgroup_broadcast) {
- CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAT ";
+ CL_gemv_compile_opts += " -DVECTOR_SUB_GROUP_BROADCAST ";
}
- backend_ctx->program_CL_gemv_32000_1_4096 = build_program_from_source(
+ prog = build_program_from_source(
backend_ctx->context, backend_ctx->device, kernel_src_CL_gemv.c_str(), CL_gemv_compile_opts);
- CL_CHECK((backend_ctx->CL_mul_mat_vec_q4_0_f32_1d_4x_flat_32000_1_4096 = clCreateKernel(backend_ctx->program_CL_gemv_32000_1_4096, "kernel_gemv_noshuffle", &err), err));
+ CL_CHECK((backend_ctx->kernel_gemv_noshuffle_q4_0_f32_32000_1_4096 = clCreateKernel(prog, "kernel_gemv_noshuffle_q4_0_f32", &err), err));
+ CL_CHECK(clReleaseProgram(prog));
GGML_LOG_CONT(".");
}
{
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src_CL_gemm {
- #include "mul_mat_Ab_Bi_8x4.cl.h"
+ #include "gemm_noshuffle_q4_0_f32.cl.h"
};
#else
- const std::string kernel_src_CL_gemm = read_file("mul_mat_Ab_Bi_8x4.cl");
+ const std::string kernel_src_CL_gemm = read_file("gemm_noshuffle_q4_0_f32.cl");
#endif
- backend_ctx->program_CL_gemm = build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src_CL_gemm.c_str(), compile_opts);
- CL_CHECK((backend_ctx->CL_mul_mat_Ab_Bi_8x4 = clCreateKernel(backend_ctx->program_CL_gemm, "kernel_mul_mat_Ab_Bi_8x4", &err), err));
+ cl_program prog = build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src_CL_gemm.c_str(), compile_opts);
+ CL_CHECK((backend_ctx->kernel_gemm_noshuffle_q4_0_f32 = clCreateKernel(prog, "kernel_gemm_noshuffle_q4_0_f32", &err), err));
+ CL_CHECK(clReleaseProgram(prog));
GGML_LOG_CONT(".");
}
// mul_mm_q8_0_f32_8x4
{
#ifdef GGML_OPENCL_EMBED_KERNELS
- const std::string kernel_src_q8_8x4_gemm {
- #include "mul_mm_q8_0_f32_8x4.cl.h"
+ const std::string kernel_src {
+ #include "gemm_noshuffle_q8_0_f32.cl.h"
};
#else
- const std::string kernel_src_q8_8x4_gemm = read_file("mul_mm_q8_0_f32_8x4.cl");
+ const std::string kernel_src = read_file("gemm_noshuffle_q8_0_f32.cl");
#endif
- backend_ctx->program_CL_gemm = build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src_q8_8x4_gemm.c_str(), compile_opts);
- CL_CHECK((backend_ctx->kernel_mul_mm_q8_0_f32_8x4 = clCreateKernel(backend_ctx->program_CL_gemm, "kernel_mul_mm_q8_0_f32_8x4", &err), err));
+ cl_program prog = build_program_from_source(backend_ctx->context, backend_ctx->device, kernel_src.c_str(), compile_opts);
+ CL_CHECK((backend_ctx->kernel_gemm_noshuffle_q8_0_f32 = clCreateKernel(prog, "kernel_gemm_noshuffle_q8_0_f32", &err), err));
+ CL_CHECK(clReleaseProgram(prog));
GGML_LOG_CONT(".");
}
#ifdef GGML_OPENCL_EMBED_KERNELS
const std::string kernel_src_CL_gemv_general {
- #include "gemv_noshuffle_general_q8_0_f32.cl.h"
+ #include "gemv_noshuffle_q8_0_f32.cl.h"
};
#else
- const std::string kernel_src_CL_gemv_general = read_file("gemv_noshuffle_general_q8_0_f32.cl");
+ const std::string kernel_src_CL_gemv_general = read_file("gemv_noshuffle_q8_0_f32.cl");
#endif
cl_program prog = build_program_from_source(
backend_ctx->context, backend_ctx->device, kernel_src_CL_gemv_general.c_str(), CL_gemv_compile_opts);
- CL_CHECK((backend_ctx->CL_mul_mat_vec_q8_0_f32 = clCreateKernel(prog, "kernel_gemv_noshuffle_q8_0_f32", &err), err));
+ CL_CHECK((backend_ctx->kernel_gemv_noshuffle_q8_0_f32 = clCreateKernel(prog, "kernel_gemv_noshuffle_q8_0_f32", &err), err));
CL_CHECK(clReleaseProgram(prog));
GGML_LOG_CONT(".");
}
// Only do transpose for large, non batched matrix
// TODO: use preallocated images instead of sub-buffer then image
if (use_adreno_kernels(backend_ctx, tensor)) {
- // <----------------------------------------------------------------------------------> //
- // start transpose
- // <----------------------------------------------------------------------------------> //
- int M = tensor->ne[1]; // ne01
- int K = tensor->ne[0]; // ne00
-
- //For matrix-vector multiplication kernel, we assume K is a multiple of 32
- GGML_ASSERT(K % 32 == 0);
- //For transpose kernels, we assume K is a multiple of 4 (satisfied by prior assert), and M is a multiple of 4
- GGML_ASSERT(M % 4 == 0);
-
- // transpose is out of place, so we need to allocate transposed buffers
- // <----------------------------------------------------------------------------------> //
- // use sub_buffer of max buffer size instead
-
- size_t q_size_bytes = K * M / 8 * sizeof(float);
- backend_ctx->prealloc_quant_trans.allocate(context, q_size_bytes);
-
- cl_buffer_region region;
- region.origin = 0;
- region.size = q_size_bytes;
- cl_mem qT_d = clCreateSubBuffer(
- backend_ctx->prealloc_quant_trans.buffer,
- 0,
- CL_BUFFER_CREATE_TYPE_REGION,
- ®ion,
- &err);
- CL_CHECK(err);
-
- bool K_tile_trans = true;
- if ((K / 32) % 4 != 0){
- K_tile_trans =false;
- }
-
- size_t d_size_bytes = M * (K / 32) * 2;
- backend_ctx->prealloc_scales_trans.allocate(context, d_size_bytes);
-
- region.origin = 0;
- region.size = d_size_bytes;
- cl_mem dT_d = clCreateSubBuffer(
- backend_ctx->prealloc_scales_trans.buffer,
- 0,
- CL_BUFFER_CREATE_TYPE_REGION,
- ®ion,
- &err);
- CL_CHECK(err);
-
- // <----------------------------------------------------------------------------------> //
-
-
- // create images from the buffers
- // <----------------------------------------------------------------------------------> //
- cl_mem q_d_image1D;
- cl_mem d_d_image1D;
- cl_mem qT_d_image1D;
- cl_mem dT_d_image1D;
-
- cl_image_format img_fmt_1d = { CL_RGBA, CL_HALF_FLOAT };
- cl_image_desc img_desc_1d;
-
- memset(&img_desc_1d, 0, sizeof(img_desc_1d));
- img_desc_1d.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
- img_desc_1d.image_width = M * K / 4 / 4;
- img_desc_1d.buffer = extra->q;
- q_d_image1D = clCreateImage(context, 0, &img_fmt_1d, &img_desc_1d, NULL, &err);
- CL_CHECK(err);
-
- img_fmt_1d = { CL_RGBA, CL_HALF_FLOAT };
- memset(&img_desc_1d, 0, sizeof(img_desc_1d));
- img_desc_1d.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
- img_desc_1d.image_width = M * K / 4 / 4;
- img_desc_1d.buffer = qT_d;
- qT_d_image1D = clCreateImage(context, 0, &img_fmt_1d, &img_desc_1d, NULL, &err);
- CL_CHECK(err);
-
- memset(&img_desc_1d, 0, sizeof(img_desc_1d));
- if (K_tile_trans) {
- img_fmt_1d = { CL_RGBA, CL_HALF_FLOAT };
- img_desc_1d.image_width = M * K / 32 / 4;
- } else {
- img_fmt_1d = { CL_R, CL_HALF_FLOAT };
- img_desc_1d.image_width = M * K / 32;
- }
- img_desc_1d.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
- img_desc_1d.buffer = extra->d;
- d_d_image1D = clCreateImage(context, 0, &img_fmt_1d, &img_desc_1d, NULL, &err);
- CL_CHECK(err);
-
- img_fmt_1d = { CL_RGBA, CL_HALF_FLOAT };
- memset(&img_desc_1d, 0, sizeof(img_desc_1d));
- img_desc_1d.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
- img_desc_1d.image_width = M * K / 32 / 4;
- img_desc_1d.buffer = dT_d;
- dT_d_image1D = clCreateImage(context, 0, &img_fmt_1d, &img_desc_1d, NULL, &err);
- CL_CHECK(err);
- // <----------------------------------------------------------------------------------> //
-
- // set up and call the transpose kernels
- // <----------------------------------------------------------------------------------> //
- // weights
- int height_q = M / 4;
- int width_q = K / 4 / 4;
- kernel = backend_ctx->kernel_transpose_16;
-
- CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &q_d_image1D));
- CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &qT_d_image1D));
- CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), &height_q));
- CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), &width_q));
-
- size_t local_size_q[3] = {4, 16, 1};
- size_t global_size_q[3] = {static_cast<size_t>(width_q), static_cast<size_t>(height_q), 1};
- CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 3, NULL, global_size_q, local_size_q, 0, NULL, &evt));
- CL_CHECK(clWaitForEvents(1, &evt));
-
- // scales
- int height_s = M / 4;
- int width_s = K / 32 / 4;
-
- kernel = backend_ctx->kernel_transpose_16;
- if (!K_tile_trans) {
- kernel = backend_ctx->kernel_transpose_16_4x1;
- width_s = K / 32;
- }
- CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_d_image1D));
- CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &dT_d_image1D));
- CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), &height_s));
- CL_CHECK(clSetKernelArg(kernel, 3, sizeof(int), &width_s));
-
- size_t local_size_s[3] = {4, 16, 1};
- size_t global_size_s[3] = {static_cast<size_t>(width_s), static_cast<size_t>(height_s), 1};
- CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 3, NULL, global_size_s, local_size_s, 0, NULL, &evt));
- CL_CHECK(clWaitForEvents(1, &evt));
- // <----------------------------------------------------------------------------------> //
+ int M = tensor->ne[1];
+ int K = tensor->ne[0];
- // copy transposed buffer contents to original buffers
- // <----------------------------------------------------------------------------------> //
- // weights
- CL_CHECK(clEnqueueCopyBuffer(queue, qT_d, extra->q, 0, 0, q_size_bytes, 0, NULL, &evt));
- CL_CHECK(clWaitForEvents(1, &evt));
+ GGML_ASSERT(K % 32 == 0);
- // scales
- CL_CHECK(clEnqueueCopyBuffer(queue, dT_d, extra->d, 0, 0, d_size_bytes, 0, NULL, &evt));
- CL_CHECK(clWaitForEvents(1, &evt));
- // <----------------------------------------------------------------------------------> //
-
- // deallocate transpose buffers
- // <----------------------------------------------------------------------------------> //
- CL_CHECK(clReleaseMemObject(qT_d));
- CL_CHECK(clReleaseMemObject(dT_d));
-
- // deallocate temporary images
- CL_CHECK(clReleaseMemObject(q_d_image1D));
- CL_CHECK(clReleaseMemObject(d_d_image1D));
- CL_CHECK(clReleaseMemObject(qT_d_image1D));
- CL_CHECK(clReleaseMemObject(dT_d_image1D));
- // <----------------------------------------------------------------------------------> //
- // end transpose
- // <----------------------------------------------------------------------------------> //
+ // Transpose q as ushort
+ transpose_2d_as_16b(backend_ctx, extra->q, extra->q, size_q, K/4, M);
+ // Transpose d as ushort
+ transpose_2d_as_16b(backend_ctx, extra->d, extra->d, size_d, K/32, M);
}
#endif // GGML_OPENCL_USE_ADRENO_KERNELS
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
if (use_adreno_kernels(backend_ctx, tensor)) {
- cl_int err;
- cl_kernel kernel;
+ ggml_cl_buffer buf_trans_q;
+ ggml_cl_buffer buf_trans_d;
+ ggml_cl_buffer buf_unpacked;
cl_int M = tensor->ne[1]; // ne01
cl_int K = tensor->ne[0]; // ne00
size_t size_d = (ggml_nelements(tensor)/ggml_blck_size(tensor->type))*sizeof(ggml_fp16_t);
GGML_ASSERT(size_d + size_q == ggml_nbytes(tensor) && "Incorrect tensor size");
- cl_mem buf_trans_q;
- cl_mem buf_trans_d;
-
- CL_CHECK((buf_trans_q = clCreateBuffer(context, CL_MEM_READ_WRITE,
- size_q, NULL, &err), err));
- CL_CHECK((buf_trans_d = clCreateBuffer(context, CL_MEM_READ_WRITE,
- size_d, NULL, &err), err));
-
- kernel = backend_ctx->kernel_transpose_16_buf;
-
- // transpose q back
- cl_int stride_k_q = K/4;
- size_t local_size_q[3] = {64, 1, 1};
- size_t global_size_q[3] = {(size_t)M, (size_t)stride_k_q, 1};
-
- CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra->q));
- CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &buf_trans_q));
- CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_int), &M));
- CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_int), &stride_k_q));
-
- CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 3, NULL,
- global_size_q, local_size_q, 0, NULL, NULL));
-
- // transpose scales back
- cl_int stride_k_d = K/32;
- size_t local_size_d[3] = {64, 1, 1};
- size_t global_size_d[3] = {(size_t)M, (size_t)stride_k_d, 1};
-
- CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra->d));
- CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &buf_trans_d));
- CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_int), &M));
- CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_int), &stride_k_d));
-
- CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 3, NULL,
- global_size_d, local_size_d, 0, NULL, NULL));
+ buf_trans_q.allocate(backend_ctx->context, size_q);
+ buf_trans_d.allocate(backend_ctx->context, size_d);
+ buf_unpacked.allocate(backend_ctx->context, ggml_nbytes(tensor));
- // unpack
- cl_mem data_device = clCreateBuffer(context, CL_MEM_READ_WRITE,
- ggml_nbytes(tensor), NULL, &err);
- CL_CHECK(err);
+ transpose_2d_as_16b(backend_ctx, extra->q, buf_trans_q.buffer, size_q, M, K/4);
+ transpose_2d_as_16b(backend_ctx, extra->d, buf_trans_d.buffer, size_d, M, K/32);
cl_uchar mask_0F = 0x0F;
cl_uchar mask_F0 = 0xF0;
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};
- kernel = backend_ctx->kernel_restore_block_q4_0_noshuffle;
- CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &buf_trans_q));
- CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &buf_trans_d));
- CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &data_device));
+ cl_kernel kernel = backend_ctx->kernel_restore_block_q4_0_noshuffle;
+ CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &buf_trans_q.buffer));
+ CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &buf_trans_d.buffer));
+ CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &buf_unpacked.buffer));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_uchar), &mask_0F));
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_uchar), &mask_F0));
- CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 3, NULL,
- global_work_size, local_work_size, 0, NULL, NULL));
-
- // read back to host
- CL_CHECK(clEnqueueReadBuffer(
- queue, data_device, CL_TRUE, offset,
- size, data, 0, NULL, NULL));
-
- CL_CHECK(clReleaseMemObject(data_device));
- CL_CHECK(clReleaseMemObject(buf_trans_q));
- CL_CHECK(clReleaseMemObject(buf_trans_d));
-
+ CL_CHECK(clEnqueueNDRangeKernel(queue, kernel, 3, NULL, global_work_size, local_work_size, 0, NULL, NULL));
+ CL_CHECK(clEnqueueReadBuffer(queue, buf_unpacked.buffer, CL_TRUE, offset, size, data, 0, NULL, NULL));
return;
}
#endif
CL_CHECK(clReleaseMemObject(D_sub_buffer));
}
+static void ggml_cl_mul_mat_q4_0_f32_adreno(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
+ GGML_ASSERT(src0);
+ GGML_ASSERT(src0->extra);
+ GGML_ASSERT(src1);
+ GGML_ASSERT(src1->extra);
+ GGML_ASSERT(dst);
+ GGML_ASSERT(dst->extra);
+
+ ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context;
+
+ ggml_tensor_extra_cl * extra1 = (ggml_tensor_extra_cl *)src1->extra;
+ ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *)dst->extra;
+ ggml_tensor_extra_cl_q4_0 * extra0_q4_0 = (ggml_tensor_extra_cl_q4_0 *)src0->extra;
+
+ cl_ulong offset1 = extra1->offset + src1->view_offs;
+ cl_ulong offsetd = extrad->offset + dst->view_offs;
+
+ const int ne00 = src0->ne[0];
+ const int ne01 = src0->ne[1];
+ const int ne02 = src0->ne[2];
+
+ const int ne10 = src1->ne[0];
+ const int ne12 = src1->ne[2];
+
+ const int ne0 = dst->ne[0];
+ const int ne1 = dst->ne[1];
+
+ GGML_ASSERT(ne00 % ggml_blck_size(src0->type) == 0);
+
+ cl_context context = backend_ctx->context;
+ cl_kernel kernel;
+
+ cl_int err;
+ cl_image_format img_fmt;
+ cl_image_desc img_desc;
+ cl_buffer_region region;
+
+ int M = ne01;
+ int N = ne1;
+ int K = ne00;
+
+ if (ne1 == 1) {
+ cl_mem q_img = nullptr;
+ cl_mem b_sub_buf = nullptr;
+ cl_mem b_img = nullptr;
+
+ // image for q
+ img_fmt = { CL_R, CL_UNSIGNED_INT32};
+ memset(&img_desc, 0, sizeof(img_desc));
+ img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+ img_desc.image_width = M * K / 2 / 4;
+ img_desc.buffer = extra0_q4_0->q;
+ CL_CHECK((q_img = clCreateImage(context, CL_MEM_READ_ONLY, &img_fmt, &img_desc, NULL, &err), err));
+
+ // subbuffer for activations
+ region.origin = offset1;
+ region.size = K * N * sizeof(float);
+ CL_CHECK((b_sub_buf = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err));
+
+ // image for activations
+ img_fmt = {CL_RGBA, CL_FLOAT};
+ memset(&img_desc, 0, sizeof(img_desc));
+ img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+ img_desc.image_width = K * N / 4;
+ img_desc.buffer = b_sub_buf;
+ CL_CHECK((b_img = clCreateImage(context, CL_MEM_READ_ONLY, &img_fmt, &img_desc, NULL, &err), err));
+
+ kernel = backend_ctx->kernel_gemv_noshuffle_q4_0_f32;
+ if (M == 4096 && K == 4096) {
+ kernel = backend_ctx->kernel_gemv_noshuffle_q4_0_f32_4096_1_4096;
+ } else if (M == 4096 && K == 11008) {
+ kernel = backend_ctx->kernel_gemv_noshuffle_q4_0_f32_4096_1_11008;
+ } else if (M == 11008 && K == 4096) {
+ kernel = backend_ctx->kernel_gemv_noshuffle_q4_0_f32_11008_1_4096;
+ } else if (M == 32000 && K == 4096) {
+ kernel = backend_ctx->kernel_gemv_noshuffle_q4_0_f32_32000_1_4096;
+ }
+
+ int r2 = 1;
+ int r3 = 1;
+
+ CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &q_img));
+ CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q4_0->d));
+ CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &b_img));
+ 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));
+
+ size_t local_work_size[3] = {64, 4, 1};
+ size_t global_work_size[3] = {(size_t)CEIL_DIV(ne01/2, 64)*64, 4, 1};
+
+ backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
+
+ CL_CHECK(clReleaseMemObject(q_img));
+ CL_CHECK(clReleaseMemObject(b_sub_buf));
+ CL_CHECK(clReleaseMemObject(b_img));
+ } else {
+ cl_mem b_sub_buf = nullptr;
+ cl_mem b_sub_buf_trans = nullptr;
+ cl_mem b_img = nullptr;
+ cl_mem b_img_trans = nullptr;
+ cl_mem d_sub_buf = nullptr;
+
+ // subbuffer for activations
+ region.origin = offset1;
+ region.size = K * N * sizeof(float);
+ CL_CHECK((b_sub_buf = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err));
+
+ // image for activations
+ img_fmt = {CL_RGBA, CL_FLOAT};
+ memset(&img_desc, 0, sizeof(img_desc));
+ img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+ img_desc.image_width = K * N / 4;
+ img_desc.buffer = b_sub_buf;
+ CL_CHECK((b_img = clCreateImage(context, CL_MEM_READ_ONLY, &img_fmt, &img_desc, NULL, &err), err));
+
+ // pad N to multiple of 8
+ int extra_elements = N % 8;
+ int padding = 0;
+ if (extra_elements > 0){
+ padding = 8 - extra_elements;
+ }
+
+ // subbuffer for transposed activations
+ region.origin = 0;
+ region.size = K * (N + padding) * sizeof(float)/2;
+ backend_ctx->prealloc_act_trans.allocate(context, region.size);
+ CL_CHECK((b_sub_buf_trans = clCreateSubBuffer(backend_ctx->prealloc_act_trans.buffer, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err));
+
+ // image for transposed activations
+ img_fmt = {CL_RGBA, CL_HALF_FLOAT};
+ memset(&img_desc, 0, sizeof(img_desc));
+ img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+ img_desc.image_width = K * (N + padding) / 4;
+ img_desc.buffer = b_sub_buf_trans;
+ CL_CHECK((b_img_trans = clCreateImage(context, 0, &img_fmt, &img_desc, NULL, &err), err));
+
+ // subbuffer for output
+ region.origin = extrad->offset; // Specify the starting offset (in bytes)
+ region.size = M * N * sizeof(float); // Specify the size of the sub-buffer
+ CL_CHECK((d_sub_buf = clCreateSubBuffer(extrad->data_device, CL_MEM_WRITE_ONLY, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err));
+
+ // transpose activations
+ int height_B = N/4;
+ if (height_B == 0) {
+ height_B = 1;
+ }
+ int width_B = K/4;
+ int padded_height_B = (N + padding)/4;
+
+ 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 local_work_size_t[2] = { 1, 16 };
+ size_t global_work_size_t[2] = { (size_t)width_B, (size_t)padded_height_B };
+ if (ne0 == 4096 && ne1 == 128 && ne10 == 4096) {
+ local_work_size_t[0]=4;
+ local_work_size_t[1]=8;
+ } else if (ne0 == 11008 && ne1 == 128 && ne10 == 4096) {
+ local_work_size_t[0]=2;
+ local_work_size_t[1]=8;
+ } else if(ne0 == 4096 && ne1 == 128 && ne10 == 11008) {
+ local_work_size_t[0]=1;
+ local_work_size_t[1]=8;
+ } else if(ne0 == 32000 && ne1 == 128 && ne10 == 4096) {
+ local_work_size_t[0]=2;
+ local_work_size_t[1]=8;
+ }
+ backend_ctx->enqueue_ndrange_kernel(kernel, 2, global_work_size_t, local_work_size_t, dst);
+
+ // gemm
+ kernel = backend_ctx->kernel_gemm_noshuffle_q4_0_f32;
+ int padded_N = N + padding;
+
+ CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q4_0->q));
+ CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q4_0->d));
+ CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &b_img_trans));
+ CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &d_sub_buf));
+ CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_int), &ne01));
+ CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_int), &padded_N));
+ CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_int), &ne00));
+ CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_int), &ne1));
+
+ size_t global_work_size[3] = {(size_t)CEIL_DIV(ne1, 8), (size_t)CEIL_DIV(ne01, 4), 1};
+ size_t local_work_size[3] = {1, 128, 1};
+ if (ne0 == 4096 && ne1 == 128 && ne10 == 4096) {
+ local_work_size[0] = 1;
+ local_work_size[1] = 128;
+ } else if (ne0 == 11008 && ne1 == 128 && ne10 == 4096) {
+ local_work_size[0] = 2;
+ local_work_size[1] = 64;
+ } else if (ne0 == 4096 && ne1 == 128 && ne10 == 11008) {
+ local_work_size[0] = 2;
+ local_work_size[1] = 64;
+ } else if (ne0 == 32000 && ne1 == 128 && ne10 == 4096) {
+ local_work_size[0] = 2;
+ local_work_size[1] = 64;
+ }
+
+ backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
+
+ CL_CHECK(clReleaseMemObject(b_sub_buf));
+ CL_CHECK(clReleaseMemObject(b_sub_buf_trans));
+ CL_CHECK(clReleaseMemObject(b_img));
+ CL_CHECK(clReleaseMemObject(b_img_trans));
+ CL_CHECK(clReleaseMemObject(d_sub_buf));
+ }
+#else
+ GGML_UNUSED(backend);
+ GGML_UNUSED(src0);
+ GGML_UNUSED(src1);
+ GGML_UNUSED(dst);
+#endif
+}
+
static void ggml_cl_mul_mat_q4_1_f32_adreno(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
GGML_ASSERT(src0);
img_desc.buffer = b_sub_buf;
CL_CHECK((b_img = clCreateImage(context, CL_MEM_READ_ONLY, &img_fmt, &img_desc, NULL, &err), err));
- kernel = backend_ctx->CL_mul_mat_vec_q8_0_f32;
+ kernel = backend_ctx->kernel_gemv_noshuffle_q8_0_f32;
int r2 = 1;
int r3 = 1;
backend_ctx->enqueue_ndrange_kernel(kernel, 2, global_work_size_t, local_work_size_t, dst);
// gemm
- kernel = backend_ctx->kernel_mul_mm_q8_0_f32_8x4;
+ kernel = backend_ctx->kernel_gemm_noshuffle_q8_0_f32;
int padded_N = N + padding;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q8_0->q));
GGML_ASSERT(dst);
GGML_ASSERT(dst->extra);
- const enum ggml_type src0t = src0 ? src0->type : GGML_TYPE_COUNT;
- const enum ggml_type src1t = src1 ? src1->type : GGML_TYPE_COUNT;
+ const enum ggml_type src0t = src0->type;
+ const enum ggml_type src1t = src1->type;
ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context;
ggml_tensor_extra_cl_q6_K * extra0_q6_K = (ggml_tensor_extra_cl_q6_K *)src0->extra;
#endif
- const int ne00 = src0 ? src0->ne[0] : 0;
- const int ne01 = src0 ? src0->ne[1] : 0;
- const int ne02 = src0 ? src0->ne[2] : 0;
- const int ne03 = src0 ? src0->ne[3] : 0;
-
- const cl_ulong nb00 = src0 ? src0->nb[0] : 0;
- const cl_ulong nb01 = src0 ? src0->nb[1] : 0;
- const cl_ulong nb02 = src0 ? src0->nb[2] : 0;
- const cl_ulong nb03 = src0 ? src0->nb[3] : 0;
-
- const int ne10 = src1 ? src1->ne[0] : 0;
- const int ne11 = src1 ? src1->ne[1] : 0;
- const int ne12 = src1 ? src1->ne[2] : 0;
- const int ne13 = src1 ? src1->ne[3] : 0;
-
- const cl_ulong nb10 = src1 ? src1->nb[0] : 0;
- const cl_ulong nb11 = src1 ? src1->nb[1] : 0;
- const cl_ulong nb12 = src1 ? src1->nb[2] : 0;
- const cl_ulong nb13 = src1 ? src1->nb[3] : 0;
-
- const int ne0 = dst ? dst->ne[0] : 0;
- const int ne1 = dst ? dst->ne[1] : 0;
+ GGML_TENSOR_LOCALS(int, ne0, src0, ne);
+ GGML_TENSOR_LOCALS(cl_ulong, nb0, src0, nb);
+ GGML_TENSOR_LOCALS(int, ne1, src1, ne);
+ GGML_TENSOR_LOCALS(cl_ulong, nb1, src1, nb);
+ GGML_TENSOR_LOCALS(int, ne, dst, ne);
+ GGML_TENSOR_LOCALS(cl_ulong, nb, dst, nb);
int r2 = ne12/ne02;
int r3 = ne13/ne03;
cl_kernel kernel;
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
- cl_context context = backend_ctx->context;
-
if(src0t == GGML_TYPE_F16 && src1t == GGML_TYPE_F32){
if (ne01 >= 64 && ne1 >= 32 && ne00 >= 16 && (ne12 % ne02) == 0 &&
// dst is wrapped with image1d_buffer, the size limit applies, also src0
}
if (ne01 && ne1 && use_adreno_kernels(backend_ctx, src0)) {
+ // NOTE: Kernels using image1d_buffer_t (e.g., src0_q) would normally require
+ // a limit check, but q4_0 / q4_1 tensors are very unlikely to exceed that
+ // limit, so the check is omitted.
- // init CL objects
- // <--------------------------------------------> //
- cl_int status;
- cl_image_format img_fmt_1d;
- cl_image_desc img_desc_1d;
- cl_buffer_region region;
- cl_mem A_image1d = nullptr;
- cl_mem B_image1d = nullptr;
- cl_mem B_sub_buffer = nullptr;
- cl_mem C_d = nullptr;
- // for B transpose
- cl_mem B_d = nullptr;
- cl_mem B_d_input_image = nullptr;
- // <--------------------------------------------> //
-
- // define matrix dimensions
- // <--------------------------------------------> //
- int M = ne01;
- int N = ne1;
- int K = ne00;
- int padding;
- // <--------------------------------------------> //
-
- // NOTE: Kernels using image1d_buffer_t (e.g., src0_q) would normally require
- // a limit check, but q4_0 / q4_1 tensors are very unlikely to exceed that
- // limit, so the check is omitted.
+ // 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);
+ return;
+ }
- // q4_1 x fp32
- if (src0t == GGML_TYPE_Q4_1 && src1t == GGML_TYPE_F32) {
+ // q4_1 x fp32
+ if (src0t == GGML_TYPE_Q4_1 && src1t == GGML_TYPE_F32) {
ggml_cl_mul_mat_q4_1_f32_adreno(backend, src0, src1, dst);
return;
- }
-
- // iq4_nl x fp32
- if (src0t == GGML_TYPE_IQ4_NL && src1t == GGML_TYPE_F32) {
- ggml_cl_mul_mat_iq4_nl_f32_adreno(backend, src0, src1, dst);
- return;
- }
+ }
- // q8_0 x fp32
- if (src0t == GGML_TYPE_Q8_0 && src1t == GGML_TYPE_F32 &&
- enable_adreno_trans_weight(backend_ctx, src0)) {
- ggml_cl_mul_mat_q8_0_f32_adreno(backend, src0, src1, dst);
+ // iq4_nl x fp32
+ if (src0t == GGML_TYPE_IQ4_NL && src1t == GGML_TYPE_F32) {
+ ggml_cl_mul_mat_iq4_nl_f32_adreno(backend, src0, src1, dst);
return;
- }
+ }
+
+ // q8_0 x fp32
+ if (src0t == GGML_TYPE_Q8_0 && src1t == GGML_TYPE_F32 &&
+ enable_adreno_trans_weight(backend_ctx, src0)) {
+ ggml_cl_mul_mat_q8_0_f32_adreno(backend, src0, src1, dst);
+ return;
+ }
- // q4_k x fp32
- if (src0t == GGML_TYPE_Q4_K && src1t == GGML_TYPE_F32) {
+ // q4_k x fp32
+ if (src0t == GGML_TYPE_Q4_K && src1t == GGML_TYPE_F32) {
ggml_cl_mul_mat_q4_k_f32_adreno(backend, src0, src1, dst);
return;
- }
-
- // q6_K x fp32
- if (src0t == GGML_TYPE_Q6_K && src1t == GGML_TYPE_F32) {
- ggml_cl_mul_mat_q6_K_f32_adreno(backend, src0, src1, dst);
- return;
- }
-
- // q5_K x fp32
- if (src0t == GGML_TYPE_Q5_K && src1t == GGML_TYPE_F32) {
- ggml_cl_mul_mat_q5_K_f32_adreno(backend, src0, src1, dst);
- return;
- }
-
- // q4_0 x fp32
- if(src0t == GGML_TYPE_Q4_0 && src1t == GGML_TYPE_F32) {
- // TODO: remove duplicate definitions of image description + format -- move to top
-
- // create an image for A
- // <--------------------------------------------> //
- if (N == 1) {
- img_fmt_1d = { CL_R, CL_UNSIGNED_INT32};
- } else {
- img_fmt_1d = { CL_R, CL_FLOAT};
- }
- memset(&img_desc_1d, 0, sizeof(img_desc_1d));
- img_desc_1d.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
- img_desc_1d.image_width = M * K / 2 / 4; // Divide by 4 for char -> float
- img_desc_1d.buffer = extra0_q4_0->q;
- A_image1d = clCreateImage(
- context,
- CL_MEM_READ_ONLY,
- &img_fmt_1d,
- &img_desc_1d,
- NULL,
- &status);
- CL_CHECK(status);
- // <--------------------------------------------> //
-
-
- // create a sub_buffer for B
- // <--------------------------------------------> //
- region.origin = (extra1->offset);
- region.size = K * N * sizeof(float);
- B_sub_buffer = clCreateSubBuffer(
- extra1->data_device,
- 0,
- CL_BUFFER_CREATE_TYPE_REGION,
- ®ion,
- &status);
- CL_CHECK(status);
- // <--------------------------------------------> //
-
- // transpose activation for Skyler's gemm
- if (N != 1) {
- //how many extra elements beyond multiple of 8
- int extra_elements = N % 8;
-
- //how much padding to add
- padding = 0;
- if (extra_elements > 0){
- padding = 8 - extra_elements;
- }
-
- // Specify the starting offset (in bytes)
- region.origin = 0;
- // Specify the size of the sub-buffer (divide by 2 for FP16)
- region.size = K * (N + padding) * sizeof(float)/2;
- backend_ctx->prealloc_act_trans.allocate(context, region.size);
-
- B_d = clCreateSubBuffer(
- backend_ctx->prealloc_act_trans.buffer,
- 0,
- CL_BUFFER_CREATE_TYPE_REGION,
- ®ion,
- &status);
- CL_CHECK(status);
-
- cl_image_format image_format_B_d_input = { CL_RGBA, CL_FLOAT };
- cl_image_desc image_desc_B_d_input = {
- CL_MEM_OBJECT_IMAGE1D_BUFFER,
- static_cast<size_t>(K * N / 4),
- 0, 0, 0, 0, 0, 0, 0, { B_sub_buffer }
- };
- B_d_input_image = clCreateImage(
- context,
- 0,
- &image_format_B_d_input,
- &image_desc_B_d_input,
- NULL,
- &status);
- CL_CHECK(status);
-
- cl_image_format image_format_B_d_output = { CL_RGBA, CL_HALF_FLOAT }; //(CL_HALF_FLOAT for FP16)
- cl_image_desc image_desc_B_d_output = {
- CL_MEM_OBJECT_IMAGE1D_BUFFER,
- static_cast<size_t>(K * (N + padding)/4),
- 0, 0, 0, 0, 0, 0, 0, { B_d }
- };
- B_image1d = clCreateImage(
- context,
- 0,
- &image_format_B_d_output,
- &image_desc_B_d_output,
- NULL,
- &status);
- CL_CHECK(status);
-
- int height_B = N/4;
- if (height_B == 0) {
- height_B = 1;
- }
- int width_B = K/4;
- int padded_height_B = (N + padding)/4;
-
- kernel = backend_ctx->kernel_transpose_32_16;
- CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &B_d_input_image));
- CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &B_image1d));
- 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 local_size_t[2] = { 1, 16 };
- //WGS tuning
- if (ne0 == 4096 && ne1 == 128 && ne10 == 4096) {
- local_size_t[0]=4;
- local_size_t[1]=8;
- } else if (ne0 == 11008 && ne1 == 128 && ne10 == 4096) {
- local_size_t[0]=2;
- local_size_t[1]=8;
- } else if(ne0 == 4096 && ne1 == 128 && ne10 == 11008) {
- local_size_t[0]=1;
- local_size_t[1]=8;
- } else if(ne0 == 32000 && ne1 == 128 && ne10 == 4096) {
- local_size_t[0]=2;
- local_size_t[1]=8;
- }
-
- size_t global_size_t[2] = {
- static_cast<size_t>(width_B),
- static_cast<size_t>(padded_height_B)
- };
-
- backend_ctx->enqueue_ndrange_kernel(kernel, 2, global_size_t, local_size_t, dst);
- } else {
- // no need to transpose B in other cases
- // create an image for B from sub_buffer
- // <--------------------------------------------> //
- img_fmt_1d = {CL_RGBA, CL_FLOAT};
-
- memset(&img_desc_1d, 0, sizeof(img_desc_1d));
- img_desc_1d.image_width = K * N / 4;
- img_desc_1d.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
- img_desc_1d.buffer = B_sub_buffer;
- B_image1d = clCreateImage(
- context,
- CL_MEM_READ_ONLY,
- &img_fmt_1d,
- &img_desc_1d,
- NULL,
- &status);
- CL_CHECK(status);
- // <--------------------------------------------> //
- }
-
- // choose gemm or gemv kernel
- // <--------------------------------------------> //
- if (N == 1) {
- kernel = backend_ctx->CL_mul_mat_vec_q4_0_f32_1d_4x_flat_general;
- if (M == 4096 && K == 4096) {
- kernel = backend_ctx->CL_mul_mat_vec_q4_0_f32_1d_4x_flat_4096_1_4096;
- } else if (M == 4096 && K == 11008) {
- kernel = backend_ctx->CL_mul_mat_vec_q4_0_f32_1d_4x_flat_4096_1_11008;
- } else if (M == 11008 && K == 4096) {
- kernel = backend_ctx->CL_mul_mat_vec_q4_0_f32_1d_4x_flat_11008_1_4096;
- } else if (M == 32000 && K == 4096) {
- kernel = backend_ctx->CL_mul_mat_vec_q4_0_f32_1d_4x_flat_32000_1_4096;
- }
- } else {
- kernel = backend_ctx->CL_mul_mat_Ab_Bi_8x4;
- }
- // <--------------------------------------------> //
-
- // set kernel args
- // <--------------------------------------------> //
- cl_uint k_arg = 0;
-
- if (N == 1) {
- CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &A_image1d));
- CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &extra0_q4_0->d));
- CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &B_image1d));
- CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_ulong), &extra1->offset));
- CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_mem), &extrad->data_device));
- CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(cl_ulong), &extrad->offset));
- CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(int), &ne00));
- CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(int), &ne01));
- CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(int), &ne02));
- CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(int), &ne10));
- CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(int), &ne12));
- CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(int), &ne0));
- CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(int), &ne1));
- CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(int), &r2));
- CL_CHECK(clSetKernelArg(kernel, k_arg++, sizeof(int), &r3));
- } else {
- region.origin = extrad->offset; // Specify the starting offset (in bytes)
- region.size = M * N * sizeof(float); // Specify the size of the sub-buffer
- C_d = clCreateSubBuffer(extrad->data_device, CL_MEM_WRITE_ONLY, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &status);
- CL_CHECK(status);
-
- int padded_N = ne1 + padding;
-
- CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q4_0->q)); //A_q_dextra0_q4_0->q
- CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q4_0->d)); //A_s_d
- CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &B_image1d)); //B_d
- CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &C_d)); //C_d
- CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), &ne01)); //M
- CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), &padded_N)); //N with padding
- CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), &ne00)); //K
- CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), &ne1)); //N without padding
- }
- // <--------------------------------------------> //
-
- // choose workgroup size
- // <--------------------------------------------> //
- size_t global_work_size[3] = {
- 64, static_cast<size_t>((M+63)/64), static_cast<size_t>((N+31)/32)};
- size_t local_work_size[3] = {64, 2, 4};
-
- global_work_size[0] = (size_t)(ceil((float)ne1/8));
- global_work_size[1] = (size_t)(ne01/4);
- global_work_size[2] = (size_t)(1);
-
- local_work_size[0] = (size_t)(1); //4x32 for FP32
- local_work_size[1] = (size_t)(128);
- local_work_size[2] = (size_t)(1);
-
- //WGS tuning
- if (ne0 == 4096 && ne1 == 128 && ne10 == 4096) {
- local_work_size[0] = 1;
- local_work_size[1] = 128;
- } else if (ne0 == 11008 && ne1 == 128 && ne10 == 4096) {
- local_work_size[0] = 2;
- local_work_size[1] = 64;
- } else if (ne0 == 4096 && ne1 == 128 && ne10 == 11008) {
- local_work_size[0] = 2;
- local_work_size[1] = 64;
- } else if (ne0 == 32000 && ne1 == 128 && ne10 == 4096) {
- local_work_size[0] = 2;
- local_work_size[1] = 64;
}
- if (N == 1) {
- size_t wavesize = backend_ctx->adreno_wave_size;
- local_work_size[0] = wavesize; // localsize
- local_work_size[1] = 4; // reduce factor
- local_work_size[2] = 1;
-
- global_work_size[0] = (((M / 2) + wavesize - 1) / wavesize) * wavesize;
- global_work_size[1] = 4; // reduce factor
- global_work_size[2] = 1;
+ // q6_K x fp32
+ if (src0t == GGML_TYPE_Q6_K && src1t == GGML_TYPE_F32) {
+ ggml_cl_mul_mat_q6_K_f32_adreno(backend, src0, src1, dst);
+ return;
}
- // <--------------------------------------------> //
-
- // enqueue kernel with profiling
- // <--------------------------------------------> //
- backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
- // <--------------------------------------------> //
-
- // deallocate sub buffers and images
- // <--------------------------------------------> //
- CL_CHECK(clReleaseMemObject(A_image1d));
- CL_CHECK(clReleaseMemObject(B_sub_buffer));
- CL_CHECK(clReleaseMemObject(B_image1d));
- if (N != 1) {
- CL_CHECK(clReleaseMemObject(B_d));
- CL_CHECK(clReleaseMemObject(B_d_input_image));
- CL_CHECK(clReleaseMemObject(C_d));
+ // q5_K x fp32
+ if (src0t == GGML_TYPE_Q5_K && src1t == GGML_TYPE_F32) {
+ ggml_cl_mul_mat_q5_K_f32_adreno(backend, src0, src1, dst);
+ return;
}
- // <--------------------------------------------> //
-
- return;
- }
} // if (ne01 && ne1)
#endif // GGML_OPENCL_USE_ADRENO_KERNELS
--- /dev/null
+// src0_q, src0_d, src1 are transposed as a preprocessing step
+// 4-bit weights are transposed in groups of 4 (unsigned short int)
+// consider weights originally "next to each other", now "on top of each other"
+// each fiber computes a 8x4 tile of output elements
+// using unshuffled weights
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
+
+#ifdef cl_qcom_reqd_sub_group_size
+#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
+#define ADRENO_GPU 1
+#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
+#endif
+
+#ifdef ADRENO_GPU
+REQD_SUBGROUP_SIZE_128
+#endif
+
+kernel void kernel_gemm_noshuffle_q4_0_f32(
+ global const ushort * src0_q, // quantized A
+ global const half * src0_d, // A scales
+ __read_only image1d_buffer_t src1, // B (1d image)
+ global float * dst, // C
+ int m, // M
+ int n, // N with padding
+ int k, // K
+ int n_no_padding // N without padding
+) {
+
+ int m_4 = m >> 2;
+ int n_4 = n >> 2;
+
+ int gy = get_global_id(0);
+ int gx = get_global_id(1);
+ int gx_2 = gx << 2;
+
+ half8 c0 = 0, c1 = 0, c2 = 0, c3 = 0; // 8x4 output elements
+ half8 B; // registers for activations
+ half4 dequantized_weights; // registers for dequantized weights
+ __global const ushort* weight_ptr = src0_q + gx_2; // pointer for weights
+ __global const half* scale_ptr = src0_d + gx_2; // pointer for scales
+
+ for(int i=0; i<k; i+=4){ //loop through K dimension
+
+ B.s0123 = read_imageh(src1, gy*2 + (i)*(n_4));
+ B.s4567 = read_imageh(src1, gy*2 + (i)*(n_4)+1);
+
+ // keep (i/4) and (i/32) in parenthesis, rounds down
+ // load 4 consecutive groups of 4 weights
+ ushort4 bits4 = vload4(0, weight_ptr + (i/4)*(m)); // (i/4) because weights grouped in 4s
+
+ // load 4 consecutive scales
+ half4 scale = vload4(0, scale_ptr + (i/32)*(m));// (i/32) because 1 scale per 32 elements
+
+ // j=0
+ dequantized_weights.s0 = ((bits4.s0 & (0x000F)) - 8) * scale.s0; // dequantize a row of the 16 weights
+ dequantized_weights.s1 = ((bits4.s1 & (0x000F)) - 8) * scale.s1;
+ dequantized_weights.s2 = ((bits4.s2 & (0x000F)) - 8) * scale.s2;
+ dequantized_weights.s3 = ((bits4.s3 & (0x000F)) - 8) * scale.s3;
+ c0 += B * dequantized_weights.s0; // vector-scalar multiplication to accumulate
+ c1 += B * dequantized_weights.s1;
+ c2 += B * dequantized_weights.s2;
+ c3 += B * dequantized_weights.s3;
+
+ // j=1
+ B.s0123 = read_imageh(src1, gy*2 + (i+1)*(n_4));
+ B.s4567 = read_imageh(src1, gy*2 + (i+1)*(n_4)+1);
+ dequantized_weights.s0 = (((bits4.s0 & (0x00F0)) >> 4) - 8) * scale.s0; // dequantize a row of the 16 weights
+ dequantized_weights.s1 = (((bits4.s1 & (0x00F0)) >> 4) - 8) * scale.s1;
+ dequantized_weights.s2 = (((bits4.s2 & (0x00F0)) >> 4) - 8) * scale.s2;
+ dequantized_weights.s3 = (((bits4.s3 & (0x00F0)) >> 4) - 8) * scale.s3;
+ c0 += B * dequantized_weights.s0; //vector-scalar multiplication to accumulate
+ c1 += B * dequantized_weights.s1;
+ c2 += B * dequantized_weights.s2;
+ c3 += B * dequantized_weights.s3;
+
+ // j=2
+ B.s0123 = read_imageh(src1, gy*2 + (i+2)*(n_4));
+ B.s4567 = read_imageh(src1, gy*2 + (i+2)*(n_4)+1);
+ dequantized_weights.s0 = (((bits4.s0 & (0x0F00)) >> 8) - 8) * scale.s0; // dequantize a row of the 16 weights
+ dequantized_weights.s1 = (((bits4.s1 & (0x0F00)) >> 8) - 8) * scale.s1;
+ dequantized_weights.s2 = (((bits4.s2 & (0x0F00)) >> 8) - 8) * scale.s2;
+ dequantized_weights.s3 = (((bits4.s3 & (0x0F00)) >> 8) - 8) * scale.s3;
+ c0 += B * dequantized_weights.s0; // vector-scalar multiplication to accumulate
+ c1 += B * dequantized_weights.s1;
+ c2 += B * dequantized_weights.s2;
+ c3 += B * dequantized_weights.s3;
+
+ // j=3
+ B.s0123 = read_imageh(src1, gy*2 + (i+3)*(n_4));
+ B.s4567 = read_imageh(src1, gy*2 + (i+3)*(n_4)+1);
+ dequantized_weights.s0 = (((bits4.s0 & (0xF000)) >> 12) - 8) * scale.s0; // dequantize a row of the 16 weights
+ dequantized_weights.s1 = (((bits4.s1 & (0xF000)) >> 12) - 8) * scale.s1;
+ dequantized_weights.s2 = (((bits4.s2 & (0xF000)) >> 12) - 8) * scale.s2;
+ dequantized_weights.s3 = (((bits4.s3 & (0xF000)) >> 12) - 8) * scale.s3;
+ c0 += B * dequantized_weights.s0; // vector-scalar multiplication to accumulate
+ c1 += B * dequantized_weights.s1;
+ c2 += B * dequantized_weights.s2;
+ c3 += B * dequantized_weights.s3;
+ }
+
+ int idx = (gy<<3)*m + (gx<<2); // vectorized store 16 elements
+
+ // conditional check if store is to a valid location. Required when N is not a multiple of 8
+ // if statements allow registers to be reused for each store
+ // provides a performance boost due to reduced register footprint, which increases number of concurrent waves
+ 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);
+ }
+}
--- /dev/null
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
+
+#ifdef cl_qcom_reqd_sub_group_size
+#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
+#define ADRENO_GPU 1
+#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
+#endif
+
+#ifdef ADRENO_GPU
+REQD_SUBGROUP_SIZE_128
+#endif
+
+kernel void kernel_gemm_noshuffle_q8_0_f32(
+ global const uint * src0_q,
+ global const half * src0_d,
+ __read_only image1d_buffer_t src1,
+ global float * dst,
+ int k,
+ int m,
+ int n,
+ int n_no_padding,
+ ulong offsetd
+) {
+
+ int m_4 = m >> 2;
+ int n_4 = n >> 2;
+
+ int gy = get_global_id(0);
+ int gx = get_global_id(1);
+ int gx_2 = gx << 2;
+ dst = (global float *)((global char*)dst + offsetd);
+
+
+ half8 c0 = 0, c1 = 0, c2 = 0, c3 = 0;
+ half8 B;
+ half4 deq;
+
+ __global const uint* wptr = src0_q + gx_2;
+ __global const half* sptr = src0_d + gx_2;
+
+ for (int i = 0; i < k; i += 4) {
+ uint4 pack4 = vload4(0, wptr + (i / 4) * m);
+ half4 scale = vload4(0, sptr + (i / 32) * m);
+
+ char4 p0 = as_char4(pack4.s0);
+ char4 p1 = as_char4(pack4.s1);
+ char4 p2 = as_char4(pack4.s2);
+ char4 p3 = as_char4(pack4.s3);
+
+ // ------------------- j = 0 (k = i+0) -------------------
+ B.s0123 = read_imageh(src1, gy * 2 + (i + 0) * n_4);
+ B.s4567 = read_imageh(src1, gy * 2 + (i + 0) * n_4 + 1);
+
+ half4 wj0 = convert_half4((char4)(p0.s0, p1.s0, p2.s0, p3.s0)) * scale;
+
+ c0 += B * wj0.s0;
+ c1 += B * wj0.s1;
+ c2 += B * wj0.s2;
+ c3 += B * wj0.s3;
+
+ // ------------------- j = 1 (k = i+1) -------------------
+ B.s0123 = read_imageh(src1, gy * 2 + (i + 1) * n_4);
+ B.s4567 = read_imageh(src1, gy * 2 + (i + 1) * n_4 + 1);
+
+ half4 wj1 = convert_half4((char4)(p0.s1, p1.s1, p2.s1, p3.s1)) * scale;
+
+ c0 += B * wj1.s0;
+ c1 += B * wj1.s1;
+ c2 += B * wj1.s2;
+ c3 += B * wj1.s3;
+
+ // ------------------- j = 2 (k = i+2) -------------------
+ B.s0123 = read_imageh(src1, gy * 2 + (i + 2) * n_4);
+ B.s4567 = read_imageh(src1, gy * 2 + (i + 2) * n_4 + 1);
+
+ half4 wj2 = convert_half4((char4)(p0.s2, p1.s2, p2.s2, p3.s2)) * scale;
+
+ c0 += B * wj2.s0;
+ c1 += B * wj2.s1;
+ c2 += B * wj2.s2;
+ c3 += B * wj2.s3;
+
+ // ------------------- j = 3 (k = i+3) -------------------
+ B.s0123 = read_imageh(src1, gy * 2 + (i + 3) * n_4);
+ B.s4567 = read_imageh(src1, gy * 2 + (i + 3) * n_4 + 1);
+
+ half4 wj3 = convert_half4((char4)(p0.s3, p1.s3, p2.s3, p3.s3)) * scale;
+
+ c0 += B * wj3.s0;
+ c1 += B * wj3.s1;
+ c2 += B * wj3.s2;
+ c3 += B * wj3.s3;
+ }
+
+ int idx = (gy << 3) * m + (gx << 2);
+
+ 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);
+ }
+}
+++ /dev/null
-#pragma OPENCL EXTENSION cl_khr_fp16 : enable
-#pragma OPENCL EXTENSION cl_khr_subgroups : enable
-
-#ifdef 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")))
-#endif
-
-// assume
-#define QK4_0 32
-#define N_SIMDGROUP 4
-
-#define dequantizeBlockAccum_ns_sgbroadcast_1_hi(total_sums, bits4, scale, y) \
- float shared_y; \
- shared_y = sub_group_broadcast(y.s0, 0); \
- total_sums.s0 += ((bits4.s0 & 0x000F) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += ((bits4.s1 & 0x000F) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s1, 0); \
- total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s2, 0); \
- total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s3, 0); \
- total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s4, 0); \
- total_sums.s0 += ((bits4.s2 & 0x000F) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += ((bits4.s3 & 0x000F) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s5, 0); \
- total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s6, 0); \
- total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s7, 0); \
- total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s0, 1); \
- total_sums.s0 += ((bits4.s4 & 0x000F) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += ((bits4.s5 & 0x000F) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s1, 1); \
- total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s2, 1); \
- total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s3, 1); \
- total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s4, 1); \
- total_sums.s0 += ((bits4.s6 & 0x000F) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += ((bits4.s7 & 0x000F) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s5, 1); \
- total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s6, 1); \
- total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s7, 1); \
- total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
-
-
-#define dequantizeBlockAccum_ns_sgbroadcast_1_lo(total_sums, bits4, scale, y) \
- shared_y = sub_group_broadcast(y.s0, 2); \
- total_sums.s0 += ((bits4.s0 & 0x000F) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += ((bits4.s1 & 0x000F) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s1, 2); \
- total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s2, 2); \
- total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s3, 2); \
- total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s4, 2); \
- total_sums.s0 += ((bits4.s2 & 0x000F) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += ((bits4.s3 & 0x000F) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s5, 2); \
- total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s6, 2); \
- total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s7, 2); \
- total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s0, 3); \
- total_sums.s0 += ((bits4.s4 & 0x000F) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += ((bits4.s5 & 0x000F) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s1, 3); \
- total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s2, 3); \
- total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s3, 3); \
- total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s4, 3); \
- total_sums.s0 += ((bits4.s6 & 0x000F) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += ((bits4.s7 & 0x000F) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s5, 3); \
- total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s6, 3); \
- total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s7, 3); \
- total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
-
-
-#define dequantizeBlockAccum_ns_sgbroadcast_8_hi(total_sums, bits4, scale, y) \
- float8 shared_y; \
- shared_y = sub_group_broadcast(y, 0); \
- total_sums.s0 += ((bits4.s0 & 0x000F) - 8) * scale.s0 * shared_y.s0; \
- total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s1; \
- total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s2; \
- total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s3; \
- total_sums.s0 += ((bits4.s2 & 0x000F) - 8) * scale.s0 * shared_y.s4; \
- total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s5; \
- total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s6; \
- total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s7; \
- total_sums.s1 += ((bits4.s1 & 0x000F) - 8) * scale.s1 * shared_y.s0; \
- total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s1; \
- total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s2; \
- total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s3; \
- total_sums.s1 += ((bits4.s3 & 0x000F) - 8) * scale.s1 * shared_y.s4; \
- total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s5; \
- total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s6; \
- total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s7; \
- shared_y = sub_group_broadcast(y, 1); \
- total_sums.s0 += ((bits4.s4 & 0x000F) - 8) * scale.s0 * shared_y.s0; \
- total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s1; \
- total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s2; \
- total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s3; \
- total_sums.s0 += ((bits4.s6 & 0x000F) - 8) * scale.s0 * shared_y.s4; \
- total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s5; \
- total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s6; \
- total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s7; \
- total_sums.s1 += ((bits4.s5 & 0x000F) - 8) * scale.s1 * shared_y.s0; \
- total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s1; \
- total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s2; \
- total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s3; \
- total_sums.s1 += ((bits4.s7 & 0x000F) - 8) * scale.s1 * shared_y.s4; \
- total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s5; \
- total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s6; \
- total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s7; \
-
-
-#define dequantizeBlockAccum_ns_sgbroadcast_8_lo(total_sums, bits4, scale, y) \
- shared_y = sub_group_broadcast(y, 2); \
- total_sums.s0 += ((bits4.s0 & 0x000F) - 8) * scale.s0 * shared_y.s0; \
- total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s1; \
- total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s2; \
- total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s3; \
- total_sums.s0 += ((bits4.s2 & 0x000F) - 8) * scale.s0 * shared_y.s4; \
- total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s5; \
- total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s6; \
- total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s7; \
- total_sums.s1 += ((bits4.s1 & 0x000F) - 8) * scale.s1 * shared_y.s0; \
- total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s1; \
- total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s2; \
- total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s3; \
- total_sums.s1 += ((bits4.s3 & 0x000F) - 8) * scale.s1 * shared_y.s4; \
- total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s5; \
- total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s6; \
- total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s7; \
- shared_y = sub_group_broadcast(y, 3); \
- total_sums.s0 += ((bits4.s4 & 0x000F) - 8) * scale.s0 * shared_y.s0; \
- total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s1; \
- total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s2; \
- total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s3; \
- total_sums.s0 += ((bits4.s6 & 0x000F) - 8) * scale.s0 * shared_y.s4; \
- total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s5; \
- total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s6; \
- total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s7; \
- total_sums.s1 += ((bits4.s5 & 0x000F) - 8) * scale.s1 * shared_y.s0; \
- total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s1; \
- total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s2; \
- total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s3; \
- total_sums.s1 += ((bits4.s7 & 0x000F) - 8) * scale.s1 * shared_y.s4; \
- total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s5; \
- total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s6; \
- total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s7; \
-
-#ifdef ADRENO_GPU
-REQD_SUBGROUP_SIZE_64
-#endif
-__kernel void kernel_gemv_noshuffle(
- __read_only image1d_buffer_t src0_q, // quantized A
- global half2 * src0_d, // A scales
- __read_only image1d_buffer_t src1, // B
- ulong offset1, // offset to B (0)
- global float * dst, // C
- ulong offsetd, // offset to C (0)
- uint K, // K
- int ne01, // M
- int ne02, // 1
- int ne10, // K
- int ne12, // 1
- int ne0, // M
- int ne1, // N
- int r2, // 1
- int r3)
-{
- uint groupId = get_local_id(1);
- uint gid = get_global_id(0);
- ushort slid = get_sub_group_local_id();
-
- __private uint4 regA;
- __private half2 regS;
- __private float8 regB;
-
- __private float2 totalSum = (float2)(0.0f);
-
- // loop along K in block granularity, skip 4 blocks every iter
- for (uint k = groupId; k < (K / QK4_0); k += N_SIMDGROUP) {
- regS = src0_d[gid + k * LINE_STRIDE_A]; // each fiber loads scale of two rows
- // first 4 fibers in each wave load 8 B values to its private scope
- if (slid < 4) {
- regB.s0123 = read_imagef(src1, (slid * 2 + k * 8));
- regB.s4567 = read_imagef(src1, (1 + slid * 2 + k * 8));
- }
-
- // load half weights for two blocks in consecutive rows
- regA.s0 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x;
- regA.s1 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x;
- regA.s2 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x;
- regA.s3 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x;
-#ifdef VECTOR_SUB_GROUP_BROADCAT
- dequantizeBlockAccum_ns_sgbroadcast_8_hi(totalSum, as_ushort8(regA), regS, regB);
-#else
- dequantizeBlockAccum_ns_sgbroadcast_1_hi(totalSum, as_ushort8(regA), regS, regB);
-#endif // VECTOR_SUB_GROUP_BROADCAT
-
- regA.s0 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x;
- regA.s1 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x;
- regA.s2 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x;
- regA.s3 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x;
-#ifdef VECTOR_SUB_GROUP_BROADCAT
- dequantizeBlockAccum_ns_sgbroadcast_8_lo(totalSum, as_ushort8(regA), regS, regB);
-#else
- dequantizeBlockAccum_ns_sgbroadcast_1_lo(totalSum, as_ushort8(regA), regS, regB);
-#endif // VECTOR_SUB_GROUP_BROADCAT
- }
-
- // reduction in local memory, assumes #wave=4
- __local float2 reduceLM[SIMDGROUP_WIDTH * 3];
- if (groupId == 1) reduceLM[SIMDGROUP_WIDTH * 0 + slid] = totalSum;
- if (groupId == 2) reduceLM[SIMDGROUP_WIDTH * 1 + slid] = totalSum;
- if (groupId == 3) reduceLM[SIMDGROUP_WIDTH * 2 + slid] = totalSum;
- barrier(CLK_LOCAL_MEM_FENCE);
- if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 0 + slid];
- if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 1 + slid];
- if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 2 + slid];
-
- // 2 outputs per fiber in wave 0
- if (groupId == 0) {
- dst = (global float*)((global char*)dst + offsetd);
- vstore2(totalSum, 0, &(dst[gid * 2]));
- }
-
-}
+++ /dev/null
-#pragma OPENCL EXTENSION cl_khr_fp16 : enable
-#pragma OPENCL EXTENSION cl_khr_subgroups : enable
-
-#ifdef 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")))
-#endif
-
-// assume
-#define QK4_0 32
-#define N_SIMDGROUP 4
-
-#define dequantizeBlockAccum_ns_sgbroadcast_1_hi(total_sums, bits4, scale, y) \
- float shared_y; \
- shared_y = sub_group_broadcast(y.s0, 0); \
- total_sums.s0 += ((bits4.s0 & 0x000F) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += ((bits4.s1 & 0x000F) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s1, 0); \
- total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s2, 0); \
- total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s3, 0); \
- total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s4, 0); \
- total_sums.s0 += ((bits4.s2 & 0x000F) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += ((bits4.s3 & 0x000F) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s5, 0); \
- total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s6, 0); \
- total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s7, 0); \
- total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s0, 1); \
- total_sums.s0 += ((bits4.s4 & 0x000F) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += ((bits4.s5 & 0x000F) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s1, 1); \
- total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s2, 1); \
- total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s3, 1); \
- total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s4, 1); \
- total_sums.s0 += ((bits4.s6 & 0x000F) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += ((bits4.s7 & 0x000F) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s5, 1); \
- total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s6, 1); \
- total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s7, 1); \
- total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
-
-
-#define dequantizeBlockAccum_ns_sgbroadcast_1_lo(total_sums, bits4, scale, y) \
- shared_y = sub_group_broadcast(y.s0, 2); \
- total_sums.s0 += ((bits4.s0 & 0x000F) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += ((bits4.s1 & 0x000F) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s1, 2); \
- total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s2, 2); \
- total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s3, 2); \
- total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s4, 2); \
- total_sums.s0 += ((bits4.s2 & 0x000F) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += ((bits4.s3 & 0x000F) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s5, 2); \
- total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s6, 2); \
- total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s7, 2); \
- total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s0, 3); \
- total_sums.s0 += ((bits4.s4 & 0x000F) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += ((bits4.s5 & 0x000F) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s1, 3); \
- total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s2, 3); \
- total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s3, 3); \
- total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s4, 3); \
- total_sums.s0 += ((bits4.s6 & 0x000F) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += ((bits4.s7 & 0x000F) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s5, 3); \
- total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s6, 3); \
- total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
- shared_y = sub_group_broadcast(y.s7, 3); \
- total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
- total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
-
-
-#define dequantizeBlockAccum_ns_sgbroadcast_8_hi(total_sums, bits4, scale, y) \
- float8 shared_y; \
- shared_y = sub_group_broadcast(y, 0); \
- total_sums.s0 += ((bits4.s0 & 0x000F) - 8) * scale.s0 * shared_y.s0; \
- total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s1; \
- total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s2; \
- total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s3; \
- total_sums.s0 += ((bits4.s2 & 0x000F) - 8) * scale.s0 * shared_y.s4; \
- total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s5; \
- total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s6; \
- total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s7; \
- total_sums.s1 += ((bits4.s1 & 0x000F) - 8) * scale.s1 * shared_y.s0; \
- total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s1; \
- total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s2; \
- total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s3; \
- total_sums.s1 += ((bits4.s3 & 0x000F) - 8) * scale.s1 * shared_y.s4; \
- total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s5; \
- total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s6; \
- total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s7; \
- shared_y = sub_group_broadcast(y, 1); \
- total_sums.s0 += ((bits4.s4 & 0x000F) - 8) * scale.s0 * shared_y.s0; \
- total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s1; \
- total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s2; \
- total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s3; \
- total_sums.s0 += ((bits4.s6 & 0x000F) - 8) * scale.s0 * shared_y.s4; \
- total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s5; \
- total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s6; \
- total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s7; \
- total_sums.s1 += ((bits4.s5 & 0x000F) - 8) * scale.s1 * shared_y.s0; \
- total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s1; \
- total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s2; \
- total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s3; \
- total_sums.s1 += ((bits4.s7 & 0x000F) - 8) * scale.s1 * shared_y.s4; \
- total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s5; \
- total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s6; \
- total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s7; \
-
-
-#define dequantizeBlockAccum_ns_sgbroadcast_8_lo(total_sums, bits4, scale, y) \
- shared_y = sub_group_broadcast(y, 2); \
- total_sums.s0 += ((bits4.s0 & 0x000F) - 8) * scale.s0 * shared_y.s0; \
- total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s1; \
- total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s2; \
- total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s3; \
- total_sums.s0 += ((bits4.s2 & 0x000F) - 8) * scale.s0 * shared_y.s4; \
- total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s5; \
- total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s6; \
- total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s7; \
- total_sums.s1 += ((bits4.s1 & 0x000F) - 8) * scale.s1 * shared_y.s0; \
- total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s1; \
- total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s2; \
- total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s3; \
- total_sums.s1 += ((bits4.s3 & 0x000F) - 8) * scale.s1 * shared_y.s4; \
- total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s5; \
- total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s6; \
- total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s7; \
- shared_y = sub_group_broadcast(y, 3); \
- total_sums.s0 += ((bits4.s4 & 0x000F) - 8) * scale.s0 * shared_y.s0; \
- total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s1; \
- total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s2; \
- total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s3; \
- total_sums.s0 += ((bits4.s6 & 0x000F) - 8) * scale.s0 * shared_y.s4; \
- total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s5; \
- total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s6; \
- total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s7; \
- total_sums.s1 += ((bits4.s5 & 0x000F) - 8) * scale.s1 * shared_y.s0; \
- total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s1; \
- total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s2; \
- total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s3; \
- total_sums.s1 += ((bits4.s7 & 0x000F) - 8) * scale.s1 * shared_y.s4; \
- total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s5; \
- total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s6; \
- total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s7; \
-
-#ifdef ADRENO_GPU
-REQD_SUBGROUP_SIZE_64
-#endif
-__kernel void kernel_gemv_noshuffle(
- __read_only image1d_buffer_t src0_q, // quantized A
- global half2 * src0_d, // A scales
- __read_only image1d_buffer_t src1, // B
- ulong offset1, // offset to B (0)
- global float * dst, // C
- ulong offsetd, // offset to C (0)
- int ne00, // K
- int ne01, // M
- int ne02, // 1
- int ne10, // K
- int ne12, // 1
- int ne0, // M
- int ne1, // N
- int r2, // 1
- int r3)
-{
- uint groupId = get_local_id(1);
- uint gid = get_global_id(0);
- ushort slid = get_sub_group_local_id();
-
- uint K = ne00;
- uint M = ne01;
-
- uint LINE_STRIDE_A = M / 2;
- uint BLOCK_STRIDE_A = N_SIMDGROUP * M;
-
- __private uint4 regA;
- __private half2 regS;
- __private float8 regB;
-
- __private float2 totalSum = (float2)(0.0f);
-
- // loop along K in block granularity, skip 4 blocks every iter
- for (uint k = groupId; k < (K / QK4_0); k += N_SIMDGROUP) {
- regS = src0_d[gid + k * LINE_STRIDE_A]; // each fiber loads scale of two rows
- // first 4 fibers in each wave load 8 B values to its private scope
- if (slid < 4) {
- regB.s0123 = read_imagef(src1, (slid * 2 + k * 8));
- regB.s4567 = read_imagef(src1, (1 + slid * 2 + k * 8));
- }
-
- // load half weights for two blocks in consecutive rows
- regA.s0 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x;
- regA.s1 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x;
- regA.s2 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x;
- regA.s3 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x;
-#ifdef VECTOR_SUB_GROUP_BROADCAT
- dequantizeBlockAccum_ns_sgbroadcast_8_hi(totalSum, as_ushort8(regA), regS, regB);
-#else
- dequantizeBlockAccum_ns_sgbroadcast_1_hi(totalSum, as_ushort8(regA), regS, regB);
-#endif // VECTOR_SUB_GROUP_BROADCAT
-
- regA.s0 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x;
- regA.s1 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x;
- regA.s2 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x;
- regA.s3 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x;
-#ifdef VECTOR_SUB_GROUP_BROADCAT
- dequantizeBlockAccum_ns_sgbroadcast_8_lo(totalSum, as_ushort8(regA), regS, regB);
-#else
- dequantizeBlockAccum_ns_sgbroadcast_1_lo(totalSum, as_ushort8(regA), regS, regB);
-#endif // VECTOR_SUB_GROUP_BROADCAT
- }
-
- // reduction in local memory, assumes #wave=4
- __local float2 reduceLM[SIMDGROUP_WIDTH * 3];
- if (groupId == 1) reduceLM[SIMDGROUP_WIDTH * 0 + slid] = totalSum;
- if (groupId == 2) reduceLM[SIMDGROUP_WIDTH * 1 + slid] = totalSum;
- if (groupId == 3) reduceLM[SIMDGROUP_WIDTH * 2 + slid] = totalSum;
- barrier(CLK_LOCAL_MEM_FENCE);
- if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 0 + slid];
- if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 1 + slid];
- if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 2 + slid];
-
- // 2 outputs per fiber in wave 0
- if (groupId == 0) {
- dst = (global float*)((global char*)dst + offsetd);
- vstore2(totalSum, 0, &(dst[gid * 2]));
- }
-
-}
+++ /dev/null
-#pragma OPENCL EXTENSION cl_khr_fp16 : enable
-#pragma OPENCL EXTENSION cl_khr_subgroups : enable
-
-#ifdef 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")))
-#endif
-
-#define QK8_0 32
-#define N_SIMDGROUP 4
-
-#define dequantizeBlockAccum_ns_sgbroadcast_1(total_sums, bits8, scale, y) \
- float shared_y; \
- char elem; \
- \
- shared_y = sub_group_broadcast(y.s0, 0); \
- elem = (char)(bits8.s0 & 0x000000FF); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s1, 0); \
- elem = (char)((bits8.s0 & 0x0000FF00) >> 8); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s2, 0); \
- elem = (char)((bits8.s0 & 0x00FF0000) >> 16); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s3, 0); \
- elem = (char)((bits8.s0 & 0xFF000000) >> 24); \
- total_sums += convert_int(elem) * scale * shared_y; \
- \
- shared_y = sub_group_broadcast(y.s4, 0); \
- elem = (char)(bits8.s1 & 0x000000FF); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s5, 0); \
- elem = (char)((bits8.s1 & 0x0000FF00) >> 8); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s6, 0); \
- elem = (char)((bits8.s1 & 0x00FF0000) >> 16); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s7, 0); \
- elem = (char)((bits8.s1 & 0xFF000000) >> 24); \
- total_sums += convert_int(elem) * scale * shared_y; \
- \
- shared_y = sub_group_broadcast(y.s0, 1); \
- elem = (char)(bits8.s2 & 0x000000FF); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s1, 1); \
- elem = (char)((bits8.s2 & 0x0000FF00) >> 8); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s2, 1); \
- elem = (char)((bits8.s2 & 0x00FF0000) >> 16); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s3, 1); \
- elem = (char)((bits8.s2 & 0xFF000000) >> 24); \
- total_sums += convert_int(elem) * scale * shared_y; \
- \
- shared_y = sub_group_broadcast(y.s4, 1); \
- elem = (char)(bits8.s3 & 0x000000FF); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s5, 1); \
- elem = (char)((bits8.s3 & 0x0000FF00) >> 8); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s6, 1); \
- elem = (char)((bits8.s3 & 0x00FF0000) >> 16); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s7, 1); \
- elem = (char)((bits8.s3 & 0xFF000000) >> 24); \
- total_sums += convert_int(elem) * scale * shared_y; \
- \
- shared_y = sub_group_broadcast(y.s0, 2); \
- elem = (char)(bits8.s4 & 0x000000FF); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s1, 2); \
- elem = (char)((bits8.s4 & 0x0000FF00) >> 8); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s2, 2); \
- elem = (char)((bits8.s4 & 0x00FF0000) >> 16); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s3, 2); \
- elem = (char)((bits8.s4 & 0xFF000000) >> 24); \
- total_sums += convert_int(elem) * scale * shared_y; \
- \
- shared_y = sub_group_broadcast(y.s4, 2); \
- elem = (char)(bits8.s5 & 0x000000FF); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s5, 2); \
- elem = (char)((bits8.s5 & 0x0000FF00) >> 8); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s6, 2); \
- elem = (char)((bits8.s5 & 0x00FF0000) >> 16); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s7, 2); \
- elem = (char)((bits8.s5 & 0xFF000000) >> 24); \
- total_sums += convert_int(elem) * scale * shared_y; \
- \
- shared_y = sub_group_broadcast(y.s0, 3); \
- elem = (char)(bits8.s6 & 0x000000FF); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s1, 3); \
- elem = (char)((bits8.s6 & 0x0000FF00) >> 8); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s2, 3); \
- elem = (char)((bits8.s6 & 0x00FF0000) >> 16); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s3, 3); \
- elem = (char)((bits8.s6 & 0xFF000000) >> 24); \
- total_sums += convert_int(elem) * scale * shared_y; \
- \
- shared_y = sub_group_broadcast(y.s4, 3); \
- elem = (char)(bits8.s7 & 0x000000FF); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s5, 3); \
- elem = (char)((bits8.s7 & 0x0000FF00) >> 8); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s6, 3); \
- elem = (char)((bits8.s7 & 0x00FF0000) >> 16); \
- total_sums += convert_int(elem) * scale * shared_y; \
- shared_y = sub_group_broadcast(y.s7, 3); \
- elem = (char)((bits8.s7 & 0xFF000000) >> 24); \
- total_sums += convert_int(elem) * scale * shared_y; \
-
-#ifdef ADRENO_GPU
-REQD_SUBGROUP_SIZE_64
-#endif
-__kernel void kernel_gemv_noshuffle_q8_0_f32(
- __read_only image1d_buffer_t src0_q, // quantized A
- global half * src0_d, // A scales
- __read_only image1d_buffer_t src1, // B
- ulong offset1, // offset to B (0)
- global float * dst, // C
- ulong offsetd, // offset to C
- int ne00, // K
- int ne01, // M
- int ne02, // 1
- int ne10, // K
- int ne12, // 1
- int ne0, // M
- int ne1, // N
- int r2, // 1
- int r3)
-{
- uint groupId = get_local_id(1);
- uint gid = get_global_id(0);
- ushort slid = get_sub_group_local_id();
-
- uint K = ne00;
- uint M = ne01;
-
- uint LINE_STRIDE_A = M;
- uint BLOCK_STRIDE_A = 8 * M; // 32 / 4 = 8
-
- __private uint8 regA;
- __private half regS;
- __private float8 regB;
-
- __private float totalSum = (float)(0.0f);
-
- // loop along K in block granularity, skip 4 blocks every iter
- #pragma unroll 1 /* tell compiler not to unroll */
- for (uint k = groupId; k < (K / QK8_0); k += N_SIMDGROUP) {
- regS = src0_d[gid + k * LINE_STRIDE_A]; // each fiber loads scale of one rows
- // first 4 fibers in each wave load 8 B values to its private scope
- if (slid < 4) {
- regB.s0123 = read_imagef(src1, (slid * 2 + k * 8));
- regB.s4567 = read_imagef(src1, (1 + slid * 2 + k * 8));
- }
-
- // load weights for one block in consecutive rows
- regA.s0 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x;
- regA.s1 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x;
- regA.s2 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x;
- regA.s3 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x;
- regA.s4 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x;
- regA.s5 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x;
- regA.s6 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x;
- regA.s7 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x;
-
- dequantizeBlockAccum_ns_sgbroadcast_1(totalSum, regA, regS, regB);
- }
-
- // reduction in local memory, assumes #wave=4
- __local float reduceLM[SIMDGROUP_WIDTH * 3];
- if (groupId == 1) reduceLM[SIMDGROUP_WIDTH * 0 + slid] = totalSum;
- if (groupId == 2) reduceLM[SIMDGROUP_WIDTH * 1 + slid] = totalSum;
- if (groupId == 3) reduceLM[SIMDGROUP_WIDTH * 2 + slid] = totalSum;
- barrier(CLK_LOCAL_MEM_FENCE);
- if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 0 + slid];
- if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 1 + slid];
- if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 2 + slid];
-
- // 1 outputs per fiber in wave 0
- if (groupId == 0) {
- dst = (global float*)((global char*)dst + offsetd);
- dst[gid] = totalSum;
- }
-}
--- /dev/null
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_khr_subgroups : enable
+
+#ifdef 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")))
+#endif
+
+// assume
+#define QK4_0 32
+#define N_SIMDGROUP 4
+
+#define dequantizeBlockAccum_ns_sgbroadcast_1_hi(total_sums, bits4, scale, y) \
+ float shared_y; \
+ shared_y = sub_group_broadcast(y.s0, 0); \
+ total_sums.s0 += ((bits4.s0 & 0x000F) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += ((bits4.s1 & 0x000F) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s1, 0); \
+ total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s2, 0); \
+ total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s3, 0); \
+ total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s4, 0); \
+ total_sums.s0 += ((bits4.s2 & 0x000F) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += ((bits4.s3 & 0x000F) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s5, 0); \
+ total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s6, 0); \
+ total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s7, 0); \
+ total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s0, 1); \
+ total_sums.s0 += ((bits4.s4 & 0x000F) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += ((bits4.s5 & 0x000F) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s1, 1); \
+ total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s2, 1); \
+ total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s3, 1); \
+ total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s4, 1); \
+ total_sums.s0 += ((bits4.s6 & 0x000F) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += ((bits4.s7 & 0x000F) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s5, 1); \
+ total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s6, 1); \
+ total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s7, 1); \
+ total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
+
+
+#define dequantizeBlockAccum_ns_sgbroadcast_1_lo(total_sums, bits4, scale, y) \
+ shared_y = sub_group_broadcast(y.s0, 2); \
+ total_sums.s0 += ((bits4.s0 & 0x000F) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += ((bits4.s1 & 0x000F) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s1, 2); \
+ total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s2, 2); \
+ total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s3, 2); \
+ total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s4, 2); \
+ total_sums.s0 += ((bits4.s2 & 0x000F) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += ((bits4.s3 & 0x000F) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s5, 2); \
+ total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s6, 2); \
+ total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s7, 2); \
+ total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s0, 3); \
+ total_sums.s0 += ((bits4.s4 & 0x000F) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += ((bits4.s5 & 0x000F) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s1, 3); \
+ total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s2, 3); \
+ total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s3, 3); \
+ total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s4, 3); \
+ total_sums.s0 += ((bits4.s6 & 0x000F) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += ((bits4.s7 & 0x000F) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s5, 3); \
+ total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s6, 3); \
+ total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s7, 3); \
+ total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
+
+
+#define dequantizeBlockAccum_ns_sgbroadcast_8_hi(total_sums, bits4, scale, y) \
+ float8 shared_y; \
+ shared_y = sub_group_broadcast(y, 0); \
+ total_sums.s0 += ((bits4.s0 & 0x000F) - 8) * scale.s0 * shared_y.s0; \
+ total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s1; \
+ total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s2; \
+ total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s3; \
+ total_sums.s0 += ((bits4.s2 & 0x000F) - 8) * scale.s0 * shared_y.s4; \
+ total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s5; \
+ total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s6; \
+ total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s7; \
+ total_sums.s1 += ((bits4.s1 & 0x000F) - 8) * scale.s1 * shared_y.s0; \
+ total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s1; \
+ total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s2; \
+ total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s3; \
+ total_sums.s1 += ((bits4.s3 & 0x000F) - 8) * scale.s1 * shared_y.s4; \
+ total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s5; \
+ total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s6; \
+ total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s7; \
+ shared_y = sub_group_broadcast(y, 1); \
+ total_sums.s0 += ((bits4.s4 & 0x000F) - 8) * scale.s0 * shared_y.s0; \
+ total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s1; \
+ total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s2; \
+ total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s3; \
+ total_sums.s0 += ((bits4.s6 & 0x000F) - 8) * scale.s0 * shared_y.s4; \
+ total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s5; \
+ total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s6; \
+ total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s7; \
+ total_sums.s1 += ((bits4.s5 & 0x000F) - 8) * scale.s1 * shared_y.s0; \
+ total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s1; \
+ total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s2; \
+ total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s3; \
+ total_sums.s1 += ((bits4.s7 & 0x000F) - 8) * scale.s1 * shared_y.s4; \
+ total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s5; \
+ total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s6; \
+ total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s7; \
+
+
+#define dequantizeBlockAccum_ns_sgbroadcast_8_lo(total_sums, bits4, scale, y) \
+ shared_y = sub_group_broadcast(y, 2); \
+ total_sums.s0 += ((bits4.s0 & 0x000F) - 8) * scale.s0 * shared_y.s0; \
+ total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s1; \
+ total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s2; \
+ total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s3; \
+ total_sums.s0 += ((bits4.s2 & 0x000F) - 8) * scale.s0 * shared_y.s4; \
+ total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s5; \
+ total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s6; \
+ total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s7; \
+ total_sums.s1 += ((bits4.s1 & 0x000F) - 8) * scale.s1 * shared_y.s0; \
+ total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s1; \
+ total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s2; \
+ total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s3; \
+ total_sums.s1 += ((bits4.s3 & 0x000F) - 8) * scale.s1 * shared_y.s4; \
+ total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s5; \
+ total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s6; \
+ total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s7; \
+ shared_y = sub_group_broadcast(y, 3); \
+ total_sums.s0 += ((bits4.s4 & 0x000F) - 8) * scale.s0 * shared_y.s0; \
+ total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s1; \
+ total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s2; \
+ total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s3; \
+ total_sums.s0 += ((bits4.s6 & 0x000F) - 8) * scale.s0 * shared_y.s4; \
+ total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s5; \
+ total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s6; \
+ total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s7; \
+ total_sums.s1 += ((bits4.s5 & 0x000F) - 8) * scale.s1 * shared_y.s0; \
+ total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s1; \
+ total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s2; \
+ total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s3; \
+ total_sums.s1 += ((bits4.s7 & 0x000F) - 8) * scale.s1 * shared_y.s4; \
+ total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s5; \
+ total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s6; \
+ total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s7; \
+
+#ifdef ADRENO_GPU
+REQD_SUBGROUP_SIZE_64
+#endif
+__kernel void kernel_gemv_noshuffle_q4_0_f32(
+ __read_only image1d_buffer_t src0_q, // quantized A
+ global half2 * src0_d, // A scales
+ __read_only image1d_buffer_t src1, // B
+ ulong offset1, // offset to B (0)
+ global float * dst, // C
+ ulong offsetd, // offset to C (0)
+ int ne00, // K
+ int ne01, // M
+ int ne02, // 1
+ int ne10, // K
+ int ne12, // 1
+ int ne0, // M
+ int ne1, // N
+ int r2, // 1
+ int r3)
+{
+ uint groupId = get_local_id(1);
+ uint gid = get_global_id(0);
+ ushort slid = get_sub_group_local_id();
+
+ uint K = ne00;
+ uint M = ne01;
+
+ uint LINE_STRIDE_A = M / 2;
+ uint BLOCK_STRIDE_A = N_SIMDGROUP * M;
+
+ __private uint4 regA;
+ __private half2 regS;
+ __private float8 regB;
+
+ __private float2 totalSum = (float2)(0.0f);
+
+ // loop along K in block granularity, skip 4 blocks every iter
+ for (uint k = groupId; k < (K / QK4_0); k += N_SIMDGROUP) {
+ regS = src0_d[gid + k * LINE_STRIDE_A]; // each fiber loads scale of two rows
+ // first 4 fibers in each wave load 8 B values to its private scope
+ if (slid < 4) {
+ regB.s0123 = read_imagef(src1, (slid * 2 + k * 8));
+ regB.s4567 = read_imagef(src1, (1 + slid * 2 + k * 8));
+ }
+
+ // load half weights for two blocks in consecutive rows
+ regA.s0 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x;
+ regA.s1 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x;
+ regA.s2 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x;
+ regA.s3 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x;
+#ifdef VECTOR_SUB_GROUP_BROADCAST
+ dequantizeBlockAccum_ns_sgbroadcast_8_hi(totalSum, as_ushort8(regA), regS, regB);
+#else
+ dequantizeBlockAccum_ns_sgbroadcast_1_hi(totalSum, as_ushort8(regA), regS, regB);
+#endif // VECTOR_SUB_GROUP_BROADCAST
+
+ regA.s0 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x;
+ regA.s1 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x;
+ regA.s2 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x;
+ regA.s3 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x;
+#ifdef VECTOR_SUB_GROUP_BROADCAST
+ dequantizeBlockAccum_ns_sgbroadcast_8_lo(totalSum, as_ushort8(regA), regS, regB);
+#else
+ dequantizeBlockAccum_ns_sgbroadcast_1_lo(totalSum, as_ushort8(regA), regS, regB);
+#endif // VECTOR_SUB_GROUP_BROADCAST
+ }
+
+ // reduction in local memory, assumes #wave=4
+ __local float2 reduceLM[SIMDGROUP_WIDTH * 3];
+ if (groupId == 1) reduceLM[SIMDGROUP_WIDTH * 0 + slid] = totalSum;
+ if (groupId == 2) reduceLM[SIMDGROUP_WIDTH * 1 + slid] = totalSum;
+ if (groupId == 3) reduceLM[SIMDGROUP_WIDTH * 2 + slid] = totalSum;
+ barrier(CLK_LOCAL_MEM_FENCE);
+ if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 0 + slid];
+ if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 1 + slid];
+ if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 2 + slid];
+
+ // 2 outputs per fiber in wave 0
+ if (groupId == 0) {
+ dst = (global float*)((global char*)dst + offsetd);
+ vstore2(totalSum, 0, &(dst[gid * 2]));
+ }
+
+}
--- /dev/null
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_khr_subgroups : enable
+
+#ifdef 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")))
+#endif
+
+// assume
+#define QK4_0 32
+#define N_SIMDGROUP 4
+
+#define dequantizeBlockAccum_ns_sgbroadcast_1_hi(total_sums, bits4, scale, y) \
+ float shared_y; \
+ shared_y = sub_group_broadcast(y.s0, 0); \
+ total_sums.s0 += ((bits4.s0 & 0x000F) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += ((bits4.s1 & 0x000F) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s1, 0); \
+ total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s2, 0); \
+ total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s3, 0); \
+ total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s4, 0); \
+ total_sums.s0 += ((bits4.s2 & 0x000F) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += ((bits4.s3 & 0x000F) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s5, 0); \
+ total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s6, 0); \
+ total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s7, 0); \
+ total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s0, 1); \
+ total_sums.s0 += ((bits4.s4 & 0x000F) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += ((bits4.s5 & 0x000F) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s1, 1); \
+ total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s2, 1); \
+ total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s3, 1); \
+ total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s4, 1); \
+ total_sums.s0 += ((bits4.s6 & 0x000F) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += ((bits4.s7 & 0x000F) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s5, 1); \
+ total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s6, 1); \
+ total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s7, 1); \
+ total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
+
+
+#define dequantizeBlockAccum_ns_sgbroadcast_1_lo(total_sums, bits4, scale, y) \
+ shared_y = sub_group_broadcast(y.s0, 2); \
+ total_sums.s0 += ((bits4.s0 & 0x000F) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += ((bits4.s1 & 0x000F) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s1, 2); \
+ total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s2, 2); \
+ total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s3, 2); \
+ total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s4, 2); \
+ total_sums.s0 += ((bits4.s2 & 0x000F) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += ((bits4.s3 & 0x000F) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s5, 2); \
+ total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s6, 2); \
+ total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s7, 2); \
+ total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s0, 3); \
+ total_sums.s0 += ((bits4.s4 & 0x000F) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += ((bits4.s5 & 0x000F) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s1, 3); \
+ total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s2, 3); \
+ total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s3, 3); \
+ total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s4, 3); \
+ total_sums.s0 += ((bits4.s6 & 0x000F) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += ((bits4.s7 & 0x000F) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s5, 3); \
+ total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s6, 3); \
+ total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y; \
+ shared_y = sub_group_broadcast(y.s7, 3); \
+ total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) - 8) * scale.s0 * shared_y; \
+ total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) - 8) * scale.s1 * shared_y; \
+
+
+#define dequantizeBlockAccum_ns_sgbroadcast_8_hi(total_sums, bits4, scale, y) \
+ float8 shared_y; \
+ shared_y = sub_group_broadcast(y, 0); \
+ total_sums.s0 += ((bits4.s0 & 0x000F) - 8) * scale.s0 * shared_y.s0; \
+ total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s1; \
+ total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s2; \
+ total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s3; \
+ total_sums.s0 += ((bits4.s2 & 0x000F) - 8) * scale.s0 * shared_y.s4; \
+ total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s5; \
+ total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s6; \
+ total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s7; \
+ total_sums.s1 += ((bits4.s1 & 0x000F) - 8) * scale.s1 * shared_y.s0; \
+ total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s1; \
+ total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s2; \
+ total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s3; \
+ total_sums.s1 += ((bits4.s3 & 0x000F) - 8) * scale.s1 * shared_y.s4; \
+ total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s5; \
+ total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s6; \
+ total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s7; \
+ shared_y = sub_group_broadcast(y, 1); \
+ total_sums.s0 += ((bits4.s4 & 0x000F) - 8) * scale.s0 * shared_y.s0; \
+ total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s1; \
+ total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s2; \
+ total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s3; \
+ total_sums.s0 += ((bits4.s6 & 0x000F) - 8) * scale.s0 * shared_y.s4; \
+ total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s5; \
+ total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s6; \
+ total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s7; \
+ total_sums.s1 += ((bits4.s5 & 0x000F) - 8) * scale.s1 * shared_y.s0; \
+ total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s1; \
+ total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s2; \
+ total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s3; \
+ total_sums.s1 += ((bits4.s7 & 0x000F) - 8) * scale.s1 * shared_y.s4; \
+ total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s5; \
+ total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s6; \
+ total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s7; \
+
+
+#define dequantizeBlockAccum_ns_sgbroadcast_8_lo(total_sums, bits4, scale, y) \
+ shared_y = sub_group_broadcast(y, 2); \
+ total_sums.s0 += ((bits4.s0 & 0x000F) - 8) * scale.s0 * shared_y.s0; \
+ total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s1; \
+ total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s2; \
+ total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s3; \
+ total_sums.s0 += ((bits4.s2 & 0x000F) - 8) * scale.s0 * shared_y.s4; \
+ total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s5; \
+ total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s6; \
+ total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s7; \
+ total_sums.s1 += ((bits4.s1 & 0x000F) - 8) * scale.s1 * shared_y.s0; \
+ total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s1; \
+ total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s2; \
+ total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s3; \
+ total_sums.s1 += ((bits4.s3 & 0x000F) - 8) * scale.s1 * shared_y.s4; \
+ total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s5; \
+ total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s6; \
+ total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s7; \
+ shared_y = sub_group_broadcast(y, 3); \
+ total_sums.s0 += ((bits4.s4 & 0x000F) - 8) * scale.s0 * shared_y.s0; \
+ total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s1; \
+ total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s2; \
+ total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s3; \
+ total_sums.s0 += ((bits4.s6 & 0x000F) - 8) * scale.s0 * shared_y.s4; \
+ total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) - 8) * scale.s0 * shared_y.s5; \
+ total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) - 8) * scale.s0 * shared_y.s6; \
+ total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) - 8) * scale.s0 * shared_y.s7; \
+ total_sums.s1 += ((bits4.s5 & 0x000F) - 8) * scale.s1 * shared_y.s0; \
+ total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s1; \
+ total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s2; \
+ total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s3; \
+ total_sums.s1 += ((bits4.s7 & 0x000F) - 8) * scale.s1 * shared_y.s4; \
+ total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) - 8) * scale.s1 * shared_y.s5; \
+ total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) - 8) * scale.s1 * shared_y.s6; \
+ total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) - 8) * scale.s1 * shared_y.s7; \
+
+#ifdef ADRENO_GPU
+REQD_SUBGROUP_SIZE_64
+#endif
+__kernel void kernel_gemv_noshuffle_q4_0_f32(
+ __read_only image1d_buffer_t src0_q, // quantized A
+ global half2 * src0_d, // A scales
+ __read_only image1d_buffer_t src1, // B
+ ulong offset1, // offset to B (0)
+ global float * dst, // C
+ ulong offsetd, // offset to C (0)
+ uint K, // K
+ int ne01, // M
+ int ne02, // 1
+ int ne10, // K
+ int ne12, // 1
+ int ne0, // M
+ int ne1, // N
+ int r2, // 1
+ int r3)
+{
+ uint groupId = get_local_id(1);
+ uint gid = get_global_id(0);
+ ushort slid = get_sub_group_local_id();
+
+ __private uint4 regA;
+ __private half2 regS;
+ __private float8 regB;
+
+ __private float2 totalSum = (float2)(0.0f);
+
+ // loop along K in block granularity, skip 4 blocks every iter
+ for (uint k = groupId; k < (K / QK4_0); k += N_SIMDGROUP) {
+ regS = src0_d[gid + k * LINE_STRIDE_A]; // each fiber loads scale of two rows
+ // first 4 fibers in each wave load 8 B values to its private scope
+ if (slid < 4) {
+ regB.s0123 = read_imagef(src1, (slid * 2 + k * 8));
+ regB.s4567 = read_imagef(src1, (1 + slid * 2 + k * 8));
+ }
+
+ // load half weights for two blocks in consecutive rows
+ regA.s0 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x;
+ regA.s1 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x;
+ regA.s2 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x;
+ regA.s3 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x;
+#ifdef VECTOR_SUB_GROUP_BROADCAST
+ dequantizeBlockAccum_ns_sgbroadcast_8_hi(totalSum, as_ushort8(regA), regS, regB);
+#else
+ dequantizeBlockAccum_ns_sgbroadcast_1_hi(totalSum, as_ushort8(regA), regS, regB);
+#endif // VECTOR_SUB_GROUP_BROADCAST
+
+ regA.s0 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x;
+ regA.s1 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x;
+ regA.s2 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x;
+ regA.s3 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x;
+#ifdef VECTOR_SUB_GROUP_BROADCAST
+ dequantizeBlockAccum_ns_sgbroadcast_8_lo(totalSum, as_ushort8(regA), regS, regB);
+#else
+ dequantizeBlockAccum_ns_sgbroadcast_1_lo(totalSum, as_ushort8(regA), regS, regB);
+#endif // VECTOR_SUB_GROUP_BROADCAST
+ }
+
+ // reduction in local memory, assumes #wave=4
+ __local float2 reduceLM[SIMDGROUP_WIDTH * 3];
+ if (groupId == 1) reduceLM[SIMDGROUP_WIDTH * 0 + slid] = totalSum;
+ if (groupId == 2) reduceLM[SIMDGROUP_WIDTH * 1 + slid] = totalSum;
+ if (groupId == 3) reduceLM[SIMDGROUP_WIDTH * 2 + slid] = totalSum;
+ barrier(CLK_LOCAL_MEM_FENCE);
+ if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 0 + slid];
+ if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 1 + slid];
+ if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 2 + slid];
+
+ // 2 outputs per fiber in wave 0
+ if (groupId == 0) {
+ dst = (global float*)((global char*)dst + offsetd);
+ vstore2(totalSum, 0, &(dst[gid * 2]));
+ }
+
+}
--- /dev/null
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_khr_subgroups : enable
+
+#ifdef 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")))
+#endif
+
+#define QK8_0 32
+#define N_SIMDGROUP 4
+
+#define dequantizeBlockAccum_ns_sgbroadcast_1(total_sums, bits8, scale, y) \
+ float shared_y; \
+ char elem; \
+ \
+ shared_y = sub_group_broadcast(y.s0, 0); \
+ elem = (char)(bits8.s0 & 0x000000FF); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s1, 0); \
+ elem = (char)((bits8.s0 & 0x0000FF00) >> 8); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s2, 0); \
+ elem = (char)((bits8.s0 & 0x00FF0000) >> 16); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s3, 0); \
+ elem = (char)((bits8.s0 & 0xFF000000) >> 24); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ \
+ shared_y = sub_group_broadcast(y.s4, 0); \
+ elem = (char)(bits8.s1 & 0x000000FF); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s5, 0); \
+ elem = (char)((bits8.s1 & 0x0000FF00) >> 8); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s6, 0); \
+ elem = (char)((bits8.s1 & 0x00FF0000) >> 16); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s7, 0); \
+ elem = (char)((bits8.s1 & 0xFF000000) >> 24); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ \
+ shared_y = sub_group_broadcast(y.s0, 1); \
+ elem = (char)(bits8.s2 & 0x000000FF); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s1, 1); \
+ elem = (char)((bits8.s2 & 0x0000FF00) >> 8); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s2, 1); \
+ elem = (char)((bits8.s2 & 0x00FF0000) >> 16); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s3, 1); \
+ elem = (char)((bits8.s2 & 0xFF000000) >> 24); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ \
+ shared_y = sub_group_broadcast(y.s4, 1); \
+ elem = (char)(bits8.s3 & 0x000000FF); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s5, 1); \
+ elem = (char)((bits8.s3 & 0x0000FF00) >> 8); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s6, 1); \
+ elem = (char)((bits8.s3 & 0x00FF0000) >> 16); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s7, 1); \
+ elem = (char)((bits8.s3 & 0xFF000000) >> 24); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ \
+ shared_y = sub_group_broadcast(y.s0, 2); \
+ elem = (char)(bits8.s4 & 0x000000FF); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s1, 2); \
+ elem = (char)((bits8.s4 & 0x0000FF00) >> 8); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s2, 2); \
+ elem = (char)((bits8.s4 & 0x00FF0000) >> 16); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s3, 2); \
+ elem = (char)((bits8.s4 & 0xFF000000) >> 24); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ \
+ shared_y = sub_group_broadcast(y.s4, 2); \
+ elem = (char)(bits8.s5 & 0x000000FF); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s5, 2); \
+ elem = (char)((bits8.s5 & 0x0000FF00) >> 8); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s6, 2); \
+ elem = (char)((bits8.s5 & 0x00FF0000) >> 16); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s7, 2); \
+ elem = (char)((bits8.s5 & 0xFF000000) >> 24); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ \
+ shared_y = sub_group_broadcast(y.s0, 3); \
+ elem = (char)(bits8.s6 & 0x000000FF); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s1, 3); \
+ elem = (char)((bits8.s6 & 0x0000FF00) >> 8); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s2, 3); \
+ elem = (char)((bits8.s6 & 0x00FF0000) >> 16); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s3, 3); \
+ elem = (char)((bits8.s6 & 0xFF000000) >> 24); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ \
+ shared_y = sub_group_broadcast(y.s4, 3); \
+ elem = (char)(bits8.s7 & 0x000000FF); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s5, 3); \
+ elem = (char)((bits8.s7 & 0x0000FF00) >> 8); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s6, 3); \
+ elem = (char)((bits8.s7 & 0x00FF0000) >> 16); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+ shared_y = sub_group_broadcast(y.s7, 3); \
+ elem = (char)((bits8.s7 & 0xFF000000) >> 24); \
+ total_sums += convert_int(elem) * scale * shared_y; \
+
+#ifdef ADRENO_GPU
+REQD_SUBGROUP_SIZE_64
+#endif
+__kernel void kernel_gemv_noshuffle_q8_0_f32(
+ __read_only image1d_buffer_t src0_q, // quantized A
+ global half * src0_d, // A scales
+ __read_only image1d_buffer_t src1, // B
+ ulong offset1, // offset to B (0)
+ global float * dst, // C
+ ulong offsetd, // offset to C
+ int ne00, // K
+ int ne01, // M
+ int ne02, // 1
+ int ne10, // K
+ int ne12, // 1
+ int ne0, // M
+ int ne1, // N
+ int r2, // 1
+ int r3)
+{
+ uint groupId = get_local_id(1);
+ uint gid = get_global_id(0);
+ ushort slid = get_sub_group_local_id();
+
+ uint K = ne00;
+ uint M = ne01;
+
+ uint LINE_STRIDE_A = M;
+ uint BLOCK_STRIDE_A = 8 * M; // 32 / 4 = 8
+
+ __private uint8 regA;
+ __private half regS;
+ __private float8 regB;
+
+ __private float totalSum = (float)(0.0f);
+
+ // loop along K in block granularity, skip 4 blocks every iter
+ #pragma unroll 1 /* tell compiler not to unroll */
+ for (uint k = groupId; k < (K / QK8_0); k += N_SIMDGROUP) {
+ regS = src0_d[gid + k * LINE_STRIDE_A]; // each fiber loads scale of one rows
+ // first 4 fibers in each wave load 8 B values to its private scope
+ if (slid < 4) {
+ regB.s0123 = read_imagef(src1, (slid * 2 + k * 8));
+ regB.s4567 = read_imagef(src1, (1 + slid * 2 + k * 8));
+ }
+
+ // load weights for one block in consecutive rows
+ regA.s0 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x;
+ regA.s1 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x;
+ regA.s2 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x;
+ regA.s3 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x;
+ regA.s4 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x;
+ regA.s5 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x;
+ regA.s6 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x;
+ regA.s7 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x;
+
+ dequantizeBlockAccum_ns_sgbroadcast_1(totalSum, regA, regS, regB);
+ }
+
+ // reduction in local memory, assumes #wave=4
+ __local float reduceLM[SIMDGROUP_WIDTH * 3];
+ if (groupId == 1) reduceLM[SIMDGROUP_WIDTH * 0 + slid] = totalSum;
+ if (groupId == 2) reduceLM[SIMDGROUP_WIDTH * 1 + slid] = totalSum;
+ if (groupId == 3) reduceLM[SIMDGROUP_WIDTH * 2 + slid] = totalSum;
+ barrier(CLK_LOCAL_MEM_FENCE);
+ if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 0 + slid];
+ if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 1 + slid];
+ if (groupId == 0) totalSum += reduceLM[SIMDGROUP_WIDTH * 2 + slid];
+
+ // 1 outputs per fiber in wave 0
+ if (groupId == 0) {
+ dst = (global float*)((global char*)dst + offsetd);
+ dst[gid] = totalSum;
+ }
+}
+++ /dev/null
-// src0_q, src0_d, src1 are transposed as a preprocessing step
-// 4-bit weights are transposed in groups of 4 (unsigned short int)
-// consider weights originally "next to each other", now "on top of each other"
-// each fiber computes a 8x4 tile of output elements
-// using unshuffled weights
-
-#pragma OPENCL EXTENSION cl_khr_fp16 : enable
-#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
-
-#ifdef cl_qcom_reqd_sub_group_size
-#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
-#define ADRENO_GPU 1
-#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
-#endif
-
-#ifdef ADRENO_GPU
-REQD_SUBGROUP_SIZE_128
-#endif
-
-kernel void kernel_mul_mat_Ab_Bi_8x4(
- global const ushort * src0_q, // quantized A
- global const half * src0_d, // A scales
- __read_only image1d_buffer_t src1, // B (1d image)
- global float * dst, // C
- int m, // M
- int n, // N with padding
- int k, // K
- int n_no_padding // N without padding
-) {
-
- int m_4 = m >> 2;
- int n_4 = n >> 2;
-
- int gy = get_global_id(0);
- int gx = get_global_id(1);
- int gx_2 = gx << 2;
-
- half8 c0 = 0, c1 = 0, c2 = 0, c3 = 0; // 8x4 output elements
- half8 B; // registers for activations
- half4 dequantized_weights; // registers for dequantized weights
- __global const ushort* weight_ptr = src0_q + gx_2; // pointer for weights
- __global const half* scale_ptr = src0_d + gx_2; // pointer for scales
-
- for(int i=0; i<k; i+=4){ //loop through K dimension
-
- B.s0123 = read_imageh(src1, gy*2 + (i)*(n_4));
- B.s4567 = read_imageh(src1, gy*2 + (i)*(n_4)+1);
-
- // keep (i/4) and (i/32) in parenthesis, rounds down
- // load 4 consecutive groups of 4 weights
- ushort4 bits4 = vload4(0, weight_ptr + (i/4)*(m)); // (i/4) because weights grouped in 4s
-
- // load 4 consecutive scales
- half4 scale = vload4(0, scale_ptr + (i/32)*(m));// (i/32) because 1 scale per 32 elements
-
- // j=0
- dequantized_weights.s0 = ((bits4.s0 & (0x000F)) - 8) * scale.s0; // dequantize a row of the 16 weights
- dequantized_weights.s1 = ((bits4.s1 & (0x000F)) - 8) * scale.s1;
- dequantized_weights.s2 = ((bits4.s2 & (0x000F)) - 8) * scale.s2;
- dequantized_weights.s3 = ((bits4.s3 & (0x000F)) - 8) * scale.s3;
- c0 += B * dequantized_weights.s0; // vector-scalar multiplication to accumulate
- c1 += B * dequantized_weights.s1;
- c2 += B * dequantized_weights.s2;
- c3 += B * dequantized_weights.s3;
-
- // j=1
- B.s0123 = read_imageh(src1, gy*2 + (i+1)*(n_4));
- B.s4567 = read_imageh(src1, gy*2 + (i+1)*(n_4)+1);
- dequantized_weights.s0 = (((bits4.s0 & (0x00F0)) >> 4) - 8) * scale.s0; // dequantize a row of the 16 weights
- dequantized_weights.s1 = (((bits4.s1 & (0x00F0)) >> 4) - 8) * scale.s1;
- dequantized_weights.s2 = (((bits4.s2 & (0x00F0)) >> 4) - 8) * scale.s2;
- dequantized_weights.s3 = (((bits4.s3 & (0x00F0)) >> 4) - 8) * scale.s3;
- c0 += B * dequantized_weights.s0; //vector-scalar multiplication to accumulate
- c1 += B * dequantized_weights.s1;
- c2 += B * dequantized_weights.s2;
- c3 += B * dequantized_weights.s3;
-
- // j=2
- B.s0123 = read_imageh(src1, gy*2 + (i+2)*(n_4));
- B.s4567 = read_imageh(src1, gy*2 + (i+2)*(n_4)+1);
- dequantized_weights.s0 = (((bits4.s0 & (0x0F00)) >> 8) - 8) * scale.s0; // dequantize a row of the 16 weights
- dequantized_weights.s1 = (((bits4.s1 & (0x0F00)) >> 8) - 8) * scale.s1;
- dequantized_weights.s2 = (((bits4.s2 & (0x0F00)) >> 8) - 8) * scale.s2;
- dequantized_weights.s3 = (((bits4.s3 & (0x0F00)) >> 8) - 8) * scale.s3;
- c0 += B * dequantized_weights.s0; // vector-scalar multiplication to accumulate
- c1 += B * dequantized_weights.s1;
- c2 += B * dequantized_weights.s2;
- c3 += B * dequantized_weights.s3;
-
- // j=3
- B.s0123 = read_imageh(src1, gy*2 + (i+3)*(n_4));
- B.s4567 = read_imageh(src1, gy*2 + (i+3)*(n_4)+1);
- dequantized_weights.s0 = (((bits4.s0 & (0xF000)) >> 12) - 8) * scale.s0; // dequantize a row of the 16 weights
- dequantized_weights.s1 = (((bits4.s1 & (0xF000)) >> 12) - 8) * scale.s1;
- dequantized_weights.s2 = (((bits4.s2 & (0xF000)) >> 12) - 8) * scale.s2;
- dequantized_weights.s3 = (((bits4.s3 & (0xF000)) >> 12) - 8) * scale.s3;
- c0 += B * dequantized_weights.s0; // vector-scalar multiplication to accumulate
- c1 += B * dequantized_weights.s1;
- c2 += B * dequantized_weights.s2;
- c3 += B * dequantized_weights.s3;
- }
-
- int idx = (gy<<3)*m + (gx<<2); // vectorized store 16 elements
-
- // conditional check if store is to a valid location. Required when N is not a multiple of 8
- // if statements allow registers to be reused for each store
- // provides a performance boost due to reduced register footprint, which increases number of concurrent waves
- 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);
- }
-}
+++ /dev/null
-#pragma OPENCL EXTENSION cl_khr_fp16 : enable
-#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
-
-#ifdef cl_qcom_reqd_sub_group_size
-#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
-#define ADRENO_GPU 1
-#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
-#endif
-
-#ifdef ADRENO_GPU
-REQD_SUBGROUP_SIZE_128
-#endif
-
-kernel void kernel_mul_mm_q8_0_f32_8x4(
- global const uint * src0_q,
- global const half * src0_d,
- __read_only image1d_buffer_t src1,
- global float * dst,
- int k,
- int m,
- int n,
- int n_no_padding,
- ulong offsetd
-) {
-
- int m_4 = m >> 2;
- int n_4 = n >> 2;
-
- int gy = get_global_id(0);
- int gx = get_global_id(1);
- int gx_2 = gx << 2;
- dst = (global float *)((global char*)dst + offsetd);
-
-
- half8 c0 = 0, c1 = 0, c2 = 0, c3 = 0;
- half8 B;
- half4 deq;
-
- __global const uint* wptr = src0_q + gx_2;
- __global const half* sptr = src0_d + gx_2;
-
- for (int i = 0; i < k; i += 4) {
- uint4 pack4 = vload4(0, wptr + (i / 4) * m);
- half4 scale = vload4(0, sptr + (i / 32) * m);
-
- char4 p0 = as_char4(pack4.s0);
- char4 p1 = as_char4(pack4.s1);
- char4 p2 = as_char4(pack4.s2);
- char4 p3 = as_char4(pack4.s3);
-
- // ------------------- j = 0 (k = i+0) -------------------
- B.s0123 = read_imageh(src1, gy * 2 + (i + 0) * n_4);
- B.s4567 = read_imageh(src1, gy * 2 + (i + 0) * n_4 + 1);
-
- half4 wj0 = convert_half4((char4)(p0.s0, p1.s0, p2.s0, p3.s0)) * scale;
-
- c0 += B * wj0.s0;
- c1 += B * wj0.s1;
- c2 += B * wj0.s2;
- c3 += B * wj0.s3;
-
- // ------------------- j = 1 (k = i+1) -------------------
- B.s0123 = read_imageh(src1, gy * 2 + (i + 1) * n_4);
- B.s4567 = read_imageh(src1, gy * 2 + (i + 1) * n_4 + 1);
-
- half4 wj1 = convert_half4((char4)(p0.s1, p1.s1, p2.s1, p3.s1)) * scale;
-
- c0 += B * wj1.s0;
- c1 += B * wj1.s1;
- c2 += B * wj1.s2;
- c3 += B * wj1.s3;
-
- // ------------------- j = 2 (k = i+2) -------------------
- B.s0123 = read_imageh(src1, gy * 2 + (i + 2) * n_4);
- B.s4567 = read_imageh(src1, gy * 2 + (i + 2) * n_4 + 1);
-
- half4 wj2 = convert_half4((char4)(p0.s2, p1.s2, p2.s2, p3.s2)) * scale;
-
- c0 += B * wj2.s0;
- c1 += B * wj2.s1;
- c2 += B * wj2.s2;
- c3 += B * wj2.s3;
-
- // ------------------- j = 3 (k = i+3) -------------------
- B.s0123 = read_imageh(src1, gy * 2 + (i + 3) * n_4);
- B.s4567 = read_imageh(src1, gy * 2 + (i + 3) * n_4 + 1);
-
- half4 wj3 = convert_half4((char4)(p0.s3, p1.s3, p2.s3, p3.s3)) * scale;
-
- c0 += B * wj3.s0;
- c1 += B * wj3.s1;
- c2 += B * wj3.s2;
- c3 += B * wj3.s3;
- }
-
- int idx = (gy << 3) * m + (gx << 2);
-
- 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);
- }
-}
overview / pkg / ggml / sources / llama.cpp / commitdiff