root/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggerganov/llama.cpp.git synced 2024-12-25 02:44:36 +00:00
Code Issues Actions 9 Packages Projects Releases Wiki Activity

cuda : fix data race in soft max (#5853)

Browse Source
This commit is contained in:
slaren 2024-03-03 14:26:18 +01:00 committed by GitHub
parent 231ae28f07
commit 67be2ce101
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
1 changed files with 1 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
ggml-cuda.cu
View File

@ -6904,6 +6904,7 @@ static __global__ void soft_max_f32(const float * x, const float * mask, const f
// find the sum of exps in the block // find the sum of exps in the block
tmp = warp_reduce_sum(tmp); tmp = warp_reduce_sum(tmp);
if (block_size > WARP_SIZE) { if (block_size > WARP_SIZE) {
__syncthreads();
if (warp_id == 0) { if (warp_id == 0) {
buf_iw[lane_id] = 0.0f; buf_iw[lane_id] = 0.0f;
} }