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

* cuda : fix bounds check for src0 rows in MMVQ kernel * Update ggml-cuda/mmvq.cu Co-authored-by: Johannes Gäßler <[email protected]> --------- Co-authored-by: Johannes Gäßler <[email protected]>
ggerganov · Jun 16, 2024 · dae7aa6 · dae7aa6
1 parent 43d078f
commit dae7aa6
Showing 1 changed file with 1 addition and 1 deletion.
diff --git a/src/ggml-cuda/mmvq.cu b/src/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];
  }
  }