From 19b7a836f6658e18e973af532a5cc6ad6b3a27f8 Mon Sep 17 00:00:00 2001 From: Georgi Gerganov Date: Tue, 11 Jun 2024 17:39:01 +0300 Subject: [PATCH] cuda : fix bounds check for src0 rows in MMVQ kernel (whisper/2231) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * cuda : fix bounds check for src0 rows in MMVQ kernel * Update ggml-cuda/mmvq.cu Co-authored-by: Johannes Gäßler --------- Co-authored-by: Johannes Gäßler --- ggml-cuda/mmvq.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml-cuda/mmvq.cu b/ggml-cuda/mmvq.cu index 5f056e91e5460..e8d157169544f 100644 --- a/ggml-cuda/mmvq.cu +++ b/ggml-cuda/mmvq.cu @@ -117,7 +117,7 @@ static __global__ void mul_mat_vec_q( tmp[j][i] = warp_reduce_sum(tmp[j][i]); } - if (threadIdx.x < rows_per_cuda_block) { + if (threadIdx.x < rows_per_cuda_block && (rows_per_cuda_block == 1 || row0 + threadIdx.x < nrows_dst)) { dst[j*nrows_dst + row0 + threadIdx.x] = tmp[j][threadIdx.x]; } }