CUDA: fix shared memory access condition for mmv (#10740)

JohannesGaessler · web-flow · commit 26a8406ba919 · 2024-12-09T20:07:12.000+01:00
diff --git a/ggml/src/ggml-cuda/mmv.cu b/ggml/src/ggml-cuda/mmv.cu
@@ -57,7 +57,7 @@ static __global__ void mul_mat_vec(
     if (block_size > WARP_SIZE) {
         buf_iw[tid/WARP_SIZE] = sumf;
         __syncthreads();
-        if (tid > WARP_SIZE) {
+        if (tid >= WARP_SIZE) {
             return;
         }
         sumf = buf_iw[tid];

Original file line number	Diff line number	Diff line change
`@@ -57,7 +57,7 @@ static __global__ void mul_mat_vec(`
`57`	`57`	`if (block_size > WARP_SIZE) {`
`58`	`58`	`buf_iw[tid/WARP_SIZE] = sumf;`
`59`	`59`	`__syncthreads();`
`60`		`- if (tid > WARP_SIZE) {`
	`60`	`+ if (tid >= WARP_SIZE) {`
`61`	`61`	`return;`
`62`	`62`	`}`
`63`	`63`	`sumf = buf_iw[tid];`