public_git/whisper.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/whisper.cpp.git synced 2026-10-21 01:31:53 +02:00
Code Issues Packages Projects Releases Wiki Activity

opencl: fix OOB read in adreno xmem GEMM (llama/30041)

Browse Source
This commit is contained in:
lhez
2026-10-19 09:47:10 +02:00
committed by Daniel Bevenius
parent 61ae70ec08
commit d05db0d10b
2 changed files with 11 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-opencl/ggml-opencl.cpp
+5 -3
View File
@@ -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;
ggml/src/ggml-opencl/kernels/gemm_xmem_f16_f32_os8.cl
+6 -1
View File
@@ -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++; }