From: Georgi Gerganov <redacted>
Date: Tue, 11 Jun 2024 14:39:01 +0000 (+0300)
Subject: cuda : fix bounds check for src0 rows in MMVQ kernel (#2231)
X-Git-Tag: upstream/1.7.4~729
X-Git-Url: https://git.djapps.eu/?a=commitdiff_plain;h=99804b0f3efc039a03d2787465fab32621332294;p=pkg%2Fggml%2Fsources%2Fwhisper.cpp

cuda : fix bounds check for src0 rows in MMVQ kernel (#2231)

* cuda : fix bounds check for src0 rows in MMVQ kernel

* Update ggml-cuda/mmvq.cu

Co-authored-by: Johannes Gäßler <redacted>

---------

Co-authored-by: Johannes Gäßler <redacted>
---

diff --git a/ggml-cuda/mmvq.cu b/ggml-cuda/mmvq.cu
index 65cc1bca..d405c124 100644
--- a/ggml-cuda/mmvq.cu
+++ b/ggml-cuda/mmvq.cu
@@ -75,7 +75,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];
         }
     }