Commit c19043f67 for llama.cpp
commit c19043f67e52f1f1e00d797246e19643bb2f168e
Author: shaofeiqi <shaoqi@qti.qualcomm.com>
Date: Sat Nov 7 22:04:10 2026 -0800
opencl: improve q8_0 and MoE small-n verification for spec (#31831)
* opencl: add kernel_gemv_noshuffle_q8_0_f32_mc
Assisted-by: Claude Opus 5.5
Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com>
Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com>
* opencl: small-n MoE GEMV for q5_1, q4_k
Assisted-by: Claude Opus 5.5
Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com>
Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com>
* opencl: adreno kernel k floor 256 for x2
Assisted-by: Claude Opus 5.5
Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com>
Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com>
* opencl: small-n MoE GEMV for more quants
Assisted-by: Claude Opus 5.5
Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com>
Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com>
* opencl: rename function
* opencl: disable bin FA for small verify batch
---------
Co-authored-by: Li He <lih@qti.qualcomm.com>
Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com>
diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp
index 1a4c03519..771fe7751 100644
--- a/ggml/src/ggml-opencl/ggml-opencl.cpp
+++ b/ggml/src/ggml-opencl/ggml-opencl.cpp
@@ -1387,6 +1387,7 @@ struct ggml_backend_opencl_context {
cl_kernel kernel_gemm_noshuffle_q8_0_q8_1_dp4a_ila_a8_bin = nullptr;
cl_kernel kernel_gemv_noshuffle_q8_0_f32;
cl_kernel kernel_gemv_noshuffle_q8_0_f32_splitk; // split-K across WGs (small-M decode)
+ cl_kernel kernel_gemv_noshuffle_q8_0_f32_mc; // multi-column (N 2..4) verify GEMV
cl_kernel kernel_gemm_noshuffle_q1_0_f32;
cl_kernel kernel_gemv_noshuffle_q1_0_f32;
cl_kernel kernel_gemv_noshuffle_q1_0_f32_nsg1 = nullptr;
@@ -6334,6 +6335,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
CL_CHECK((backend_ctx->kernel_gemv_noshuffle_q8_0_f32 = clCreateKernel(prog, "kernel_gemv_noshuffle_q8_0_f32", &err), err));
CL_CHECK((backend_ctx->kernel_gemv_noshuffle_q8_0_f32_splitk = clCreateKernel(prog, "kernel_gemv_noshuffle_q8_0_f32_splitk", &err), err));
+ CL_CHECK((backend_ctx->kernel_gemv_noshuffle_q8_0_f32_mc = clCreateKernel(prog, "kernel_gemv_noshuffle_q8_0_f32_mc", &err), err));
CL_CHECK(clReleaseProgram(prog));
GGML_LOG_CONT(".");
}
@@ -11117,6 +11119,13 @@ inline bool use_adreno_kernels(const ggml_backend_opencl_context *backend_ctx, c
threshold_ne0 = 128;
threshold_ne1 = 128;
}
+ // K floor 256 on X2E: the gemma-4 E2B/E4B MTP drafters' weights are all K = 256.
+ // Default on for X2E, GGML_OPENCL_K_FLOOR_256=0/1 forces it.
+ static const char * k_floor_256_env = getenv("GGML_OPENCL_K_FLOOR_256");
+ const bool k_floor_256 = k_floor_256_env ? atoi(k_floor_256_env) != 0 : backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E;
+ if (k_floor_256 && threshold_ne0 > 256) {
+ threshold_ne0 = 256;
+ }
bool threashold_ok = tensor->ne[0] >= threshold_ne0 && tensor->ne[1] >= threshold_ne1 &&
tensor->ne[2] == 1 && tensor->ne[3] == 1;
@@ -11497,6 +11506,13 @@ static bool use_fa_bin_kernels_prefill(const ggml_backend_opencl_context * backe
const int dk = q->ne[0];
const int dv = v->ne[0];
+ // do not choose bin FA kernels to avoid repacking the KV cache for small
+ // verify batches on X2
+ if (backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E && is_mixed &&
+ (dk == 256 || dk == 512) && n_q <= 4) {
+ return false;
+ }
+
constexpr bool prefill_only = true;
return (backend_ctx->gpu_family == GPU_FAMILY::ADRENO &&
@@ -26194,7 +26210,13 @@ static void ggml_cl_mul_mat_q8_0_f32_adreno(ggml_backend_t backend, const ggml_t
int N = ne1;
int K = ne00;
- if (ne1 == 1) {
+ // Multi-column verify GEMV for ne1 2..4, reuses the ne1 == 1 image setup.
+ // Default on for X2E, GGML_OPENCL_Q8_0_MC=0/1 forces it.
+ static const char * q8_mc_env = getenv("GGML_OPENCL_Q8_0_MC");
+ const bool q8_mc = q8_mc_env ? atoi(q8_mc_env) != 0 : backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E;
+ const bool use_q8_mc = q8_mc && ne1 >= 2 && ne1 <= 4;
+
+ if (ne1 == 1 || use_q8_mc) {
cl_mem q_img = nullptr;
cl_mem b_sub_buf = nullptr;
cl_mem b_img = nullptr;
@@ -26220,6 +26242,30 @@ static void ggml_cl_mul_mat_q8_0_f32_adreno(ggml_backend_t backend, const ggml_t
img_desc.buffer = b_sub_buf;
CL_CHECK((b_img = clCreateImage(context, CL_MEM_READ_ONLY, &img_fmt, &img_desc, NULL, &err), err));
+ if (use_q8_mc) {
+ kernel = backend_ctx->kernel_gemv_noshuffle_q8_0_f32_mc;
+
+ CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &q_img));
+ CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_mem), &extra0_q8_0->d));
+ CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &b_img));
+ CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_mem), &extrad->data_device));
+ CL_CHECK(clSetKernelArg(kernel, 4, sizeof(cl_ulong), &offsetd));
+ CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), &ne00));
+ CL_CHECK(clSetKernelArg(kernel, 6, sizeof(int), &ne01));
+ CL_CHECK(clSetKernelArg(kernel, 7, sizeof(int), &ne1));
+
+ size_t wavesize = backend_ctx->adreno_wave_size;
+ size_t local_work_size[] = { wavesize, 4, 1 };
+ size_t global_work_size[] = { CEIL_DIV(M, wavesize)*wavesize, 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_img));
+ CL_CHECK(clReleaseMemObject(b_sub_buf));
+ return;
+ }
+
// Split-K for small-M decode GEMVs. The base kernel puts one output row
// per lane and splits K only inside one workgroup, so M is the sole source
// of workgroup parallelism: gpt-oss's K and V projections are M=512 = 8
@@ -34170,6 +34216,23 @@ static void moe_router_reoerder(ggml_backend_t backend, const ggml_tensor * src,
CL_CHECK(clReleaseMemObject(emap_buf));
}
+// Use the per-token MoE GEMV instead of the tiled expert GEMM for small batches.
+// The GEMV costs one weight pass per token, the GEMM pads every expert to a
+// 32-token tile, so the GEMV wins while tokens rarely share an expert.
+// Default on for X2E up to 4 tokens while 2 * tokens * used < experts.
+// GGML_OPENCL_MOE_GEMV_MAX_TOK overrides with a plain token limit (1 = decode only).
+static bool use_adreno_moe_gemv_small_batch(const ggml_backend_opencl_context * backend_ctx, int n_tokens, int n_used, int n_expert) {
+ if (n_tokens <= 1) {
+ return true;
+ }
+ static const char * env = getenv("GGML_OPENCL_MOE_GEMV_MAX_TOK");
+ if (env) {
+ return n_tokens <= std::max(1, atoi(env));
+ }
+ return backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E && n_tokens <= 4 &&
+ 2 * (int64_t) n_tokens * n_used < (int64_t) n_expert;
+}
+
static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
GGML_ASSERT(src0);
GGML_ASSERT(src0->extra);
@@ -34299,7 +34362,7 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
size_t local_size[3] = {64, 2, 1};
size_t global_size[3] = {64, 2, 1};
- if (ne12 == 1) { // for gemv
+ if (use_adreno_moe_gemv_small_batch(backend_ctx, ne12, ne20, ne02)) { // for gemv
kernel = backend_ctx->kernel_gemv_moe_q4_0_f32_ns;
cl_mem src1_sub_buffer, buf_src1_image, buf_src2;
@@ -34307,14 +34370,14 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
// create a sub_buffer for src2
cl_buffer_region region;
region.origin = offset2;
- region.size = ne20 * ne21 * sizeof(int);
+ region.size = (size_t) nb21 * (ne21 - 1) + ne20 * sizeof(int);
buf_src2 = clCreateSubBuffer(extra2->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &status);
CL_CHECK(status);
// set thread grid
global_size[0] = static_cast<size_t>(((ne01 + 63) / 64) * 64);
global_size[1] = 4;
- global_size[2] = static_cast<size_t>(ne20);
+ global_size[2] = static_cast<size_t>(ne20 * ne12);
local_size[1] = 4;
// create a sub_buffer for src1
@@ -34331,6 +34394,7 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
// Set kernel args
int arg_idx = 0;
+ const int ids_stride = (int) (nb21 / sizeof(int));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q4_0->q));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q4_0->d));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &buf_src1_image));
@@ -34340,6 +34404,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne00));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne01));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne11));
+ CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne20));
+ CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ids_stride));
// launch kernel
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_size, local_size, dst);
@@ -34608,7 +34674,7 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
size_t local_size[3] = {64, 2, 1};
size_t global_size[3] = {64, 2, 1};
- if (ne12 == 1) { // for gemv
+ if (use_adreno_moe_gemv_small_batch(backend_ctx, ne12, ne20, ne02)) { // for gemv
kernel = backend_ctx->kernel_gemv_moe_q4_1_f32_ns;
cl_mem src1_sub_buffer, buf_src1_image, buf_src2;
@@ -34616,14 +34682,14 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
// create a sub_buffer for src2
cl_buffer_region region;
region.origin = offset2;
- region.size = ne20 * ne21 * sizeof(int);
+ region.size = (size_t) nb21 * (ne21 - 1) + ne20 * sizeof(int);
buf_src2 = clCreateSubBuffer(extra2->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &status);
CL_CHECK(status);
// set thread grid
global_size[0] = static_cast<size_t>(((ne01 + 63) / 64) * 64);
global_size[1] = 4;
- global_size[2] = static_cast<size_t>(ne20);
+ global_size[2] = static_cast<size_t>(ne20 * ne12);
local_size[1] = 4;
// create a sub_buffer for src1
@@ -34640,6 +34706,7 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
// Set kernel args
int arg_idx = 0;
+ const int ids_stride = (int) (nb21 / sizeof(int));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q4_1->q));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q4_1->d));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q4_1->m));
@@ -34650,6 +34717,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne00));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne01));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne11));
+ CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne20));
+ CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ids_stride));
// launch kernel
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_size, local_size, dst);
@@ -34794,7 +34863,7 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
size_t local_size[3] = {64, 2, 1};
size_t global_size[3] = {64, 2, 1};
- if (ne12 == 1) { // for gemv
+ if (use_adreno_moe_gemv_small_batch(backend_ctx, ne12, ne20, ne02)) { // for gemv
kernel = backend_ctx->kernel_gemv_moe_q5_0_f32_ns;
cl_mem src1_sub_buffer, buf_src1_image, buf_src2;
@@ -34802,14 +34871,14 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
// create a sub_buffer for src2
cl_buffer_region region;
region.origin = offset2;
- region.size = ne20 * ne21 * sizeof(int);
+ region.size = (size_t) nb21 * (ne21 - 1) + ne20 * sizeof(int);
buf_src2 = clCreateSubBuffer(extra2->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &status);
CL_CHECK(status);
// set thread grid
global_size[0] = static_cast<size_t>(((ne01 + 63) / 64) * 64);
global_size[1] = 4;
- global_size[2] = static_cast<size_t>(ne20);
+ global_size[2] = static_cast<size_t>(ne20 * ne12);
local_size[1] = 4;
// create a sub_buffer for src1
@@ -34826,6 +34895,7 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
// Set kernel args
int arg_idx = 0;
+ const int ids_stride = (int) (nb21 / sizeof(int));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q5_0->qs));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q5_0->qh));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q5_0->d));
@@ -34836,6 +34906,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne00));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne01));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne11));
+ CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne20));
+ CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ids_stride));
// launch kernel
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_size, local_size, dst);
@@ -35047,7 +35119,7 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
size_t local_size[3] = {64, 2, 1};
size_t global_size[3] = {64, 2, 1};
- if (ne12 == 1) { // for gemv
+ if (use_adreno_moe_gemv_small_batch(backend_ctx, ne12, ne20, ne02)) { // for gemv
kernel = backend_ctx->kernel_gemv_moe_q5_1_f32_ns;
cl_mem src1_sub_buffer, buf_src1_image, buf_src2;
@@ -35055,14 +35127,14 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
// create a sub_buffer for src2
cl_buffer_region region;
region.origin = offset2;
- region.size = ne20 * ne21 * sizeof(int);
+ region.size = (size_t) nb21 * (ne21 - 1) + ne20 * sizeof(int);
buf_src2 = clCreateSubBuffer(extra2->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &status);
CL_CHECK(status);
// set thread grid
global_size[0] = static_cast<size_t>(((ne01 + 63) / 64) * 64);
global_size[1] = 4;
- global_size[2] = static_cast<size_t>(ne20);
+ global_size[2] = static_cast<size_t>(ne20 * ne12);
local_size[1] = 4;
// create a sub_buffer for src1
@@ -35079,6 +35151,7 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
// Set kernel args
int arg_idx = 0;
+ const int ids_stride = (int) (nb21 / sizeof(int));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q5_1->qs));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q5_1->qh));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q5_1->d));
@@ -35090,6 +35163,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne00));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne01));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne11));
+ CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne20));
+ CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ids_stride));
// launch kernel
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_size, local_size, dst);
@@ -35479,7 +35554,7 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
size_t local_size[3] = {64, 2, 1};
size_t global_size[3] = {64, 2, 1};
- if (ne12 == 1) { // for gemv
+ if (use_adreno_moe_gemv_small_batch(backend_ctx, ne12, ne20, ne02)) { // for gemv
kernel = backend_ctx->kernel_gemv_moe_q4_k_f32_ns;
// Weight-as-texture MoE decode GEMV
@@ -35499,14 +35574,14 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
// create a sub_buffer for src2
cl_buffer_region region;
region.origin = offset2;
- region.size = ne20 * ne21 * sizeof(int);
+ region.size = (size_t) nb21 * (ne21 - 1) + ne20 * sizeof(int);
buf_src2 = clCreateSubBuffer(extra2->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &status);
CL_CHECK(status);
// set thread grid
global_size[0] = static_cast<size_t>(((ne01 + 63) / 64) * 64);
global_size[1] = 4;
- global_size[2] = static_cast<size_t>(ne20);
+ global_size[2] = static_cast<size_t>(ne20 * ne12);
local_size[1] = 4;
// create a sub_buffer for src1
@@ -35523,6 +35598,7 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
// Set kernel args
int arg_idx = 0;
+ const int ids_stride = (int) (nb21 / sizeof(int));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), use_moe_decode_wimg ? &extra0_q4_K->q_img : &extra0_q4_K->q));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q4_K->d));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q4_K->dm));
@@ -35534,6 +35610,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne00));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne01));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne11));
+ CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne20));
+ CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ids_stride));
// launch kernel
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_size, local_size, dst);
@@ -35749,7 +35827,7 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
size_t local_size[3] = {64, 2, 1};
size_t global_size[3] = {64, 2, 1};
- if (ne12 == 1) { // for gemv
+ if (use_adreno_moe_gemv_small_batch(backend_ctx, ne12, ne20, ne02)) { // for gemv
kernel = backend_ctx->kernel_gemv_moe_q5_k_f32_ns;
cl_mem src1_sub_buffer, buf_src1_image, buf_src2;
@@ -35757,14 +35835,14 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
// create a sub_buffer for src2
cl_buffer_region region;
region.origin = offset2;
- region.size = ne20 * ne21 * sizeof(int);
+ region.size = (size_t) nb21 * (ne21 - 1) + ne20 * sizeof(int);
buf_src2 = clCreateSubBuffer(extra2->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &status);
CL_CHECK(status);
// set thread grid
global_size[0] = static_cast<size_t>(((ne01 + 63) / 64) * 64);
global_size[1] = 4;
- global_size[2] = static_cast<size_t>(ne20);
+ global_size[2] = static_cast<size_t>(ne20 * ne12);
local_size[1] = 4;
// create a sub_buffer for src1
@@ -35781,6 +35859,7 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
// Set kernel args
int arg_idx = 0;
+ const int ids_stride = (int) (nb21 / sizeof(int));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q5_K->q));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q5_K->qh));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q5_K->d));
@@ -35793,6 +35872,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne00));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne01));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne11));
+ CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne20));
+ CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ids_stride));
// launch kernel
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_size, local_size, dst);
@@ -36006,7 +36087,7 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
size_t local_size[3] = {64, 2, 1};
size_t global_size[3] = {64, 2, 1};
- if (ne12 == 1) { // for gemv
+ if (use_adreno_moe_gemv_small_batch(backend_ctx, ne12, ne20, ne02)) { // for gemv
kernel = backend_ctx->kernel_gemv_moe_q6_k_f32_ns;
cl_mem src1_sub_buffer, buf_src1_image, buf_src2;
@@ -36014,14 +36095,14 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
// create a sub_buffer for src2
cl_buffer_region region;
region.origin = offset2;
- region.size = ne20 * ne21 * sizeof(int);
+ region.size = (size_t) nb21 * (ne21 - 1) + ne20 * sizeof(int);
buf_src2 = clCreateSubBuffer(extra2->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &status);
CL_CHECK(status);
// set thread grid
global_size[0] = static_cast<size_t>(((ne01 + 63) / 64) * 64);
global_size[1] = 4;
- global_size[2] = static_cast<size_t>(ne20);
+ global_size[2] = static_cast<size_t>(ne20 * ne12);
local_size[1] = 4;
// create a sub_buffer for src1
@@ -36038,6 +36119,7 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
// Set kernel args
int arg_idx = 0;
+ const int ids_stride = (int) (nb21 / sizeof(int));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q6_K->ql));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q6_K->qh));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(cl_mem), &extra0_q6_K->s));
@@ -36049,6 +36131,8 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne00));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne01));
CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne11));
+ CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ne20));
+ CL_CHECK(clSetKernelArg(kernel, arg_idx++, sizeof(int), &ids_stride));
// launch kernel
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_size, local_size, dst);
diff --git a/ggml/src/ggml-opencl/kernels/gemv_moe_q4_0_f32_ns.cl b/ggml/src/ggml-opencl/kernels/gemv_moe_q4_0_f32_ns.cl
index f90970b21..47ffed112 100644
--- a/ggml/src/ggml-opencl/kernels/gemv_moe_q4_0_f32_ns.cl
+++ b/ggml/src/ggml-opencl/kernels/gemv_moe_q4_0_f32_ns.cl
@@ -30,7 +30,9 @@ __kernel void kernel_gemv_moe_q4_0_f32_ns(
ulong offsetd,
int ne00,
int ne01,
- int ne11
+ int ne11,
+ int ne20,
+ int ids_stride
) {
uint i01 = get_global_id(0);
uint i20 = get_global_id(2);
@@ -41,9 +43,10 @@ __kernel void kernel_gemv_moe_q4_0_f32_ns(
return;
}
- uint i11 = i20 % ne11;
+ // i20 runs over n_tokens x n_expert_used; src1 has ne11 = 1 (shared) or ne20 rows per token
+ uint i11 = (i20 / ne20) * ne11 + (i20 % ne20) % ne11;
- uint expert_id = src2[i20];
+ uint expert_id = src2[(i20 / ne20) * ids_stride + i20 % ne20];
if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
diff --git a/ggml/src/ggml-opencl/kernels/gemv_moe_q4_1_f32_ns.cl b/ggml/src/ggml-opencl/kernels/gemv_moe_q4_1_f32_ns.cl
index 44bc15ff9..4041887f0 100644
--- a/ggml/src/ggml-opencl/kernels/gemv_moe_q4_1_f32_ns.cl
+++ b/ggml/src/ggml-opencl/kernels/gemv_moe_q4_1_f32_ns.cl
@@ -31,7 +31,9 @@ __kernel void kernel_gemv_moe_q4_1_f32_ns(
ulong offsetd,
int ne00,
int ne01,
- int ne11
+ int ne11,
+ int ne20,
+ int ids_stride
) {
uint i01 = get_global_id(0);
uint i20 = get_global_id(2);
@@ -42,9 +44,10 @@ __kernel void kernel_gemv_moe_q4_1_f32_ns(
return;
}
- uint i11 = i20 % ne11;
+ // i20 runs over n_tokens x n_expert_used; src1 has ne11 = 1 (shared) or ne20 rows per token
+ uint i11 = (i20 / ne20) * ne11 + (i20 % ne20) % ne11;
- uint expert_id = src2[i20];
+ uint expert_id = src2[(i20 / ne20) * ids_stride + i20 % ne20];
if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
diff --git a/ggml/src/ggml-opencl/kernels/gemv_moe_q4_k_f32_ns.cl b/ggml/src/ggml-opencl/kernels/gemv_moe_q4_k_f32_ns.cl
index 290a31593..13787568e 100644
--- a/ggml/src/ggml-opencl/kernels/gemv_moe_q4_k_f32_ns.cl
+++ b/ggml/src/ggml-opencl/kernels/gemv_moe_q4_k_f32_ns.cl
@@ -47,7 +47,9 @@ __kernel void kernel_gemv_moe_q4_k_f32_ns(
ulong offsetd,
int ne00,
int ne01,
- int ne11
+ int ne11,
+ int ne20,
+ int ids_stride
) {
uint i01 = get_global_id(0);
uint i20 = get_global_id(2);
@@ -58,9 +60,10 @@ __kernel void kernel_gemv_moe_q4_k_f32_ns(
return;
}
- uint i11 = i20 % ne11;
+ // i20 runs over n_tokens x n_expert_used; src1 has ne11 = 1 (shared) or ne20 rows per token
+ uint i11 = (i20 / ne20) * ne11 + (i20 % ne20) % ne11;
- uint expert_id = src2[i20];
+ uint expert_id = src2[(i20 / ne20) * ids_stride + i20 % ne20];
if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
@@ -169,7 +172,9 @@ __kernel void kernel_gemv_moe_q4_k_f32_ns_wimg(
ulong offsetd,
int ne00,
int ne01,
- int ne11
+ int ne11,
+ int ne20,
+ int ids_stride
) {
uint i01 = get_global_id(0);
uint i20 = get_global_id(2);
@@ -180,9 +185,10 @@ __kernel void kernel_gemv_moe_q4_k_f32_ns_wimg(
return;
}
- uint i11 = i20 % ne11;
+ // i20 runs over n_tokens x n_expert_used; src1 has ne11 = 1 (shared) or ne20 rows per token
+ uint i11 = (i20 / ne20) * ne11 + (i20 % ne20) % ne11;
- uint expert_id = src2[i20];
+ uint expert_id = src2[(i20 / ne20) * ids_stride + i20 % ne20];
if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
diff --git a/ggml/src/ggml-opencl/kernels/gemv_moe_q5_0_f32_ns.cl b/ggml/src/ggml-opencl/kernels/gemv_moe_q5_0_f32_ns.cl
index 83b186b44..f98e7bc39 100644
--- a/ggml/src/ggml-opencl/kernels/gemv_moe_q5_0_f32_ns.cl
+++ b/ggml/src/ggml-opencl/kernels/gemv_moe_q5_0_f32_ns.cl
@@ -31,7 +31,9 @@ __kernel void kernel_gemv_moe_q5_0_f32_ns(
ulong offsetd,
uint ne00,
uint ne01,
- uint ne11
+ uint ne11,
+ uint ne20,
+ uint ids_stride
) {
uint i01 = get_global_id(0);
uint i20 = get_global_id(2);
@@ -42,9 +44,10 @@ __kernel void kernel_gemv_moe_q5_0_f32_ns(
return;
}
- uint i11 = i20 % ne11;
+ // i20 runs over n_tokens x n_expert_used; src1 has ne11 = 1 (shared) or ne20 rows per token
+ uint i11 = (i20 / ne20) * ne11 + (i20 % ne20) % ne11;
- uint expert_id = src2[i20];
+ uint expert_id = src2[(i20 / ne20) * ids_stride + i20 % ne20];
if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
diff --git a/ggml/src/ggml-opencl/kernels/gemv_moe_q5_1_f32_ns.cl b/ggml/src/ggml-opencl/kernels/gemv_moe_q5_1_f32_ns.cl
index 2eba34e68..d5ec3c34d 100644
--- a/ggml/src/ggml-opencl/kernels/gemv_moe_q5_1_f32_ns.cl
+++ b/ggml/src/ggml-opencl/kernels/gemv_moe_q5_1_f32_ns.cl
@@ -32,7 +32,9 @@ __kernel void kernel_gemv_moe_q5_1_f32_ns(
ulong offsetd,
uint ne00,
uint ne01,
- uint ne11
+ uint ne11,
+ uint ne20,
+ uint ids_stride
) {
uint i01 = get_global_id(0);
uint i20 = get_global_id(2);
@@ -43,9 +45,10 @@ __kernel void kernel_gemv_moe_q5_1_f32_ns(
return;
}
- uint i11 = i20 % ne11;
+ // i20 runs over n_tokens x n_expert_used; src1 has ne11 = 1 (shared) or ne20 rows per token
+ uint i11 = (i20 / ne20) * ne11 + (i20 % ne20) % ne11;
- uint expert_id = src2[i20];
+ uint expert_id = src2[(i20 / ne20) * ids_stride + i20 % ne20];
if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
diff --git a/ggml/src/ggml-opencl/kernels/gemv_moe_q5_k_f32_ns.cl b/ggml/src/ggml-opencl/kernels/gemv_moe_q5_k_f32_ns.cl
index a959f72c3..cb41cd1d0 100644
--- a/ggml/src/ggml-opencl/kernels/gemv_moe_q5_k_f32_ns.cl
+++ b/ggml/src/ggml-opencl/kernels/gemv_moe_q5_k_f32_ns.cl
@@ -48,7 +48,9 @@ __kernel void kernel_gemv_moe_q5_k_f32_ns(
ulong offsetd,
int ne00,
int ne01,
- int ne11
+ int ne11,
+ int ne20,
+ int ids_stride
) {
uint i01 = get_global_id(0);
uint i20 = get_global_id(2);
@@ -59,9 +61,10 @@ __kernel void kernel_gemv_moe_q5_k_f32_ns(
return;
}
- uint i11 = i20 % ne11;
+ // i20 runs over n_tokens x n_expert_used; src1 has ne11 = 1 (shared) or ne20 rows per token
+ uint i11 = (i20 / ne20) * ne11 + (i20 % ne20) % ne11;
- uint expert_id = src2[i20];
+ uint expert_id = src2[(i20 / ne20) * ids_stride + i20 % ne20];
if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
diff --git a/ggml/src/ggml-opencl/kernels/gemv_moe_q6_k_f32_ns.cl b/ggml/src/ggml-opencl/kernels/gemv_moe_q6_k_f32_ns.cl
index dcd34fbc2..66cd9a6cc 100644
--- a/ggml/src/ggml-opencl/kernels/gemv_moe_q6_k_f32_ns.cl
+++ b/ggml/src/ggml-opencl/kernels/gemv_moe_q6_k_f32_ns.cl
@@ -31,7 +31,9 @@ __kernel void kernel_gemv_moe_q6_k_f32_ns(
ulong offsetd,
int ne00,
int ne01,
- int ne11
+ int ne11,
+ int ne20,
+ int ids_stride
) {
uint i01 = get_global_id(0);
uint i20 = get_global_id(2);
@@ -42,9 +44,10 @@ __kernel void kernel_gemv_moe_q6_k_f32_ns(
return;
}
- uint i11 = i20 % ne11;
+ // i20 runs over n_tokens x n_expert_used; src1 has ne11 = 1 (shared) or ne20 rows per token
+ uint i11 = (i20 / ne20) * ne11 + (i20 % ne20) % ne11;
- uint expert_id = src2[i20];
+ uint expert_id = src2[(i20 / ne20) * ids_stride + i20 % ne20];
if ((int)expert_id == -1) { // skipped slot, the dst row is zeroed by kernel_moe_zero_dst
return;
}
diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q8_0_f32.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q8_0_f32.cl
index 6109c967a..3c854c32a 100644
--- a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q8_0_f32.cl
+++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q8_0_f32.cl
@@ -201,6 +201,88 @@ __kernel void kernel_gemv_noshuffle_q8_0_f32_splitk(
}
}
+// Multi-column (n_cols in [2..4]) variant of the decode GEMV below, for the spec/MTP
+// verify batch: each block's weights are loaded once and reused across the columns.
+// Per column the operand order matches the decode GEMV.
+#define MC_COL_Q8_0(ts, c) \
+ { if (slid < 4) { regB.s0123 = read_imagef(src1, (c)*COL_STRIDE + slid*2 + k*8); \
+ regB.s4567 = read_imagef(src1, (c)*COL_STRIDE + 1 + slid*2 + k*8); } \
+ dequantizeBlockAccum_ns_sgbroadcast_1(ts, regA, scale, regB); }
+#ifdef ADRENO_GPU
+REQD_SUBGROUP_SIZE_64
+#endif
+__kernel void kernel_gemv_noshuffle_q8_0_f32_mc(
+ __read_only image1d_buffer_t src0_q, // quantized A
+ global half * src0_d, // A scales
+ __read_only image1d_buffer_t src1, // B, n_cols columns of K
+ global float * dst, // C, column-major [M x n_cols]
+ ulong offsetd,
+ int ne00, // K
+ int ne01, // M
+ int n_cols)
+{
+ 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;
+ // Keep tail lanes active for subgroup ops, but read a valid weight row.
+ uint gid_s = min(gid, LINE_STRIDE_A - 1);
+ uint BLOCK_STRIDE_A = 8 * M;
+ uint COL_STRIDE = K / 4; // float4 pixels per activation column
+
+ __private uint8 regA;
+ __private float8 regB;
+
+ __private float ts0 = 0.0f;
+ __private float ts1 = 0.0f;
+ __private float ts2 = 0.0f;
+ __private float ts3 = 0.0f;
+
+ #pragma unroll 1
+ for (uint k = groupId; k < (K / QK8_0); k += N_SIMDGROUP) {
+ float scale = convert_float(src0_d[gid_s + k * LINE_STRIDE_A]);
+
+ regA.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x;
+ regA.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x;
+ regA.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x;
+ regA.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x;
+ regA.s4 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x;
+ regA.s5 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x;
+ regA.s6 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x;
+ regA.s7 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x;
+
+ MC_COL_Q8_0(ts0, 0);
+ MC_COL_Q8_0(ts1, 1);
+ if (n_cols > 2) MC_COL_Q8_0(ts2, 2);
+ if (n_cols > 3) MC_COL_Q8_0(ts3, 3);
+ }
+
+ // reduction in local memory, assumes #wave=4
+ __local float4 reduceLM[SIMDGROUP_WIDTH * 3];
+ float4 acc = (float4)(ts0, ts1, ts2, ts3);
+ if (groupId == 1) reduceLM[SIMDGROUP_WIDTH * 0 + slid] = acc;
+ if (groupId == 2) reduceLM[SIMDGROUP_WIDTH * 1 + slid] = acc;
+ if (groupId == 3) reduceLM[SIMDGROUP_WIDTH * 2 + slid] = acc;
+ barrier(CLK_LOCAL_MEM_FENCE);
+ if (groupId == 0) {
+ acc += reduceLM[SIMDGROUP_WIDTH * 0 + slid];
+ acc += reduceLM[SIMDGROUP_WIDTH * 1 + slid];
+ acc += reduceLM[SIMDGROUP_WIDTH * 2 + slid];
+ dst = (global float*)((global char*)dst + offsetd);
+ if (gid < M) {
+ dst[0 * M + gid] = acc.s0;
+ dst[1 * M + gid] = acc.s1;
+ if (n_cols > 2) dst[2 * M + gid] = acc.s2;
+ if (n_cols > 3) dst[3 * M + gid] = acc.s3;
+ }
+ }
+}
+#undef MC_COL_Q8_0
+
#ifdef ADRENO_GPU
REQD_SUBGROUP_SIZE_64
#endif