Revert "CUDA: fix race condition in MMQ stream-k fixup (ggml-org#13299)"

Nexesenex · Nexesenex · commit 3eee3f0a817f · 2025-05-09T01:33:31.000+02:00
This reverts commit 93c4e23.
diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh
@@ -2959,7 +2959,6 @@ static __global__ void mul_mat_q_stream_k_fixup(
     for (int j = threadIdx.y*WARP_SIZE + threadIdx.x; j < mmq_x; j += nwarps*WARP_SIZE) {
         ids_dst_shared[j] = ids_dst[col_low + j];
     }
-    __syncthreads();
 
     const int offset_dst = it*mmq_y;
     dst += offset_dst;

Original file line number	Diff line number	Diff line change
`@@ -2959,7 +2959,6 @@ static __global__ void mul_mat_q_stream_k_fixup(`
`2959`	`2959`	`for (int j = threadIdx.yWARP_SIZE + threadIdx.x; j < mmq_x; j += nwarpsWARP_SIZE) {`
`2960`	`2960`	`ids_dst_shared[j] = ids_dst[col_low + j];`
`2961`	`2961`	`}`
`2962`		`- __syncthreads();`
`2963`	`2962`
`2964`	`2963`	`const int offset_dst = it*mmq_y;`
`2965`	`2964`	`dst += offset_dst;`