Revert "CUDA: fix race conditions FlashAttention kernels (ggml-org#13438)"

Nexesenex · Nexesenex · commit 6a199dd8aeb4 · 2025-05-14T20:24:32.000+02:00
This reverts commit 0208355.
diff --git a/ggml/src/ggml-cuda/fattn-mma-f16.cuh b/ggml/src/ggml-cuda/fattn-mma-f16.cuh
@@ -874,8 +874,6 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
             }
         }
 
-        __syncthreads();
-
         // Write back combined meta data:
 #pragma unroll
         for (int imeta = 0; imeta < nmeta; ++imeta) {
diff --git a/ggml/src/ggml-cuda/fattn-vec-f16.cuh b/ggml/src/ggml-cuda/fattn-vec-f16.cuh
@@ -168,7 +168,6 @@ static __global__ void flash_attn_vec_ext_f16(
     for (int j = 0; j < ncols; ++j) {
         KQ[j*D + tid] = -HALF_MAX_HALF;
     }
-    __syncthreads();
 
     half2 VKQ[ncols] = {{0.0f, 0.0f}};
 

Original file line number	Diff line number	Diff line change
`@@ -874,8 +874,6 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(`
`874`	`874`	`}`
`875`	`875`	`}`
`876`	`876`
`877`		`- __syncthreads();`
`878`		`-`
`879`	`877`	`// Write back combined meta data:`
`880`	`878`	`#pragma unroll`
`881`	`879`	`for (int imeta = 0; imeta < nmeta; ++imeta) {`
Original file line number	Diff line number	Diff line change
`@@ -168,7 +168,6 @@ static __global__ void flash_attn_vec_ext_f16(`
`168`	`168`	`for (int j = 0; j < ncols; ++j) {`
`169`	`169`	`KQ[j*D + tid] = -HALF_MAX_HALF;`
`170`	`170`	`}`
`171`		`- __syncthreads();`
`172`	`171`
`173`	`172`	`half2 VKQ[ncols] = {{0.0f, 0.0f}};`
`174`	`173`