Added missing __syncthreads();

JohannesGaessler · JohannesGaessler · commit 1e735d2dbfda · 2023-05-12T21:37:34.000+02:00
diff --git a/ggml-cuda.cu b/ggml-cuda.cu
@@ -229,6 +229,7 @@ template <int block_size, int qk, dequantize_kernel_t dequantize_kernel> static
     }
 
     // sum up partial sums and write back result
+    __syncthreads();
     for (int s=block_size/2; s>0; s>>=1) {
         if (tid < s) {
             tmp[tid] += tmp[tid + s];

Original file line number	Diff line number	Diff line change
`@@ -229,6 +229,7 @@ template <int block_size, int qk, dequantize_kernel_t dequantize_kernel> static`
`229`	`229`	`}`
`230`	`230`
`231`	`231`	`// sum up partial sums and write back result`
	`232`	`+ __syncthreads();`
`232`	`233`	`for (int s=block_size/2; s>0; s>>=1) {`
`233`	`234`	`if (tid < s) {`
`234`	`235`	`tmp[tid] += tmp[tid + s];`