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

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 Hasan Mukhlis
parent 5b2daffcde
commit a0a1ca04d8
1 changed files with 1 additions and 0 deletions
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
tmp = warp_reduce_sum(tmp);
if (block_size > WARP_SIZE) {
__syncthreads();
if (warp_id == 0) {
buf_iw[lane_id] = 0.0f;
}