ver4a/llama.cpp
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity

CUDA: fix race conditions FlashAttention kernels (#13438)

Browse source
This commit is contained in:
Johannes Gäßler 2025-05-10 22:22:48 +02:00 • committed by GitHub
parent d2a4ef05c6
commit 0208355f42
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
2 changed files with 3 additions and 0 deletions
Download patch file Download diff file Expand all files Collapse all files

2
ggml/src/ggml-cuda/fattn-mma-f16.cuh
View file

@ -874,6 +874,8 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
} }
} }
__syncthreads();
// Write back combined meta data: // Write back combined meta data:
#pragma unroll #pragma unroll
for (int imeta = 0; imeta < nmeta; ++imeta) { for (int imeta = 0; imeta < nmeta; ++imeta) {

1
ggml/src/ggml-cuda/fattn-vec-f16.cuh
View file

@ -168,6 +168,7 @@ static __global__ void flash_attn_vec_ext_f16(
for (int j = 0; j < ncols; ++j) { for (int j = 0; j < ncols; ++j) {
KQ[j*D + tid] = -HALF_MAX_HALF; KQ[j*D + tid] = -HALF_MAX_HALF;
} }
__syncthreads();
half2 VKQ[ncols] = {{0.0f, 0.0f}}; half2 VKQ[ncols] = {{0.0f, 0.0f}};