Commit 2207c8e57 for llama.cpp
commit 2207c8e57cc74c35ae0e128483d0cabbc6d78e96
Author: lhez <lih@qti.qualcomm.com>
Date: Tue Oct 6 09:15:25 2026 -0700
opencl: fix OOB read in adreno xmem GEMM (#30041)
diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp
index 4f1f937bb..347c76931 100644
--- a/ggml/src/ggml-opencl/ggml-opencl.cpp
+++ b/ggml/src/ggml-opencl/ggml-opencl.cpp
@@ -18864,9 +18864,11 @@ static void ggml_cl_mul_mat_f16_f32_adreno_xmem(
const int kpack = K / 4;
const int npack = CEIL_DIV(M, 4);
const int os = 8;
+ // Pad weights to the 32-row tiles read by the xmem kernel.
+ const int npack_padded = CEIL_DIV(npack, os)*os;
const size_t xmem_bytes = 6144;
- const size_t weight_bytes = static_cast<size_t>(kpack) * static_cast<size_t>(npack) * 4u * sizeof(cl_half4);
+ const size_t weight_bytes = static_cast<size_t>(kpack) * static_cast<size_t>(npack_padded) * 4u * sizeof(cl_half4);
backend_ctx->prealloc_adreno_xmem_const.allocate(backend_ctx->context, xmem_bytes);
@@ -18899,14 +18901,14 @@ static void ggml_cl_mul_mat_f16_f32_adreno_xmem(
CL_CHECK(clSetKernelArg(prepack, 3, sizeof(int), &K));
CL_CHECK(clSetKernelArg(prepack, 4, sizeof(int), &M));
CL_CHECK(clSetKernelArg(prepack, 5, sizeof(int), &kpack));
- CL_CHECK(clSetKernelArg(prepack, 6, sizeof(int), &npack));
+ CL_CHECK(clSetKernelArg(prepack, 6, sizeof(int), &npack_padded));
CL_CHECK(clSetKernelArg(prepack, 7, sizeof(int), &os));
size_t lws = 256;
size_t max_wg = backend_ctx->get_kernel_workgroup_size(prepack);
if (lws > max_wg) {
lws = max_wg;
}
- size_t gws = CEIL_DIV(static_cast<size_t>(kpack) * static_cast<size_t>(npack), lws) * lws;
+ size_t gws = CEIL_DIV(static_cast<size_t>(kpack) * static_cast<size_t>(npack_padded), lws) * lws;
backend_ctx->enqueue_ndrange_kernel(prepack, 1, &gws, &lws, dst);
cl_kernel pack_src = backend_ctx->kernel_adreno_xmem_pack_src_f32;
diff --git a/ggml/src/ggml-opencl/kernels/gemm_xmem_f16_f32_os8.cl b/ggml/src/ggml-opencl/kernels/gemm_xmem_f16_f32_os8.cl
index df9d9aed0..0f17df47a 100644
--- a/ggml/src/ggml-opencl/kernels/gemm_xmem_f16_f32_os8.cl
+++ b/ggml/src/ggml-opencl/kernels/gemm_xmem_f16_f32_os8.cl
@@ -101,7 +101,7 @@ __kernel void kernel_gemm_xmem_f16_f32_os8(
const int X = get_group_id(1)*get_local_size(0) + get_local_id(0);
const int Z = get_group_id(0)*get_local_size(2) + get_local_id(2);
- if (X >= N || Z*8 >= npack) {
+ if (Z*8 >= npack) {
return;
}
@@ -198,6 +198,11 @@ __kernel void kernel_gemm_xmem_f16_f32_os8(
r7 += src1.w * weights_cache[15].scdef;
} while (coord_s < kpack);
+ // Keep all lanes active until the subgroup loads and syncs are done.
+ if (X >= N) {
+ return;
+ }
+
int coord_s_out = Z*8;
if (coord_s_out < npack) { write_imageh(dst_img, (int2)(X, coord_s_out), r0); coord_s_out++; }
if (coord_s_out < npack) { write_imageh(dst_img, (int2)(X, coord_s_out), r1); coord_s_out++; }