Commit d930fe3c2 for llama.cpp
commit d930fe3c26d97d0c32919ca884723384d27d41dc
Author: lhez <lih@qti.qualcomm.com>
Date: Mon Oct 26 22:52:07 2026 -0700
opencl: refine mul mat id -1 index (#31170)
Assisted-by: Codex
diff --git a/docs/backend/OPENCL.md b/docs/backend/OPENCL.md
index 337b0c82a..acaaedfbf 100644
--- a/docs/backend/OPENCL.md
+++ b/docs/backend/OPENCL.md
@@ -72,6 +72,8 @@ However, A6x GPUs in phones are likely not supported due to the outdated driver
| MXFP4 | Support |
| IQ4_NL | Support |
+`MUL_MAT_ID` accepts an expert id of `-1` and writes zeros to that slot. `ADD_ID` adds no bias for `-1`, including in fused MoE operations.
+
## Model Preparation
Since common quantizations are supported now, it is recommanded to download GGUF models directly from Huggingface.
diff --git a/ggml/src/ggml-opencl/CMakeLists.txt b/ggml/src/ggml-opencl/CMakeLists.txt
index 9b977fae7..1963c6943 100644
--- a/ggml/src/ggml-opencl/CMakeLists.txt
+++ b/ggml/src/ggml-opencl/CMakeLists.txt
@@ -134,6 +134,7 @@ set(GGML_OPENCL_KERNELS
mul_mv_id_q8_0_f32_flat
mul_mv_id_mxfp4_f32
mul_mv_id_mxfp4_f32_flat
+ mul_mat_id_utils
gemm_moe_q4_0_f32_ns
gemm_moe_q4_0_q8_1_dp4a
gemv_moe_q4_0_f32_ns
diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp
index cb4cfdd75..67248264e 100644
--- a/ggml/src/ggml-opencl/ggml-opencl.cpp
+++ b/ggml/src/ggml-opencl/ggml-opencl.cpp
@@ -1843,6 +1843,21 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
GGML_LOG_CONT(".");
}
+ // mul_mat_id_utils
+ {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+ const std::string kernel_src {
+ #include "mul_mat_id_utils.cl.h"
+ };
+#else
+ const std::string kernel_src = read_file("mul_mat_id_utils.cl");
+#endif
+ cl_program prog = build_program_from_source(backend_ctx, kernel_src.c_str(), compile_opts);
+ CL_CHECK((backend_ctx->kernel_moe_zero_dst = clCreateKernel(prog, "kernel_moe_zero_dst", &err), err));
+ CL_CHECK(clReleaseProgram(prog));
+ GGML_LOG_CONT(".");
+ }
+
// moe_add_id_glu
{
#ifdef GGML_OPENCL_EMBED_KERNELS
@@ -5890,7 +5905,6 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
build_program_from_source(backend_ctx, kernel_src.c_str(), CL_moe_compile_opts);
CL_CHECK((backend_ctx->kernel_moe_histogram = clCreateKernel(prog, "kernel_moe_histogram", &err), err));
- CL_CHECK((backend_ctx->kernel_moe_zero_dst = clCreateKernel(prog, "kernel_moe_zero_dst", &err), err));
CL_CHECK((backend_ctx->kernel_moe_scan = clCreateKernel(prog, "kernel_moe_scan", &err), err));
CL_CHECK((backend_ctx->kernel_moe_fill = clCreateKernel(prog, "kernel_moe_fill", &err), err));
CL_CHECK((backend_ctx->kernel_moe_scatter = clCreateKernel(prog, "kernel_moe_scatter", &err), err));
@@ -28284,35 +28298,21 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
// the matmul kernels only write the rows an expert owns, so zero the skipped rows here
{
- cl_int err;
-
- cl_buffer_region region;
- region.origin = offset2;
- region.size = nb21 * ne21;
- cl_mem buf_ids = clCreateSubBuffer(extra2->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err);
- CL_CHECK(err);
-
- region.origin = offsetd;
- region.size = (size_t) ne0 * ne1 * ne2 * sizeof(float);
- cl_mem buf_dst = clCreateSubBuffer(extrad->data_device, 0, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &err);
- CL_CHECK(err);
-
- const cl_uint ids_stride = nb21 / sizeof(int);
-
+ const cl_ulong nb1 = dst->nb[1];
+ const cl_ulong nb2 = dst->nb[2];
cl_kernel kernel_zero = backend_ctx->kernel_moe_zero_dst;
- CL_CHECK(clSetKernelArg(kernel_zero, 0, sizeof(cl_mem), &buf_ids));
- CL_CHECK(clSetKernelArg(kernel_zero, 1, sizeof(cl_mem), &buf_dst));
- CL_CHECK(clSetKernelArg(kernel_zero, 2, sizeof(cl_uint), &ne21));
- CL_CHECK(clSetKernelArg(kernel_zero, 3, sizeof(cl_uint), &ne20));
- CL_CHECK(clSetKernelArg(kernel_zero, 4, sizeof(cl_uint), &ids_stride));
- CL_CHECK(clSetKernelArg(kernel_zero, 5, sizeof(cl_uint), &ne0));
+ CL_CHECK(clSetKernelArg(kernel_zero, 0, sizeof(cl_mem), &extra2->data_device));
+ CL_CHECK(clSetKernelArg(kernel_zero, 1, sizeof(cl_ulong), &offset2));
+ CL_CHECK(clSetKernelArg(kernel_zero, 2, sizeof(cl_ulong), &nb21));
+ CL_CHECK(clSetKernelArg(kernel_zero, 3, sizeof(cl_mem), &extrad->data_device));
+ CL_CHECK(clSetKernelArg(kernel_zero, 4, sizeof(cl_ulong), &offsetd));
+ CL_CHECK(clSetKernelArg(kernel_zero, 5, sizeof(cl_ulong), &nb1));
+ CL_CHECK(clSetKernelArg(kernel_zero, 6, sizeof(cl_ulong), &nb2));
+ CL_CHECK(clSetKernelArg(kernel_zero, 7, sizeof(int), &ne0));
size_t zero_local[3] = {1, 1, 64};
- size_t zero_global[3] = {(size_t) ne21, (size_t) ne20, 64};
+ size_t zero_global[3] = {(size_t) ne20, (size_t) ne21, 64};
backend_ctx->enqueue_ndrange_kernel(kernel_zero, 3, zero_global, zero_local, dst);
-
- CL_CHECK(clReleaseMemObject(buf_ids));
- CL_CHECK(clReleaseMemObject(buf_dst));
}
int sgs = 32; // subgroup size
diff --git a/ggml/src/ggml-opencl/kernels/moe_add_id_glu.cl b/ggml/src/ggml-opencl/kernels/moe_add_id_glu.cl
index 6a8e4fb17..df9919ec1 100644
--- a/ggml/src/ggml-opencl/kernels/moe_add_id_glu.cl
+++ b/ggml/src/ggml-opencl/kernels/moe_add_id_glu.cl
@@ -57,13 +57,17 @@ kernel void kernel_add_id_add_id_swiglu_oai(
global const float * g_row = (global const float *)(src_g + i1*nb01_g + i2*nb02_g);
global const float * u_row = (global const float *)(src_u + i1*nb01_u + i2*nb02_u);
- global const float * gb_row = (global const float *)(src_gb + i11*nb11_g);
- global const float * ub_row = (global const float *)(src_ub + i11*nb11_u);
global float * d_row = (global float *)(dst + i1*nbd1 + i2*nbd2);
for (int i0 = get_local_id(0); i0 < ne0; i0 += get_local_size(0)) {
- float x0 = g_row[i0] + gb_row[i0];
- float x1 = u_row[i0] + ub_row[i0];
+ float x0 = g_row[i0];
+ float x1 = u_row[i0];
+ if (i11 != -1) {
+ global const float * gb_row = (global const float *)(src_gb + i11*nb11_g);
+ global const float * ub_row = (global const float *)(src_ub + i11*nb11_u);
+ x0 += gb_row[i0];
+ x1 += ub_row[i0];
+ }
x0 = min(x0, limit);
x1 = max(min(x1, limit), -limit);
diff --git a/ggml/src/ggml-opencl/kernels/moe_combine.cl b/ggml/src/ggml-opencl/kernels/moe_combine.cl
index acd08dbe6..cae7a7f2a 100644
--- a/ggml/src/ggml-opencl/kernels/moe_combine.cl
+++ b/ggml/src/ggml-opencl/kernels/moe_combine.cl
@@ -42,8 +42,11 @@ __kernel void kernel_moe_combine_bias_f32(
float4 acc = (float4)(0.0f);
for (int e = 0; e < k; ++e) {
const int i11 = *((__global const int *)(I + (ulong)e*sizeof(int)));
- __global const float * Brow = (__global const float *)(B + (ulong)i11*nb_b1) + r4*4u;
- const float4 v = vload4(0, E + (uint)e*e1) + vload4(0, Brow);
+ float4 v = vload4(0, E + (uint)e*e1);
+ if (i11 != -1) {
+ __global const float * Brow = (__global const float *)(B + (ulong)i11*nb_b1) + r4*4u;
+ v += vload4(0, Brow);
+ }
acc = mad(v, (float4)(W[(uint)e*w1]), acc);
}
diff --git a/ggml/src/ggml-opencl/kernels/moe_sort_by_expert.cl b/ggml/src/ggml-opencl/kernels/moe_sort_by_expert.cl
index 5ea31afb8..e8ac31af2 100644
--- a/ggml/src/ggml-opencl/kernels/moe_sort_by_expert.cl
+++ b/ggml/src/ggml-opencl/kernels/moe_sort_by_expert.cl
@@ -80,32 +80,6 @@ __kernel void kernel_moe_scatter(
emap[tile_idx] = val;
}
-// the gemm only writes the rows an expert owns, so the skipped rows (negative ids) are zeroed here
-__kernel void kernel_moe_zero_dst(
- __global const int * input,
- __global float * dst,
- uint N,
- uint topK,
- uint n_experts,
- uint ne0
-) {
- uint n = get_global_id(0);
- uint k = get_global_id(1);
-
- if (n >= N || k >= topK) {
- return;
- }
-
- if (input[n * n_experts + k] != -1) {
- return;
- }
-
- __global float * dst_row = dst + ((ulong)n * topK + k) * ne0;
- for (uint i = get_local_id(2); i < ne0; i += get_local_size(2)) {
- dst_row[i] = 0.0f;
- }
-}
-
// Deterministic replacement for kernel_moe_scatter.
//
// kernel_moe_scatter takes each token's slot from atomic_inc(slot_counter[expert]),
diff --git a/ggml/src/ggml-opencl/kernels/mul_mat_id_utils.cl b/ggml/src/ggml-opencl/kernels/mul_mat_id_utils.cl
new file mode 100644
index 000000000..79ec21163
--- /dev/null
+++ b/ggml/src/ggml-opencl/kernels/mul_mat_id_utils.cl
@@ -0,0 +1,22 @@
+kernel void kernel_moe_zero_dst(
+ global const char * ids,
+ ulong offset_ids,
+ ulong nb_ids1,
+ global char * dst,
+ ulong offset_dst,
+ ulong nb1,
+ ulong nb2,
+ int ne0
+) {
+ const size_t i1 = get_group_id(0);
+ const size_t i2 = get_group_id(1);
+ const int expert = *((global const int *)(ids + offset_ids + i1*sizeof(int) + i2*nb_ids1));
+ if (expert != -1) {
+ return;
+ }
+
+ global float * row = (global float *)(dst + offset_dst + i1*nb1 + i2*nb2);
+ for (size_t i0 = get_local_id(2); i0 < (size_t)ne0; i0 += get_local_size(2)) {
+ row[i0] = 0.0f;
+ }
+}