Public/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-09-19 09:15:04 +02:00
Code Issues Packages Projects Releases Wiki Activity

opencl: add bin kernel kernel_gemm_noshuffle_q6_k_f32_32b_trans_ila_a8_bin (#28678)

Browse Source 
* opencl: add A8 Q6_K non-MoE binary kernel

* opencl: fix layout compatibility
This commit is contained in:
shaofeiqi
2026-09-18 10:50:15 -07:00
committed by GitHub
parent 4fea119de3
commit ec92815050
3 changed files with 388 additions and 16 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-opencl/CMakeLists.txt
+1
View File
@@ -191,6 +191,7 @@ set(GGML_OPENCL_KERNELS
gemv_noshuffle_q6_k_f32_tiled
gemm_noshuffle_q6_k_f32
gemm_noshuffle_q6_k_f32_tiled
gemv_noshuffle_q6_k_f32_32b_trans
gemv_noshuffle_q5_k_f32
gemm_noshuffle_q5_k_f32
mul
ggml/src/ggml-opencl/ggml-opencl.cpp
+259 -16
View File
@@ -1246,6 +1246,8 @@ struct ggml_backend_opencl_context {
cl_kernel kernel_gemv_noshuffle_q6_K_f32_mc3; // multi-column (N=3) verify GEMV
cl_kernel kernel_gemm_noshuffle_q6_K_f32;
cl_kernel kernel_gemm_noshuffle_q6_K_f32_cok;
cl_kernel kernel_gemm_noshuffle_q6_k_f32_32b_trans_ila_a8_bin;
cl_kernel kernel_gemv_noshuffle_q6_k_f32_32b_trans;
cl_kernel kernel_gemv_noshuffle_q5_k_f32;
cl_kernel kernel_gemv_noshuffle_q5_k_f32_mc3; // multi-column (N=3) verify GEMV (spec/MTP)
cl_kernel kernel_gemm_noshuffle_q5_k_f32;
@@ -4367,6 +4369,43 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
}
}
backend_ctx->kernel_gemv_noshuffle_q6_k_f32_32b_trans = nullptr;
backend_ctx->kernel_gemm_noshuffle_q6_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_q6_k_f32_32b_trans.cl.h"
};
#else
const std::string kernel_src = read_file("gemv_noshuffle_q6_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_q6_k_f32_32b_trans =
clCreateKernel(prog, "kernel_gemv_noshuffle_q6_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_q6_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_q6_k_f32_32b_trans_ila_a8_bin =
clCreateKernel(bin_prog, "kernel_gemm_noshuffle_q6_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";
@@ -7294,6 +7333,8 @@ struct ggml_tensor_extra_cl_q6_K {
cl_mem ql_img = nullptr;
// Upper 2 bits of quantized weights.
cl_mem qh = nullptr;
// Upper 2 bits as image1d_buffer_t
cl_mem qh_img = nullptr;
// Scales for each block.
cl_mem s = nullptr;
// Scales for each super block.
@@ -7329,6 +7370,10 @@ struct ggml_tensor_extra_cl_q6_K {
CL_CHECK(clReleaseMemObject(ql_img));
ql_img = nullptr;
}
if (qh_img != nullptr) {
CL_CHECK(clReleaseMemObject(qh_img));
qh_img = nullptr;
}
size_ql = 0;
size_qh = 0;
@@ -8566,6 +8611,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_q6_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_q6_k_f32_32b_trans ||
!backend_ctx->kernel_gemm_noshuffle_q6_k_f32_32b_trans_ila_a8_bin) {
return false;
}
return (tensor->ne[0] % 256 == 0) && (tensor->ne[1] % 64 == 0) &&
!use_q6k_tiled(backend_ctx, tensor) && !use_flat_gemv_for_large_m_q6_K(backend_ctx, tensor);
#else
GGML_UNUSED(backend_ctx);
GGML_UNUSED(tensor);
return false;
#endif
}
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 ||
@@ -11181,18 +11241,39 @@ static void ggml_backend_opencl_buffer_set_tensor(ggml_backend_buffer_t buffer,
cl_int M = tensor->ne[1]; // ne01
cl_int K = tensor->ne[0]; // ne00
// Transpose ql as ushort
transpose_2d_as_16b(backend_ctx,
extra->ql, extra->ql, size_ql, K/4, M);
if (use_q6_k_bin_kernels(backend_ctx, tensor)) {
GGML_ASSERT(K % 256 == 0);
GGML_ASSERT(M % 64 == 0);
// Transpose qh as uchar
transpose_2d_as_8b(backend_ctx,
extra->qh, extra->qh, size_qh, K/4, M);
transpose_2d_as_32b(backend_ctx, extra->ql, extra->ql, size_ql, K/8, M);
transpose_2d_as_32b(backend_ctx, extra->qh, extra->qh, size_qh, K/16, M);
// Transpose s as ushort
transpose_2d_as_16b(backend_ctx,
extra->s, extra->s, size_s, K/16/2, M);
cl_image_format wimg_fmt = { CL_R, CL_UNSIGNED_INT32 };
cl_image_desc wimg_desc;
memset(&wimg_desc, 0, sizeof(wimg_desc));
wimg_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
wimg_desc.image_width = static_cast<size_t>(ggml_nelements(tensor) / 8);
wimg_desc.buffer = extra->ql;
CL_CHECK((extra->ql_img = clCreateImage(context, CL_MEM_READ_ONLY, &wimg_fmt, &wimg_desc, NULL, &err), err));
memset(&wimg_desc, 0, sizeof(wimg_desc));
wimg_desc.image_type = CL_MEM_OBJECT_IMAGE1D_BUFFER;
wimg_desc.image_width = static_cast<size_t>(ggml_nelements(tensor) / 16);
wimg_desc.buffer = extra->qh;
CL_CHECK((extra->qh_img = clCreateImage(context, CL_MEM_READ_ONLY, &wimg_fmt, &wimg_desc, NULL, &err), err));
} else {
// Transpose ql as ushort
transpose_2d_as_16b(backend_ctx,
extra->ql, extra->ql, size_ql, K/4, M);
// Transpose qh as uchar
transpose_2d_as_8b(backend_ctx,
extra->qh, extra->qh, size_qh, K/4, M);
// Transpose s as ushort
transpose_2d_as_16b(backend_ctx,
extra->s, extra->s, size_s, K/16/2, M);
}
// Transpose d as ushort
transpose_2d_as_16b(backend_ctx,
extra->d, extra->d, size_d, K/256, M);
@@ -12317,15 +12398,24 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
buf_trans_ql.allocate(backend_ctx->context, size_ql);
buf_trans_qh.allocate(backend_ctx->context, size_qh);
buf_trans_s.allocate(backend_ctx->context, size_s);
buf_trans_d.allocate(backend_ctx->context, size_d);
buf_unpacked.allocate(backend_ctx->context, ggml_nbytes(tensor));
// transpose ql, qh, s and d back
transpose_2d_as_16b(backend_ctx, extra->ql, buf_trans_ql.buffer, size_ql, M, K/4);
transpose_2d_as_8b(backend_ctx, extra->qh, buf_trans_qh.buffer, size_qh, M, K/4);
transpose_2d_as_16b(backend_ctx, extra->s, buf_trans_s.buffer, size_s, M, K/16/2);
transpose_2d_as_16b(backend_ctx, extra->d, buf_trans_d.buffer, size_d, M, K/256);
cl_mem s_buffer;
if (use_q6_k_bin_kernels(backend_ctx, tensor)) {
transpose_2d_as_32b(backend_ctx, extra->ql, buf_trans_ql.buffer, size_ql, M, K/8);
transpose_2d_as_32b(backend_ctx, extra->qh, buf_trans_qh.buffer, size_qh, M, K/16);
// s is left row-major, untransposed, for the binary layout.
s_buffer = extra->s;
} else {
// transpose ql, qh, s and d back
buf_trans_s.allocate(backend_ctx->context, size_s);
transpose_2d_as_16b(backend_ctx, extra->ql, buf_trans_ql.buffer, size_ql, M, K/4);
transpose_2d_as_8b(backend_ctx, extra->qh, buf_trans_qh.buffer, size_qh, M, K/4);
transpose_2d_as_16b(backend_ctx, extra->s, buf_trans_s.buffer, size_s, M, K/16/2);
s_buffer = buf_trans_s.buffer;
}
transpose_2d_as_16b(backend_ctx, extra->d, buf_trans_d.buffer, size_d, M, K/256);
// unpack
cl_uchar mask = 0xFF;
@@ -12333,7 +12423,7 @@ static void ggml_backend_opencl_buffer_get_tensor(ggml_backend_buffer_t buffer,
cl_kernel kernel = backend_ctx->kernel_restore_block_q6_K_noshuffle;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &buf_trans_ql.buffer));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &buf_trans_qh.buffer));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &buf_trans_s.buffer));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &s_buffer));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &buf_trans_d.buffer));
CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_mem), &buf_unpacked.buffer));
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(cl_uchar), &mask));
@@ -21111,6 +21201,145 @@ static void ggml_cl_mul_mat_q4_k_f32_adreno(ggml_backend_t backend, const ggml_t
#endif
}
#ifdef GGML_OPENCL_USE_ADRENO_KERNELS
static void ggml_cl_mul_mat_q6_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_q6_K * extra0_q6_K = (ggml_tensor_extra_cl_q6_K *)src0->extra;
ggml_tensor_extra_cl * extra1 = (ggml_tensor_extra_cl *)src1->extra;
ggml_tensor_extra_cl * extrad = (ggml_tensor_extra_cl *)dst->extra;
cl_ulong offset1 = extra1->offset + src1->view_offs;
cl_ulong offsetd = extrad->offset + dst->view_offs;
const int ne00 = src0->ne[0];
const int ne01 = src0->ne[1];
const int ne1 = dst->ne[1];
GGML_ASSERT(ne00 % ggml_blck_size(src0->type) == 0);
cl_context context = backend_ctx->context;
cl_kernel kernel;
cl_int err;
cl_buffer_region region;
cl_image_format img_fmt;
cl_image_desc img_desc;
const int M = ne01;
const int N = ne1;
const int K = ne00;
if (ne1 == 1) {
cl_mem b_sub_buf = nullptr;
cl_mem b_img = nullptr;
region.origin = offset1;
region.size = (size_t)K * N * sizeof(float);
CL_CHECK((b_sub_buf = clCreateSubBuffer(extra1->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, &region, &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_q6_k_f32_32b_trans;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q6_K->ql_img));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q6_K->qh_img));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_q6_K->s));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra0_q6_K->d));
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_img));
CL_CHECK(clReleaseMemObject(b_sub_buf));
} 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, &region, &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, &region, &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_q6_k_f32_32b_trans_ila_a8_bin;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0_q6_K->ql_img));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q6_K->qh));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra0_q6_K->s));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extra0_q6_K->d));
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_q6_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);
@@ -21159,6 +21388,20 @@ static void ggml_cl_mul_mat_q6_K_f32_adreno(ggml_backend_t backend, const ggml_t
// (the #1 MTP bottleneck; mc3 above can't, it reads the noshuffle layout).
const bool use_q6k_tiled_mc = q6k_mc3 && (ne1 == 3) && (ne01 >= 32768) && use_q6k_tiled(backend_ctx, src0);
const bool use_bin = use_q6_k_bin_kernels(backend_ctx, src0);
if (use_bin) {
if (use_q6k_mc3 || use_q6k_tiled_mc) {
static bool warned = false;
if (!warned) {
GGML_LOG_WARN("ggml_opencl: GGML_OPENCL_Q6K_MC3 is bypassed by Q6_K binary kernels\n");
warned = true;
}
}
ggml_cl_mul_mat_q6_K_f32_adreno_ila(backend, src0, src1, dst);
return;
}
if (ne1 == 1 || use_q6k_mc3 || use_q6k_tiled_mc) {
cl_mem ql_img = nullptr;
cl_mem qh_img = nullptr;
ggml/src/ggml-opencl/kernels/gemv_noshuffle_q6_k_f32_32b_trans.cl
+128
View File
@@ -0,0 +1,128 @@
#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 N_SIMDGROUP 8
#define SIMDGROUP_WIDTH 64
static inline float8 q6_k_to_fp32_packed8(ushort2 ql8, ushort qh8, float d_scale) {
float8 fp32x8;
fp32x8.s0 = ((float)(( ql8.s0 & 0x000F) | ((uint)((qh8 ) & 0x3) << 4)) - 32.f) * d_scale;
fp32x8.s1 = ((float)((( ql8.s0 >> 4) & 0x000F) | ((uint)((qh8 >> 2) & 0x3) << 4)) - 32.f) * d_scale;
fp32x8.s2 = ((float)((( ql8.s0 >> 8) & 0x000F) | ((uint)((qh8 >> 4) & 0x3) << 4)) - 32.f) * d_scale;
fp32x8.s3 = ((float)((( ql8.s0 >> 12)& 0x000F) | ((uint)((qh8 >> 6) & 0x3) << 4)) - 32.f) * d_scale;
fp32x8.s4 = ((float)(( ql8.s1 & 0x000F) | ((uint)((qh8 >> 8) & 0x3) << 4)) - 32.f) * d_scale;
fp32x8.s5 = ((float)((( ql8.s1 >> 4) & 0x000F) | ((uint)((qh8 >>10) & 0x3) << 4)) - 32.f) * d_scale;
fp32x8.s6 = ((float)((( ql8.s1 >> 8) & 0x000F) | ((uint)((qh8 >>12) & 0x3) << 4)) - 32.f) * d_scale;
fp32x8.s7 = ((float)((( ql8.s1 >> 12)& 0x000F) | ((uint)((qh8 >>14) & 0x3) << 4)) - 32.f) * d_scale;
return fp32x8;
}
__attribute__((qcom_reqd_sub_group_size("half")))
__kernel void kernel_gemv_noshuffle_q6_k_f32_32b_trans(
__read_only image1d_buffer_t src0_ql,
__read_only image1d_buffer_t src0_qh,
__global char * src0_s,
__global half * src0_d,
__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_superblocks = ne00 / QK_K;
int num_subblocks = ne00 / 32; // 2 sub-blocks of 16 processed per iter below
int scales_per_row = num_superblocks * 16;
__private float sum = 0.0f;
// Loop over 32-element groups (2 sub-blocks of 16 each), N_SIMDGROUP groups per iter.
for (uint ib = sgid; ib < num_subblocks; ib += N_SIMDGROUP) {
uint sb = ib / 8; // super-block index
uint j = ib % 8; // 32-element group within super-block (0..7)
// Load d for this super-block.
half d_val = src0_d[sb * ne01 + i01];
// Load 2 sub-block scales (int8), one per 16 elements.
global const char * sc = src0_s + i01 * scales_per_row + sb * 16;
float scale0 = (float)d_val * (float)sc[j * 2];
float scale1 = (float)d_val * (float)sc[j * 2 + 1];
// Load 4 uints of ql (32 elements, 4-bit each = 128 bits), column-major stride ne01.
uint ql_base = (ib * 4) * ne01 + i01;
uint4 regQL;
regQL.s0 = read_imageui(src0_ql, ql_base).x;
regQL.s1 = read_imageui(src0_ql, ql_base + ne01).x;
regQL.s2 = read_imageui(src0_ql, ql_base + ne01 * 2).x;
regQL.s3 = read_imageui(src0_ql, ql_base + ne01 * 3).x;
// Load 2 uints of qh (32 elements, 2-bit each = 64 bits), column-major stride ne01.
uint qh_base = (ib * 2) * ne01 + i01;
uint2 regQH;
regQH.s0 = read_imageui(src0_qh, qh_base).x;
regQH.s1 = read_imageui(src0_qh, qh_base + ne01).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 y4v = 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);
// Dequantize elements 0..7 (scale0).
float8 fp32x8 = q6_k_to_fp32_packed8(as_ushort2(regQL.s0), (ushort)(regQH.s0 & 0xFFFF), scale0);
float4 acc = y0 * fp32x8.lo;
acc += y1 * fp32x8.hi;
// Dequantize elements 8..15 (scale0).
fp32x8 = q6_k_to_fp32_packed8(as_ushort2(regQL.s1), (ushort)(regQH.s0 >> 16), scale0);
acc += y2 * fp32x8.lo;
acc += y3 * fp32x8.hi;
// Dequantize elements 16..23 (scale1).
fp32x8 = q6_k_to_fp32_packed8(as_ushort2(regQL.s2), (ushort)(regQH.s1 & 0xFFFF), scale1);
acc += y4v * fp32x8.lo;
acc += y5 * fp32x8.hi;
// Dequantize elements 24..31 (scale1).
fp32x8 = q6_k_to_fp32_packed8(as_ushort2(regQL.s3), (ushort)(regQH.s1 >> 16), scale1);
acc += y6 * fp32x8.lo;
acc += y7 * fp32x8.hi;
sum += ((acc.s0 + acc.s1) + (acc.s2 + acc.s3));
}
// reduction in local memory, assumes #subgroups=4
__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;
}
}