Update ggml-cuda/mmvq.cu

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
cuda : fix bounds check for src0 rows in MMVQ kernel
2025-06-24 17:15:19 +00:00 · 2024-06-11 17:37:32 +03:00 · 2024-06-11 11:30:12 +03:00
1 changed files with 1 additions and 1 deletions
--- 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];
        }
    }
Author	SHA1	Message	Date
Georgi Gerganov	13c5446759	Update ggml-cuda/mmvq.cu Co-authored-by: Johannes Gäßler <johannesg@5d6.de>	2024-06-11 17:37:32 +03:00
Georgi Gerganov	9df6298a91	cuda : fix bounds check for src0 rows in MMVQ kernel	2024-06-11 11:30:12 +03:00