CUDA: Fix data-races when reusing SMEM in block_reduce

[ORippler]

Overview

block_reduce introduced in #18785 introduces data-races when SMEM is reused across block_reduce calls. This PR fixes it by double-buffering/adding __syncthreads() as needed. We may look to alternatively always synchronize to make block_reduce inherently safe/self-contained, but that is not needed when not reusing SMEM.

Additional information

• cuda: fix race condition in async tensor copy on virtual devices #26344 surfaced some reproducibility issues in the cuda-backend.

Failing softmax tests when using `compute-sanitizer --tool=racecheck` before this PR

(base) osimons@ub-osimons:~/llama.cpp$ compute-sanitizer --tool=racecheck ./build-x64-linux-gcc-reldbg/bin/test-backend-ops -o SOFT_MAX
========= COMPUTE-SANITIZER
ggml_cuda_init: found 1 CUDA devices (Total VRAM: 97250 MiB):
  Device 0: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition, compute capability 12.0, VMM: yes, VRAM: 97250 MiB
Testing 2 devices

Backend 1/2: CUDA0
  Device description: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition
  Device memory: 97250 MB (96593 MB free)

  SOFT_MAX(type=f32,ne=[16,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0): OK
  SOFT_MAX(type=f32,ne=[15,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0): OK
  SOFT_MAX(type=f32,ne=[16,1024,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0): OK
  SOFT_MAX(type=f32,ne=[15,1023,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0): OK
========= Error: Race reported between Read access at T3 block_reduce<(block_reduce_method)0, (unsigned int)1024, float>(T3, T3 *)+0x940 in common.cuh:643
=========     and Write access at T3 block_reduce<(block_reduce_method)1, (unsigned int)1024, float>(T3, T3 *)+0xaf0 in common.cuh:638 [3516 hazards]
=========
[SOFT_MAX] ERR = 0.850562643 > 0.000001000   SOFT_MAX(type=f32,ne=[1024,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0): FAIL
========= Error: Race reported between Read access at T3 block_reduce<(block_reduce_method)0, (unsigned int)0, float>(T3, T3 *)+0xd00 in common.cuh:643
=========     and Write access at T3 block_reduce<(block_reduce_method)1, (unsigned int)0, float>(T3, T3 *)+0xff0 in common.cuh:638 [3248 hazards]
=========
[SOFT_MAX] ERR = 0.861961235 > 0.000001000   SOFT_MAX(type=f32,ne=[1023,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0): FAIL

Passing softmax tests when using `compute-sanitizer --tool=racecheck` after this PR

(base) osimons@ub-osimons:~/llama.cpp$ compute-sanitizer --tool=racecheck ./build-x64-linux-gcc-reldbg/bin/test-backend-ops -o SOFT_MAX
========= COMPUTE-SANITIZER
ggml_cuda_init: found 1 CUDA devices (Total VRAM: 97250 MiB):
Device 0: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition, compute capability 12.0, VMM: yes, VRAM: 97250 MiB
Testing 2 devices

Backend 1/2: CUDA0
Device description: NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition
Device memory: 97250 MB (96593 MB free)

SOFT_MAX(type=f32,ne=[16,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0): OK
SOFT_MAX(type=f32,ne=[15,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0): OK
SOFT_MAX(type=f32,ne=[16,1024,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0): OK
SOFT_MAX(type=f32,ne=[15,1023,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0): OK
SOFT_MAX(type=f32,ne=[1024,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0): OK
SOFT_MAX(type=f32,ne=[1023,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0): OK
SOFT_MAX(type=f32,ne=[1024,1024,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0): OK
SOFT_MAX(type=f32,ne=[1023,1023,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=1.000000,max_bias=0.000000,inplace=0): OK
SOFT_MAX(type=f32,ne=[16,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=0.100000,max_bias=0.000000,inplace=0): OK
SOFT_MAX(type=f32,ne=[15,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=0.100000,max_bias=0.000000,inplace=0): OK
SOFT_MAX(type=f32,ne=[16,1024,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=0.100000,max_bias=0.000000,inplace=0): OK
SOFT_MAX(type=f32,ne=[15,1023,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=0.100000,max_bias=0.000000,inplace=0): OK
SOFT_MAX(type=f32,ne=[1024,16,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=0.100000,max_bias=0.000000,inplace=0): OK
SOFT_MAX(type=f32,ne=[1023,15,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=0.100000,max_bias=0.000000,inplace=0): OK
SOFT_MAX(type=f32,ne=[1024,1024,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=0.100000,max_bias=0.000000,inplace=0): OK
SOFT_MAX(type=f32,ne=[1023,1023,1,1],mask=0,sinks=0,m_prec=f32,nr23=[1,1],scale=0.100000,max_bias=0.000000,inplace=0): OK

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: Yes - paired with codex on this

[Kononnable]

@ORippler I can confirm that with this changes alone I get deterministic results from llama-completion runs.
It seems that the full stream sync on virtual device I proposed in #26344 just lowered the chances for negative effects of race condition fixed by this PR.