wmantly/llama.cpp-xrip
1
0 Fork 0
Code Issues Pull Requests Actions 32 Packages Projects Releases Wiki Activity

opencl: refine mul mat id -1 index (#31170)

Browse Source 
Assisted-by: Codex
This commit is contained in:
lhez authored and GitHub committed 2026-10-27 07:52:07 +02:00
1 parent a2b3e121ce
commit d930fe3c26
7 files changed
+64 -58
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

docs/backend/OPENCL.md
+2
View File
@@ -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.
ggml/src/ggml-opencl/CMakeLists.txt
+1
View File
@@ -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
ggml/src/ggml-opencl/ggml-opencl.cpp
+26 -26
View File
@@ -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, &region, &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, &region, &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
ggml/src/ggml-opencl/kernels/moe_add_id_glu.cl
+8 -4
View File
@@ -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);
ggml/src/ggml-opencl/kernels/moe_combine.cl
+5 -2
View File
@@ -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);
}
ggml/src/ggml-opencl/kernels/moe_sort_by_expert.cl
-26
View File
@@ -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]),
ggml/src/ggml-opencl/kernels/mul_mat_id_utils.cl
+22
View File
@@ -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;
}
}