Commit 07fc97716 for llama.cpp
commit 07fc97716fa3dab457797fcfe51705034ce8957a
Author: shaofeiqi <shaoqi@qti.qualcomm.com>
Date: Fri Sep 11 22:10:08 2026 -0700
opencl: add bin kernel `kernel_gemm_noshuffle_q4_k_f32_32b_trans_ila_a8_bin` (#28677)
* opencl: add A8 Q4_K non-MoE binary kernel
* opencl: fix layout compatibility
* opencl: rename binary kernel selection helpers
---------
Co-authored-by: Li He <lih@qti.qualcomm.com>
diff --git a/ggml/src/ggml-opencl/CMakeLists.txt b/ggml/src/ggml-opencl/CMakeLists.txt
index 716577bb7..45a7075b2 100644
--- a/ggml/src/ggml-opencl/CMakeLists.txt
+++ b/ggml/src/ggml-opencl/CMakeLists.txt
@@ -185,6 +185,7 @@ set(GGML_OPENCL_KERNELS
gemv_noshuffle_q4_k_f32_o4
gemv_noshuffle_q4_k_f32_tiled
gemm_noshuffle_q4_k_f32
+ gemv_noshuffle_q4_k_f32_32b_trans
gemv_noshuffle_q6_k_f32
gemv_noshuffle_q6_k_f32_o4
gemv_noshuffle_q6_k_f32_tiled
diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp
index 231be2cf3..d6820b37d 100644
--- a/ggml/src/ggml-opencl/ggml-opencl.cpp
+++ b/ggml/src/ggml-opencl/ggml-opencl.cpp
@@ -1185,6 +1185,8 @@ struct ggml_backend_opencl_context {
cl_kernel kernel_convert_block_q4_k_tiled_ns; // tiled-wide convert (opt-in)
cl_kernel kernel_gemv_noshuffle_q4_k_f32_mc3; // multi-column (N=3) verify GEMV
cl_kernel kernel_gemm_noshuffle_q4_k_f32;
+ cl_kernel kernel_gemm_noshuffle_q4_k_f32_32b_trans_ila_a8_bin;
+ cl_kernel kernel_gemv_noshuffle_q4_k_f32_32b_trans;
cl_kernel kernel_gemm_noshuffle_q4_k_q8_1_dp4a = nullptr; // dp4a (int8) dense prefill GEMM
cl_kernel kernel_gemm_noshuffle_q4_k_q8_1_dp4a_wimg = nullptr; // dp4a dense prefill GEMM, weights via texture (X1 opt-in)
cl_kernel kernel_gemm_noshuffle_q5_k_q8_1_dp4a = nullptr; // dp4a (int8) dense q5_K prefill GEMM
@@ -4260,6 +4262,43 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
GGML_LOG_CONT(".");
}
+ backend_ctx->kernel_gemv_noshuffle_q4_k_f32_32b_trans = nullptr;
+ backend_ctx->kernel_gemm_noshuffle_q4_k_f32_32b_trans_ila_a8_bin = nullptr;
+ if (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E) {
+ {
+ std::string opts = std::string("-cl-std=") + opencl_c_std +
+ " -cl-mad-enable "
+ " -DSIMDGROUP_WIDTH=" +
+ std::to_string(backend_ctx->adreno_wave_size);
+#ifdef GGML_OPENCL_EMBED_KERNELS
+ const std::string kernel_src {
+ #include "gemv_noshuffle_q4_k_f32_32b_trans.cl.h"
+ };
+#else
+ const std::string kernel_src = read_file("gemv_noshuffle_q4_k_f32_32b_trans.cl");
+#endif
+ cl_program prog = build_program_from_source(backend_ctx, kernel_src.c_str(), opts);
+ CL_CHECK((backend_ctx->kernel_gemv_noshuffle_q4_k_f32_32b_trans =
+ clCreateKernel(prog, "gemv_noshuffle_q4_k_f32_32b_trans", &err), err));
+ CL_CHECK(clReleaseProgram(prog));
+ GGML_LOG_CONT(".");
+ }
+
+ if (use_adreno_bin_kernels(backend_ctx)) {
+ size_t bin_size = 0;
+ const char * kernel_bin = (const char *)backend_ctx->get_adreno_bin_kernel("gemm_noshuffle_q4_k_f32_32b_trans_ila_a8", &bin_size);
+ if (kernel_bin && bin_size > 0) {
+ cl_program bin_prog =
+ build_program_from_binary(backend_ctx->context, backend_ctx->device, kernel_bin, "", bin_size);
+
+ CL_CHECK((backend_ctx->kernel_gemm_noshuffle_q4_k_f32_32b_trans_ila_a8_bin =
+ clCreateKernel(bin_prog, "kernel_gemm_noshuffle_q4_k_f32_32b_trans_ila_a8", &err), err));
+ CL_CHECK(clReleaseProgram(bin_prog));
+ GGML_LOG_CONT(".");
+ }
+ }
+ }
+
std::string CL_moe_compile_opts = std::string("-cl-std=") + opencl_c_std +
" -cl-mad-enable "
" -cl-fast-relaxed-math";
@@ -7722,6 +7761,7 @@ static void ggml_cl_moe_combine_fused(ggml_backend_t backend, const ggml_tensor
}
inline bool use_q4k_tiled(const ggml_backend_opencl_context *backend_ctx, const ggml_tensor *tensor); // defined below (used by the GLU-subgraph fuse check)
+inline bool use_q4_k_bin_kernels(const ggml_backend_opencl_context *backend_ctx, const ggml_tensor *tensor);
inline bool use_adreno_kernels(const ggml_backend_opencl_context *backend_ctx, const ggml_tensor *tensor); // defined below
static bool ggml_opencl_can_fuse(const ggml_backend_opencl_context * backend_ctx, const struct ggml_cgraph * cgraph, int node_idx, std::initializer_list<enum ggml_op> ops) {
@@ -7776,6 +7816,10 @@ static bool ggml_opencl_can_fuse(const ggml_backend_opencl_context * backend_ctx
if (use_q4k_tiled(backend_ctx, gate->src[0]) || use_q4k_tiled(backend_ctx, up->src[0])) {
return false;
}
+ // q4_K bin kernel requires 32b transposed layout, not compatible with the fused gemv
+ if (use_q4_k_bin_kernels(backend_ctx, gate->src[0]) || use_q4_k_bin_kernels(backend_ctx, up->src[0])) {
+ return false;
+ }
// that noshuffle layout is only produced at set_tensor time when
// use_adreno_kernels() accepts the weight (ne0 >= 512 && ne1 >= 512).
// Smaller weights stay in the plain q4_K layout, which this kernel would
@@ -8349,7 +8393,7 @@ inline bool enable_adreno_trans_weight_q5_K(const ggml_backend_opencl_context *b
qh_img_width <= backend_ctx->image_max_buffer_size;
}
-inline bool use_q4_0_ila_kernels(const ggml_backend_opencl_context *backend_ctx, const ggml_tensor *tensor) {
+inline bool use_q4_0_bin_kernels(const ggml_backend_opencl_context *backend_ctx, const ggml_tensor *tensor) {
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
if (!backend_ctx->kernel_gemv_noshuffle_q4_0_f32_32b_trans ||
!backend_ctx->kernel_gemm_noshuffle_q4_0_f32_32b_trans_ila_a8_bin) {
@@ -8442,6 +8486,21 @@ static inline bool use_flat_gemv_for_large_m_q6_K(const ggml_backend_opencl_cont
&& tensor->ne[2] == 1 && tensor->ne[3] == 1;
}
+inline bool use_q4_k_bin_kernels(const ggml_backend_opencl_context *backend_ctx, const ggml_tensor *tensor) {
+#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
+ if (!backend_ctx->kernel_gemv_noshuffle_q4_k_f32_32b_trans ||
+ !backend_ctx->kernel_gemm_noshuffle_q4_k_f32_32b_trans_ila_a8_bin) {
+ return false;
+ }
+ return (tensor->ne[0] % 256 == 0) && (tensor->ne[1] % 64 == 0) &&
+ !use_q4k_tiled(backend_ctx, tensor) && !use_flat_gemv_for_large_m_q4_K(backend_ctx, tensor);
+#else
+ GGML_UNUSED(backend_ctx);
+ GGML_UNUSED(tensor);
+ return false;
+#endif
+}
+
static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_tensor * op) {
ggml_backend_opencl_device_context * dev_ctx = (ggml_backend_opencl_device_context *)dev->context;
ggml_backend_opencl_context * backend_ctx = dev_ctx->backend_ctx;
@@ -9625,7 +9684,7 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
GGML_ASSERT(K % 32 == 0);
- if (use_q4_0_ila_kernels(backend_ctx, tensor)) {
+ if (use_q4_0_bin_kernels(backend_ctx, tensor)) {
cl_int err;
cl_image_format wimg_fmt;
cl_image_desc wimg_desc;
@@ -10576,8 +10635,25 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
GGML_ASSERT(K % 32 == 0);
- // Transpose q, d, dm as ushort
- transpose_2d_as_16b(backend_ctx, extra->q, extra->q, size_q, K/4, M);
+ if (use_q4_k_bin_kernels(backend_ctx, tensor)) {
+ cl_int err;
+ cl_image_format wimg_fmt;
+ cl_image_desc wimg_desc;
+
+ // transpose quants as 32-bit words (M-first)
+ GGML_ASSERT(M % 64 == 0);
+ transpose_2d_as_32b(backend_ctx, extra->q, extra->q, size_q, K/8, M);
+
+ wimg_fmt = { CL_R, CL_UNSIGNED_INT32 };
+ memset(&wimg_desc, 0, sizeof(wimg_desc));
+ wimg_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+ wimg_desc.image_width = (size_t)M * K / 8;
+ wimg_desc.buffer = extra->q;
+ CL_CHECK((extra->q_img = clCreateImage(context, CL_MEM_READ_ONLY, &wimg_fmt, &wimg_desc, NULL, &err), err));
+ } else {
+ // Transpose q as ushort
+ transpose_2d_as_16b(backend_ctx, extra->q, extra->q, size_q, K/4, M);
+ }
transpose_2d_as_16b(backend_ctx, extra->d, extra->d, size_d, K/256, M);
transpose_2d_as_16b(backend_ctx, extra->dm, extra->dm, size_dm, K/256, M);
@@ -11180,7 +11256,7 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
buf_trans_d.allocate(backend_ctx->context, size_d);
buf_unpacked.allocate(backend_ctx->context, ggml_nbytes(tensor));
- if (use_q4_0_ila_kernels(backend_ctx, tensor)) {
+ if (use_q4_0_bin_kernels(backend_ctx, tensor)) {
transpose_2d_as_32b(backend_ctx, extra->q, buf_trans_q.buffer, size_q, M, K / 8);
} else {
transpose_2d_as_16b(backend_ctx, extra->q, buf_trans_q.buffer, size_q, M, K / 4);
@@ -11855,7 +11931,11 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
buf_trans_s.allocate(backend_ctx->context, size_s);
// Transpose q, d, dm, s back
- transpose_2d_as_16b(backend_ctx, extra->q, buf_trans_q.buffer, size_q, M, K/4);
+ if (use_q4_k_bin_kernels(backend_ctx, tensor)) {
+ transpose_2d_as_32b(backend_ctx, extra->q, buf_trans_q.buffer, size_q, M, K/8);
+ } else {
+ 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/256);
transpose_2d_as_16b(backend_ctx, extra->dm, buf_trans_dm.buffer, size_dm, M, K/256);
transpose_2d_as_8b (backend_ctx, extra->s, buf_trans_s.buffer, size_s, M, K/256*12, true, true);
@@ -18639,9 +18719,9 @@ static void ggml_cl_mul_mat_q4_0_f32_adreno(ggml_backend_t backend, const ggml_t
static const bool q40_mc3 = (getenv("GGML_OPENCL_Q40_MC3") != nullptr);
const bool use_q40_mc3 = q40_mc3 && (ne1 >= 2 && ne1 <= 4) && (ne01 < 32768);
- const bool use_ila = use_q4_0_ila_kernels(backend_ctx, src0);
+ const bool use_bin = use_q4_0_bin_kernels(backend_ctx, src0);
- if (use_ila) {
+ if (use_bin) {
if (use_q40_mc3) {
static bool warned = false;
if (!warned) {
@@ -20196,6 +20276,145 @@ static void ggml_cl_mul_mat_q8_0_f32_adreno(ggml_backend_t backend, const ggml_t
#endif
}
+#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
+static void ggml_cl_mul_mat_q4_k_f32_adreno_ila(ggml_backend_t backend, const ggml_tensor * src0,
+ const ggml_tensor * src1, ggml_tensor * dst) {
+ GGML_ASSERT(src0);
+ GGML_ASSERT(src0->extra);
+ GGML_ASSERT(src1);
+ GGML_ASSERT(src1->extra);
+ GGML_ASSERT(dst);
+ GGML_ASSERT(dst->extra);
+
+ ggml_backend_opencl_context *backend_ctx = (ggml_backend_opencl_context *)backend->context;
+
+ ggml_tensor_extra_cl * extra1 = (ggml_tensor_extra_cl *)src1->extra;
+ ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *)dst->extra;
+ ggml_tensor_extra_cl_q4_K * extra0_q4_k = (ggml_tensor_extra_cl_q4_K *)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 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 b_sub_buf = nullptr;
+ cl_mem b_img = nullptr;
+
+ region.origin = offset1;
+ region.size = (size_t)K * N * sizeof(float);
+ CL_CHECK((b_sub_buf = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err));
+
+ img_fmt = { CL_RGBA, CL_FLOAT };
+ memset(&img_desc, 0, sizeof(img_desc));
+ img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+ img_desc.image_width = (size_t)K * N / 4;
+ img_desc.buffer = b_sub_buf;
+ CL_CHECK((b_img = clCreateImage(context, CL_MEM_READ_ONLY, &img_fmt, &img_desc, NULL, &err), err));
+
+ kernel = backend_ctx->kernel_gemv_noshuffle_q4_k_f32_32b_trans;
+ CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q4_k->q_img));
+ CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q4_k->d));
+ CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_q4_k->dm));
+ CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra0_q4_k->s));
+ CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &b_img));
+ CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), &extrad->data_device));
+ CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_ulong), &offsetd));
+ CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_int), &ne00));
+ CL_CHECK(clSetKernelArg(kernel, 8, sizeof(cl_int), &ne01));
+
+ size_t local_work_size[3] = { 64, 8, 1 };
+ size_t global_work_size[3] = { (size_t)ne01, 8, 1 };
+ backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
+
+ CL_CHECK(clReleaseMemObject(b_sub_buf));
+ CL_CHECK(clReleaseMemObject(b_img));
+ } else {
+ const int gemm_tile_n = 64;
+ int N_pad = CEIL_DIV(N, gemm_tile_n) * gemm_tile_n;
+
+ cl_mem b_sub_buf = nullptr;
+ cl_mem b_padded = nullptr;
+ cl_mem b_buf = nullptr;
+ if (N_pad == N) {
+ region.origin = offset1;
+ region.size = (size_t)K * N * sizeof(float);
+ CL_CHECK((b_sub_buf = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err));
+ b_buf = b_sub_buf;
+ } else {
+ CL_CHECK((b_padded = clCreateBuffer(context, CL_MEM_READ_WRITE, (size_t)K * N_pad * sizeof(float), NULL, &err), err));
+ const float zero = 0.0f;
+ CL_CHECK(clEnqueueFillBuffer(backend_ctx->queue, b_padded, &zero, sizeof(zero), 0, (size_t)K * N_pad * sizeof(float), 0, NULL, NULL));
+ CL_CHECK(clEnqueueCopyBuffer(backend_ctx->queue, extra1->data_device, b_padded, offset1, 0, (size_t)K * N * sizeof(float), 0, NULL, NULL));
+ b_buf = b_padded;
+ }
+
+ img_fmt = { CL_R, CL_FLOAT };
+ memset(&img_desc, 0, sizeof(img_desc));
+ img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+ img_desc.image_width = (size_t)K * N_pad;
+ img_desc.buffer = b_buf;
+ cl_mem b_img;
+ CL_CHECK((b_img = clCreateImage(context, CL_MEM_READ_ONLY, &img_fmt, &img_desc, NULL, &err), err));
+
+ region.origin = offsetd;
+ region.size = (size_t)M * N * sizeof(float);
+ cl_mem d_sub_buf;
+ CL_CHECK((d_sub_buf = clCreateSubBuffer(extrad->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err), err));
+ img_fmt = { CL_R, CL_FLOAT };
+ memset(&img_desc, 0, sizeof(img_desc));
+ img_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
+ img_desc.image_width = (size_t)M * N;
+ img_desc.buffer = d_sub_buf;
+ cl_mem d_img;
+ CL_CHECK((d_img = clCreateImage(context, CL_MEM_WRITE_ONLY, &img_fmt, &img_desc, NULL, &err), err));
+
+ kernel = backend_ctx->kernel_gemm_noshuffle_q4_k_f32_32b_trans_ila_a8_bin;
+ CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q4_k->q_img));
+ CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q4_k->d));
+ CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_q4_k->dm));
+ CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra0_q4_k->s));
+ CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &b_img));
+ CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_mem), &d_img));
+ CL_CHECK(clSetKernelArg(kernel, 6, sizeof(cl_uint), &ne00));
+ CL_CHECK(clSetKernelArg(kernel, 7, sizeof(cl_uint), &ne01));
+ CL_CHECK(clSetKernelArg(kernel, 8, sizeof(int), &N));
+
+ size_t local_work_size[3] = { 64, 2, 2 };
+ size_t m_tiles = (size_t)CEIL_DIV(M, 64);
+ size_t global_work_size[3] = { 64, m_tiles, (size_t)CEIL_DIV(N_pad, gemm_tile_n) };
+ backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
+
+ CL_CHECK(clReleaseMemObject(b_img));
+ if (b_sub_buf) {
+ CL_CHECK(clReleaseMemObject(b_sub_buf));
+ }
+ if (b_padded) {
+ CL_CHECK(clReleaseMemObject(b_padded));
+ }
+ CL_CHECK(clReleaseMemObject(d_img));
+ CL_CHECK(clReleaseMemObject(d_sub_buf));
+ }
+}
+#endif // GGML_OPENCL_USE_ADRENO_KERNELS
+
static void ggml_cl_mul_mat_q4_k_f32_adreno(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
GGML_ASSERT(src0);
@@ -20248,6 +20467,20 @@ static void ggml_cl_mul_mat_q4_k_f32_adreno(ggml_backend_t backend, const ggml_t
// unified routes batched Q6_K lm_head to CPU). Per-layer mc3 is byte-identical.
const bool use_mc3 = q4k_mc3 && (ne1 == 3) && (ne01 < 32768);
+ const bool use_bin = use_q4_k_bin_kernels(backend_ctx, src0);
+
+ if (use_bin) {
+ if (use_mc3) {
+ static bool warned = false;
+ if (!warned) {
+ GGML_LOG_WARN("ggml_opencl: GGML_OPENCL_Q4K_MC3 is bypassed by Q4_K binary kernels\n");
+ warned = true;
+ }
+ }
+ ggml_cl_mul_mat_q4_k_f32_adreno_ila(backend, src0, src1, dst);
+ return;
+ }
+
if (ne1 == 1 || use_mc3) {
cl_mem q_img = nullptr;
cl_mem b_sub_buf = nullptr;
diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_k_f32_32b_trans.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_k_f32_32b_trans.cl
new file mode 100644
index 000000000..2dbd943fd
--- /dev/null
+++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_k_f32_32b_trans.cl
@@ -0,0 +1,134 @@
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+#pragma OPENCL EXTENSION cl_khr_subgroups : enable
+#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
+
+#define QK_K 256
+#define K_SCALE_SIZE 12
+#define N_SIMDGROUP 8
+#define SIMDGROUP_WIDTH 64
+
+inline void get_scale_min_k4(
+ int j,
+ global const uchar * q,
+ uint stride,
+ uchar * d,
+ uchar * m
+) {
+ if (j < 4) {
+ *d = q[j*stride] & 63;
+ *m = q[(j+4)*stride] & 63;
+ } else {
+ *d = (q[(j+4)*stride] & 0x0F) | ((q[(j-4)*stride] & 0xC0) >> 2);
+ *m = ((q[(j+4)*stride] >> 4) & 0x0F) | ((q[j*stride] & 0xC0) >> 2);
+ }
+}
+
+static inline float8 q4_k_to_fp32_packed8(ushort2 q4x8, float scale, float minv) {
+ float8 fp32x8;
+ fp32x8.s0 = (q4x8.s0 & 0x000F) * scale - minv;
+ fp32x8.s1 = ((q4x8.s0 & 0x00F0) >> 4) * scale - minv;
+ fp32x8.s2 = ((q4x8.s0 & 0x0F00) >> 8) * scale - minv;
+ fp32x8.s3 = ((q4x8.s0 & 0xF000) >> 12) * scale - minv;
+ fp32x8.s4 = (q4x8.s1 & 0x000F) * scale - minv;
+ fp32x8.s5 = ((q4x8.s1 & 0x00F0) >> 4) * scale - minv;
+ fp32x8.s6 = ((q4x8.s1 & 0x0F00) >> 8) * scale - minv;
+ fp32x8.s7 = ((q4x8.s1 & 0xF000) >> 12) * scale - minv;
+ return fp32x8;
+}
+
+__attribute__((qcom_reqd_sub_group_size("half")))
+__kernel void gemv_noshuffle_q4_k_f32_32b_trans(
+ read_only image1d_buffer_t src0_q,
+ __global half * src0_d,
+ __global half * src0_dm,
+ __global uchar * src0_s,
+ __read_only image1d_buffer_t src1,
+ __global float * dst,
+ ulong offsetd,
+ int ne00,
+ int ne01
+) {
+ uint i01 = get_global_id(0);
+ uint sgid = get_local_id(1);
+ uint slid = get_sub_group_local_id();
+
+ int num_subblocks = ne00 / 32;
+
+ __private float sum = 0.0f;
+
+ // Loop over sub-blocks of 32 elements, N_SIMDGROUP sub-blocks per iter
+ for (uint ib = sgid; ib < num_subblocks; ib += N_SIMDGROUP) {
+ uint sb = ib / 8;
+ uint j = ib % 8;
+
+ // Load d and dmin for this super-block
+ half d_val = src0_d[sb * ne01 + i01];
+ half dm_val = src0_dm[sb * ne01 + i01];
+
+ // Load sub-block scale and min. s is transposed [nb][12][M]; stride ne01 per code.
+ global const uchar * sc = src0_s + sb * K_SCALE_SIZE * ne01 + i01;
+ uchar sv, mn;
+ get_scale_min_k4(j, sc, ne01, &sv, &mn);
+
+ float scale = (float)d_val * (float)sv;
+ float minv = (float)dm_val * (float)mn;
+
+ // Load 4 uints of quants (32 nibbles = 32 elements), column-major stride ne01
+ uint q_base = ib * ne01 * 4 + i01;
+
+ uint4 regQ;
+ regQ.s0 = read_imageui(src0_q, q_base).x;
+ regQ.s1 = read_imageui(src0_q, q_base + ne01).x;
+ regQ.s2 = read_imageui(src0_q, q_base + ne01 * 2).x;
+ regQ.s3 = read_imageui(src0_q, q_base + ne01 * 3).x;
+
+ // Load activations: 32 floats = 8 float4s
+ uint y_offset = ib * 8;
+
+ float4 y_local = (slid < 8) ? read_imagef(src1, (y_offset + slid)) : (float4)0.0f;
+ float4 y0 = sub_group_broadcast(y_local, 0);
+ float4 y1 = sub_group_broadcast(y_local, 1);
+ float4 y2 = sub_group_broadcast(y_local, 2);
+ float4 y3 = sub_group_broadcast(y_local, 3);
+ float4 y4 = sub_group_broadcast(y_local, 4);
+ float4 y5 = sub_group_broadcast(y_local, 5);
+ float4 y6 = sub_group_broadcast(y_local, 6);
+ float4 y7 = sub_group_broadcast(y_local, 7);
+
+ float8 fp32x8 = q4_k_to_fp32_packed8(as_ushort2(regQ.s0), scale, minv);
+ float4 acc = y0 * fp32x8.lo;
+ acc += y1 * fp32x8.hi;
+
+ fp32x8 = q4_k_to_fp32_packed8(as_ushort2(regQ.s1), scale, minv);
+ acc += y2 * fp32x8.lo;
+ acc += y3 * fp32x8.hi;
+
+ fp32x8 = q4_k_to_fp32_packed8(as_ushort2(regQ.s2), scale, minv);
+ acc += y4 * fp32x8.lo;
+ acc += y5 * fp32x8.hi;
+
+ fp32x8 = q4_k_to_fp32_packed8(as_ushort2(regQ.s3), scale, minv);
+ acc += y6 * fp32x8.lo;
+ acc += y7 * fp32x8.hi;
+
+ sum += ((acc.s0 + acc.s1) + (acc.s2 + acc.s3));
+ }
+
+ // reduction in local memory over N_SIMDGROUP subgroups
+ __local float reduceLM[SIMDGROUP_WIDTH * (N_SIMDGROUP - 1)];
+ if (sgid > 0) {
+ reduceLM[SIMDGROUP_WIDTH * (sgid - 1) + slid] = sum;
+ }
+ barrier(CLK_LOCAL_MEM_FENCE);
+ if (sgid == 0) {
+ for (uint i = 0; i < N_SIMDGROUP - 1; ++i) {
+ sum += reduceLM[SIMDGROUP_WIDTH * i + slid];
+ }
+ }
+
+ // 1 output per thread in subgroup 0
+ if (sgid == 0) {
+ dst = dst + (offsetd >> 2);
+ dst[i01] = sum;
+ }
+}