From d05db0d10b3ff661d82797601a70a4025526324f Mon Sep 17 00:00:00 2001 From: lhez Date: Tue, 6 Oct 2026 09:15:25 -0700 Subject: [PATCH] opencl: fix OOB read in adreno xmem GEMM (llama/30041) --- ggml/src/ggml-opencl/ggml-opencl.cpp | 8 +++++--- ggml/src/ggml-opencl/kernels/gemm_xmem_f16_f32_os8.cl | 7 ++++++- 2 files changed, 11 insertions(+), 4 deletions(-) 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(kpack) * static_cast(npack) * 4u * sizeof(cl_half4); + const size_t weight_bytes = static_cast(kpack) * static_cast(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(kpack) * static_cast(npack), lws) * lws; + size_t gws = CEIL_DIV(static_cast(kpack) * static_cast(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++; }