Revert "CUDA: fix race condition in FA vector kernels (ggml-org#13742)"

Nexesenex · Nexesenex · commit d676ef95f32c · 2025-05-25T04:42:14.000+02:00
This reverts commit ffd0eae.
diff --git a/ggml/src/ggml-cuda/fattn-vec-f16.cuh b/ggml/src/ggml-cuda/fattn-vec-f16.cuh
@@ -212,7 +212,6 @@ static __global__ void flash_attn_vec_ext_f16(
                 }
             }
             if (__all_sync(0xFFFFFFFF, skip)) {
-                __syncthreads();
                 continue;
             }
 #endif // GGML_USE_HIP
diff --git a/ggml/src/ggml-cuda/fattn-vec-f32.cuh b/ggml/src/ggml-cuda/fattn-vec-f32.cuh
@@ -217,7 +217,6 @@ static __global__ void flash_attn_vec_ext_f32(
                 }
             }
             if (__all_sync(0xFFFFFFFF, skip)) {
-                __syncthreads();
                 continue;
             }
 #endif // GGML_USE_HIP

Original file line number	Diff line number	Diff line change
`@@ -212,7 +212,6 @@ static __global__ void flash_attn_vec_ext_f16(`
`212`	`212`	`}`
`213`	`213`	`}`
`214`	`214`	`if (__all_sync(0xFFFFFFFF, skip)) {`
`215`		`- __syncthreads();`
`216`	`215`	`continue;`
`217`	`216`	`}`
`218`	`217`	`#endif // GGML_USE_HIP`
Original file line number	Diff line number	Diff line change
`@@ -217,7 +217,6 @@ static __global__ void flash_attn_vec_ext_f32(`
`217`	`217`	`}`
`218`	`218`	`}`
`219`	`219`	`if (__all_sync(0xFFFFFFFF, skip)) {`
`220`		`- __syncthreads();`
`221`	`220`	`continue;`
`222`	`221`	`}`
`223`	`222`	`#endif // GGML_USE_HIP`