Commit 1ae047c92 for llama.cpp
commit 1ae047c92654dc34824d1dfa00b8e3aa82de5c8a
Author: lhez <lih@qti.qualcomm.com>
Date: Sat Nov 14 14:32:48 2026 -0800
opencl: fixes for q4_k mv oob access, f16 swiglu_oai, cpy with strided dst, batched group_norm and bf16 sigmoid (#32251)
* opencl: fix q4_K mv oob read
Assisted-by: Codex
* opencl: fix kernel dispatch and oob writes for f16 swiglu_oai
Assisted-by: Codex
* opencl: fix the row broadcast for add/mul
Assisted-by: Codex
* opencl: respect destination tensor strides for cpy
Assisted-by: Codex
* opencl: handle batches in ne3 for group_norm
Assisted-by: Codex
* opencl: avoid double rounding for bf16 sigmoid
Assisted-by: Codex
diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp
index a60929abc..10165b403 100644
--- a/ggml/src/ggml-opencl/ggml-opencl.cpp
+++ b/ggml/src/ggml-opencl/ggml-opencl.cpp
@@ -878,7 +878,7 @@ struct ggml_backend_opencl_context {
cl_kernel kernel_arange;
cl_kernel kernel_clamp, kernel_clamp_nc;
cl_kernel kernel_geglu, kernel_reglu, kernel_swiglu, kernel_swiglu_oai, kernel_swiglu_clamp, kernel_geglu_erf,
- kernel_geglu_quick, kernel_geglu_f16, kernel_reglu_f16, kernel_swiglu_f16, kernel_swiglu_clamp_f16,
+ kernel_geglu_quick, kernel_geglu_f16, kernel_reglu_f16, kernel_swiglu_f16, kernel_swiglu_oai_f16, kernel_swiglu_clamp_f16,
kernel_geglu_erf_f16, kernel_geglu_quick_f16;
cl_kernel kernel_norm, kernel_norm_mul_add;
cl_kernel kernel_rms_norm, kernel_rms_norm_mul;
@@ -2349,6 +2349,7 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
CL_CHECK((backend_ctx->kernel_geglu_f16 = clCreateKernel(backend_ctx->program_glu, "kernel_geglu_f16", &err), err));
CL_CHECK((backend_ctx->kernel_reglu_f16 = clCreateKernel(backend_ctx->program_glu, "kernel_reglu_f16", &err), err));
CL_CHECK((backend_ctx->kernel_swiglu_f16 = clCreateKernel(backend_ctx->program_glu, "kernel_swiglu_f16", &err), err));
+ CL_CHECK((backend_ctx->kernel_swiglu_oai_f16 = clCreateKernel(backend_ctx->program_glu, "kernel_swiglu_oai_f16", &err), err));
CL_CHECK((backend_ctx->kernel_swiglu_clamp_f16 = clCreateKernel(backend_ctx->program_glu, "kernel_swiglu_clamp_f16", &err), err));
CL_CHECK((backend_ctx->kernel_geglu_erf_f16 = clCreateKernel(backend_ctx->program_glu, "kernel_geglu_erf_f16", &err), err));
CL_CHECK((backend_ctx->kernel_geglu_quick_f16 = clCreateKernel(backend_ctx->program_glu, "kernel_geglu_quick_f16", &err), err));
@@ -18506,7 +18507,7 @@ static void ggml_cl_add(ggml_backend_t backend, const ggml_tensor * src0, const
GGML_ASSERT(src0->type == GGML_TYPE_F32 && src1->type == GGML_TYPE_F32);
if (bcast_row) {
kernel = backend_ctx->kernel_add_row;
- const int ne = ne00 / 4;
+ const int ne = ne10 / 4;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0->data_device));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offset0));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra1->data_device));
@@ -18554,7 +18555,7 @@ static void ggml_cl_add(ggml_backend_t backend, const ggml_tensor * src0, const
const int type_src1 = (src1->type == GGML_TYPE_F32);
if (bcast_row) {
kernel = backend_ctx->kernel_add_row_f16;
- const int ne = ne00 / 4;
+ const int ne = ne10 / 4;
CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &extra0->data_device));
CL_CHECK(clSetKernelArg(kernel, 1, sizeof(cl_ulong), &offset0));
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extra1->data_device));
@@ -18754,7 +18755,7 @@ static void ggml_cl_mul(ggml_backend_t backend, const ggml_tensor * src0, const
GGML_ASSERT(ne11 == 1);
bcast_row = true;
- int ne = ne00 / 4;
+ int ne = ne10 / 4;
if (src0->type == GGML_TYPE_F32) {
kernel = backend_ctx->kernel_mul_row;
@@ -19663,6 +19664,10 @@ static void ggml_cl_sigmoid(ggml_backend_t backend, const ggml_tensor * src0, co
CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_ulong), &nb1));
CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_ulong), &nb2));
CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_ulong), &nb3));
+ if (src0->type != GGML_TYPE_F32) {
+ const int is_bf16 = dst->type == GGML_TYPE_BF16;
+ CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &is_bf16));
+ }
size_t global_work_size[] = {(size_t)ne01 * 64, (size_t)ne02, (size_t)ne03};
size_t local_work_size[] = {64, 1, 1};
@@ -19687,6 +19692,11 @@ static void ggml_cl_sigmoid(ggml_backend_t backend, const ggml_tensor * src0, co
CL_CHECK(clSetKernelArg(kernel, 2, sizeof(cl_mem), &extrad->data_device));
CL_CHECK(clSetKernelArg(kernel, 3, sizeof(cl_ulong), &offsetd));
+ if (src0->type != GGML_TYPE_F32) {
+ const int is_bf16 = dst->type == GGML_TYPE_BF16;
+ CL_CHECK(clSetKernelArg(kernel, 4, sizeof(int), &is_bf16));
+ }
+
const int64_t n = ggml_nelements(dst);
size_t global_work_size[] = {(size_t)n, 1, 1};
@@ -20299,7 +20309,7 @@ static void ggml_cl_group_norm(ggml_backend_t backend, const ggml_tensor * src0,
CL_CHECK(clSetKernelArg(kernel, 5, sizeof(int), &group_size));
CL_CHECK(clSetKernelArg(kernel, 6, sizeof(float), &eps));
- size_t global_work_size[] = {(size_t)n_groups*sgs, 1, 1};
+ size_t global_work_size[] = {(size_t)n_groups*sgs, (size_t)src0->ne[3], 1};
size_t local_work_size[] = {(size_t)sgs, 1, 1};
backend_ctx->enqueue_ndrange_kernel(kernel, 3, global_work_size, local_work_size, dst);
@@ -39234,7 +39244,11 @@ static void ggml_cl_glu(ggml_backend_t backend, const ggml_tensor * src0, const
}
break;
case GGML_GLU_OP_SWIGLU_OAI:
- kernel = backend_ctx->kernel_swiglu_oai;
+ if (dst->type == GGML_TYPE_F32) {
+ kernel = backend_ctx->kernel_swiglu_oai;
+ } else {
+ kernel = backend_ctx->kernel_swiglu_oai_f16;
+ }
break;
case GGML_GLU_OP_SWIGLU_CLAMP:
if (dst->type == GGML_TYPE_F32) {
diff --git a/ggml/src/ggml-opencl/kernels/cpy.cl b/ggml/src/ggml-opencl/kernels/cpy.cl
index 18400d37f..656aa38d5 100644
--- a/ggml/src/ggml-opencl/kernels/cpy.cl
+++ b/ggml/src/ggml-opencl/kernels/cpy.cl
@@ -33,18 +33,16 @@ kernel void kernel_cpy_f16_f16(
int i02 = get_group_id(1);
int i01 = get_group_id(0);
- int n = i03*ne02*ne01*ne00 + i02*ne01*ne00 + i01*ne00;
-
- int i3 = n / (ne2*ne1*ne0);
- int i2 = (n - i3*ne2*ne1*ne0) / (ne1*ne0);
- int i1 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0) / ne0;
- int i0 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0 - i1*ne0);
-
- global half * dst_data = (global half *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
-
for (int i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
+ const ulong n = (((ulong)i03 * ne02 + i02) * ne01 + i01) * ne00 + i00;
+ const ulong i0 = n % ne0;
+ const ulong i1 = n / ne0 % ne1;
+ const ulong i2 = n / ne0 / ne1 % ne2;
+ const ulong i3 = n / ne0 / ne1 / ne2;
+
global const half * src = (global half *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
- dst_data[i00] = src[0];
+ global half * out = (global half *)((global char *) dst + i0*nb0 + i1*nb1 + i2*nb2 + i3*nb3);
+ out[0] = src[0];
}
}
@@ -78,18 +76,16 @@ kernel void kernel_cpy_f16_f32(
int i02 = get_group_id(1);
int i01 = get_group_id(0);
- int n = i03*ne02*ne01*ne00 + i02*ne01*ne00 + i01*ne00;
-
- int i3 = n / (ne2*ne1*ne0);
- int i2 = (n - i3*ne2*ne1*ne0) / (ne1*ne0);
- int i1 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0) / ne0;
- int i0 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0 - i1*ne0);
-
- global float * dst_data = (global float *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
-
for (int i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
- global half * src = (global half *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
- dst_data[i00] = src[0];
+ const ulong n = (((ulong)i03 * ne02 + i02) * ne01 + i01) * ne00 + i00;
+ const ulong i0 = n % ne0;
+ const ulong i1 = n / ne0 % ne1;
+ const ulong i2 = n / ne0 / ne1 % ne2;
+ const ulong i3 = n / ne0 / ne1 / ne2;
+
+ global const half * src = (global half *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
+ global float * out = (global float *)((global char *) dst + i0*nb0 + i1*nb1 + i2*nb2 + i3*nb3);
+ out[0] = src[0];
}
}
@@ -122,19 +118,16 @@ kernel void kernel_cpy_f32_f16(
int i02 = get_group_id(1);
int i01 = get_group_id(0);
- int n = i03*ne02*ne01*ne00 + i02*ne01*ne00 + i01*ne00;
-
- int i3 = n / (ne2*ne1*ne0);
- int i2 = (n - i3*ne2*ne1*ne0) / (ne1*ne0);
- int i1 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0) / ne0;
- int i0 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0 - i1*ne0);
-
- global half * dst_data = (global half *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
-
for (int i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
- global const float * src = (global float *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
+ const ulong n = (((ulong)i03 * ne02 + i02) * ne01 + i01) * ne00 + i00;
+ const ulong i0 = n % ne0;
+ const ulong i1 = n / ne0 % ne1;
+ const ulong i2 = n / ne0 / ne1 % ne2;
+ const ulong i3 = n / ne0 / ne1 / ne2;
- dst_data[i00] = src[0];
+ global const float * src = (global float *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
+ global half * out = (global half *)((global char *) dst + i0*nb0 + i1*nb1 + i2*nb2 + i3*nb3);
+ out[0] = src[0];
}
}
@@ -167,19 +160,16 @@ kernel void kernel_cpy_f32_f32(
int i02 = get_group_id(1);
int i01 = get_group_id(0);
- int n = i03*ne02*ne01*ne00 + i02*ne01*ne00 + i01*ne00;
-
- int i3 = n / (ne2*ne1*ne0);
- int i2 = (n - i3*ne2*ne1*ne0) / (ne1*ne0);
- int i1 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0) / ne0;
- int i0 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0 - i1*ne0);
-
- global float * dst_data = (global float *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
-
for (int i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
- global const float * src = (global float *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
+ const ulong n = (((ulong)i03 * ne02 + i02) * ne01 + i01) * ne00 + i00;
+ const ulong i0 = n % ne0;
+ const ulong i1 = n / ne0 % ne1;
+ const ulong i2 = n / ne0 / ne1 % ne2;
+ const ulong i3 = n / ne0 / ne1 / ne2;
- dst_data[i00] = src[0];
+ global const float * src = (global float *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
+ global float * out = (global float *)((global char *) dst + i0*nb0 + i1*nb1 + i2*nb2 + i3*nb3);
+ out[0] = src[0];
}
}
@@ -226,20 +216,16 @@ kernel void kernel_cpy_f32_f32_pack(
int i02 = t % ne02;
int i03 = t / ne02;
- // linear index of the first element of this row, unflattened over dst dims
- long n = (long)row * ne00;
- int i3 = (int)(n / ((long)ne2*ne1*ne0));
- long rm = n - (long)i3*ne2*ne1*ne0;
- int i2 = (int)(rm / ((long)ne1*ne0));
- rm -= (long)i2*ne1*ne0;
- int i1 = (int)(rm / ne0);
- int i0 = (int)(rm - (long)i1*ne0);
-
- global float * dst_data = (global float *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
-
for (int i00 = lane; i00 < ne00; i00 += tpr) {
+ const ulong n = (ulong)row * ne00 + i00;
+ const ulong i0 = n % ne0;
+ const ulong i1 = n / ne0 % ne1;
+ const ulong i2 = n / ne0 / ne1 % ne2;
+ const ulong i3 = n / ne0 / ne1 / ne2;
+
global const float * src = (global float *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
- dst_data[i00] = src[0];
+ global float * out = (global float *)((global char *) dst + i0*nb0 + i1*nb1 + i2*nb2 + i3*nb3);
+ out[0] = src[0];
}
}
#endif // GGML_CL_NO_CPY_PACK
@@ -273,19 +259,16 @@ kernel void kernel_cpy_i32_i32(
int i02 = get_group_id(1);
int i01 = get_group_id(0);
- int n = i03*ne02*ne01*ne00 + i02*ne01*ne00 + i01*ne00;
-
- int i3 = n / (ne2*ne1*ne0);
- int i2 = (n - i3*ne2*ne1*ne0) / (ne1*ne0);
- int i1 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0) / ne0;
- int i0 = (n - i3*ne2*ne1*ne0 - i2*ne1*ne0 - i1*ne0);
-
- global int * dst_data = (global int *) ((global char *) dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
-
for (int i00 = get_local_id(0); i00 < ne00; i00 += get_local_size(0)) {
- global const int * src = (global int *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
+ const ulong n = (((ulong)i03 * ne02 + i02) * ne01 + i01) * ne00 + i00;
+ const ulong i0 = n % ne0;
+ const ulong i1 = n / ne0 % ne1;
+ const ulong i2 = n / ne0 / ne1 % ne2;
+ const ulong i3 = n / ne0 / ne1 / ne2;
- dst_data[i00] = src[0];
+ global const int * src = (global int *)((global char *) src0 + i03*nb03 + i02*nb02 + i01*nb01 + i00*nb00);
+ global int * out = (global int *)((global char *) dst + i0*nb0 + i1*nb1 + i2*nb2 + i3*nb3);
+ out[0] = src[0];
}
}
diff --git a/ggml/src/ggml-opencl/kernels/glu.cl b/ggml/src/ggml-opencl/kernels/glu.cl
index 30bad00f7..d8d85c709 100644
--- a/ggml/src/ggml-opencl/kernels/glu.cl
+++ b/ggml/src/ggml-opencl/kernels/glu.cl
@@ -243,6 +243,44 @@ kernel void kernel_swiglu_oai(
}
}
+kernel void kernel_swiglu_oai_f16(
+ global char * src0,
+ ulong offset0,
+ global char * src1,
+ ulong offset1,
+ global char * dst,
+ ulong offsetd,
+ ulong nb01,
+ ulong nb11,
+ int ne0,
+ ulong nb1,
+ int ne00_off,
+ int ne10_off,
+ float limit,
+ float alpha
+) {
+ src0 = (global char*)((global char*)src0 + offset0);
+ src1 = (global char*)((global char*)src1 + offset1);
+ dst = (global char*)((global char*)dst + offsetd);
+
+ global half * src0_row = (global half *) ((global char *) src0 + get_group_id(0)*nb01) + ne00_off;
+ global half * src1_row = (global half *) ((global char *) src1 + get_group_id(0)*nb11) + ne10_off;
+ global half * dst_row = (global half *) ((global char *) dst + get_group_id(0)*nb1);
+
+ for (int i0 = get_local_id(0); i0 < ne0; i0 += get_local_size(0)) {
+ float x0 = src0_row[i0];
+ float x1 = src1_row[i0];
+
+ x0 = min(x0, limit);
+ x1 = max(min(x1, limit), -limit);
+
+ float out_glu = x0 / (1.0f + exp(-x0 * alpha));
+ out_glu = out_glu * (1.0f + x1);
+
+ dst_row[i0] = out_glu;
+ }
+}
+
//------------------------------------------------------------------------------
// swiglu_clamp
//------------------------------------------------------------------------------
diff --git a/ggml/src/ggml-opencl/kernels/group_norm.cl b/ggml/src/ggml-opencl/kernels/group_norm.cl
index 8e4fa0ed1..3eba7add4 100644
--- a/ggml/src/ggml-opencl/kernels/group_norm.cl
+++ b/ggml/src/ggml-opencl/kernels/group_norm.cl
@@ -36,6 +36,9 @@ kernel void kernel_group_norm(
src0 = (global float *)((global char *)src0 + offset0);
dst = (global float *)((global char *)dst + offsetd);
+ src0 += get_group_id(1) * ne;
+ dst += get_group_id(1) * ne;
+
int start = get_group_id(0) * group_size;
int end = start + group_size;
diff --git a/ggml/src/ggml-opencl/kernels/mul_mv_q4_k_f32_flat.cl b/ggml/src/ggml-opencl/kernels/mul_mv_q4_k_f32_flat.cl
index 5316bd363..1c191f233 100644
--- a/ggml/src/ggml-opencl/kernels/mul_mv_q4_k_f32_flat.cl
+++ b/ggml/src/ggml-opencl/kernels/mul_mv_q4_k_f32_flat.cl
@@ -145,7 +145,7 @@ kernel void kernel_mul_mv_q4_K_f32_flat(
global half * d = blk_d + ib;
global half * dm = blk_dm + ib;
- for (int row = 0; row < N_DST; row++) {
+ for (int row = 0; row < N_DST && first_row + row < ne01; row++) {
sc16[0] = sc[0] & kmask1;
sc16[1] = sc[2] & kmask1;
sc16[2] = ((sc[4] >> 0) & kmask2) | ((sc[0] & kmask3) >> 2);
diff --git a/ggml/src/ggml-opencl/kernels/sigmoid.cl b/ggml/src/ggml-opencl/kernels/sigmoid.cl
index cfe9f7295..c912d69f9 100644
--- a/ggml/src/ggml-opencl/kernels/sigmoid.cl
+++ b/ggml/src/ggml-opencl/kernels/sigmoid.cl
@@ -16,16 +16,28 @@ kernel void kernel_sigmoid_f32(
dst[get_global_id(0)] = 1.0f / (1.0f + exp(-src0[get_global_id(0)]));
}
+inline half sigmoid_f16(half x, int is_bf16) {
+ float y = 1.0f / (1.0f + exp(-(float)x));
+ if (is_bf16) {
+ // Round before F16 storage to avoid double rounding on BF16 readback.
+ uint bits = as_uint(y);
+ bits += 0x7fffu + ((bits >> 16) & 1u);
+ y = as_float(bits & 0xffff0000u);
+ }
+ return y;
+}
+
kernel void kernel_sigmoid_f16(
global half * src0,
ulong offset0,
global half * dst,
- ulong offsetd
+ ulong offsetd,
+ int is_bf16
) {
src0 = (global half*)((global char*)src0 + offset0);
dst = (global half*)((global char*)dst + offsetd);
- dst[get_global_id(0)] = 1.0f / (1.0f + exp(-src0[get_global_id(0)]));
+ dst[get_global_id(0)] = sigmoid_f16(src0[get_global_id(0)], is_bf16);
}
kernel void kernel_sigmoid_f32_nc(
@@ -71,7 +83,8 @@ kernel void kernel_sigmoid_f16_nc(
ulong nb0,
ulong nb1,
ulong nb2,
- ulong nb3
+ ulong nb3,
+ int is_bf16
) {
src0 = src0 + offset0;
dst = dst + offsetd;
@@ -84,6 +97,6 @@ kernel void kernel_sigmoid_f16_nc(
global const half * x = (global const half *)(src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00);
global half * y = (global half *)(dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
- *y = 1.0f / (1.0f + exp(-(float)*x));
+ *y = sigmoid_f16(*x, is_bf16);
}
}