Revert "Fix race in the CUDA DeepSeek FA kernel (#406)"

Iwan Kawrakow · Iwan Kawrakow · commit 504fb890d90e · 2025-05-11T12:22:19.000+03:00
This reverts commit 36e6e88. I should have tested. We get NaNs.
diff --git a/ggml/src/ggml-cuda/fattn-new-mma.cu b/ggml/src/ggml-cuda/fattn-new-mma.cu
@@ -898,8 +898,6 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
             KQ_crs += __shfl_xor_sync(0xFFFFFFFF, KQ_crs, offset, WARP_SIZE);
         }
 
-        __syncthreads();
-
         // Write back combined meta data:
 #pragma unroll
         for (int imeta = 0; imeta < nmeta; ++imeta) {

Original file line number	Diff line number	Diff line change
`@@ -898,8 +898,6 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(`
`898`	`898`	`KQ_crs += __shfl_xor_sync(0xFFFFFFFF, KQ_crs, offset, WARP_SIZE);`
`899`	`899`	`}`
`900`	`900`
`901`		`- __syncthreads();`
`902`		`-`
`903`	`901`	`// Write back combined meta data:`
`904`	`902`	`#pragma unroll`
`905`	`903`	`for (int imeta = 0; imeta < nmeta; ++imeta) {`