fix out-of-bounds writes

JohannesGaessler · JohannesGaessler · commit 054d4ea97c2e · 2024-06-09T09:40:53.000+02:00
diff --git a/ggml-cuda/mmq.cuh b/ggml-cuda/mmq.cuh
@@ -1349,6 +1349,10 @@ static __device__ __forceinline__ void mmq_write_back_mma(const float * __restri
 
             const int i = blockIdx.x*mmq_y + i0 + mma_C::get_i(l);
 
+            if (need_check && i >= ne0) {
+                continue;
+            }
+
             dst[j*ne0 + i] = sum[(j0/mma_C::J)*mma_C::ne + l];
         }
     }

Original file line number	Diff line number	Diff line change
`@@ -1349,6 +1349,10 @@ static __device__ __forceinline__ void mmq_write_back_mma(const float * __restri`
`1349`	`1349`
`1350`	`1350`	`const int i = blockIdx.x*mmq_y + i0 + mma_C::get_i(l);`
`1351`	`1351`
	`1352`	`+ if (need_check && i >= ne0) {`
	`1353`	`+ continue;`
	`1354`	`+ }`
	`1355`	`+`
`1352`	`1356`	`dst[jne0 + i] = sum[(j0/mma_C::J)mma_C::ne + l];`
`1353`	`1357`	`}`
`1354`	`1358`	`}`