From d930fe3c26d97d0c32919ca884723384d27d41dc Mon Sep 17 00:00:00 2001 From: lhez Date: Mon, 26 Oct 2026 22:52:07 -0700 Subject: [PATCH] opencl: refine mul mat id -1 index (#31170) Assisted-by: Codex --- docs/backend/OPENCL.md | 2 + ggml/src/ggml-opencl/CMakeLists.txt | 1 + ggml/src/ggml-opencl/ggml-opencl.cpp | 52 +++++++++---------- .../src/ggml-opencl/kernels/moe_add_id_glu.cl | 12 +++-- ggml/src/ggml-opencl/kernels/moe_combine.cl | 7 ++- .../ggml-opencl/kernels/moe_sort_by_expert.cl | 26 ---------- .../ggml-opencl/kernels/mul_mat_id_utils.cl | 22 ++++++++ 7 files changed, 64 insertions(+), 58 deletions(-) create mode 100644 ggml/src/ggml-opencl/kernels/mul_mat_id_utils.cl 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; + } +}