From 9df6298a91f3bfb24217f1f05b2a16d6742c8f8b Mon Sep 17 00:00:00 2001
From: Georgi Gerganov <ggerganov@gmail.com>
Date: Tue, 11 Jun 2024 11:30:12 +0300
Subject: [PATCH] cuda : fix bounds check for src0 rows in MMVQ kernel

---
 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 65cc1bca..8516252c 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 && row0 + threadIdx.x < nrows_dst) {
             dst[j*nrows_dst + row0 + threadIdx.x] = tmp[j][threadIdx.x];
         }
     }