]> git.djapps.eu Git - pkg/ggml/sources/whisper.cpp/commit
CUDA: Fix ssm_scan_f32 data-races (llama/24360)
authorOliver Simons <redacted>
Wed, 10 Jun 2026 12:27:08 +0000 (14:27 +0200)
committerGeorgi Gerganov <redacted>
Mon, 15 Jun 2026 07:33:53 +0000 (10:33 +0300)
commitef85b26d9f0bfeba3548ea6ceb213a5191ef4c11
tree82cdb14f9c6d8ad905519c9ad5cdaba6154c1f98
parentdc794303d86cfc650f41e2545ae8cf19a7dc5548
CUDA: Fix ssm_scan_f32 data-races (llama/24360)

* Add missing syncthreads before resuing cub_temp_storage

__syncthreads() is required before being allowed to resue TempStorage
smem:
https://nvidia.github.io/cccl/unstable/cub/api/classcub_1_1BlockLoad.html#_CPPv4I0EN3cub9BlockLoad4LoadEv20RandomAccessIteratorRA14ItemsPerThread_1Ti

* Add one more missing __syncthreads

Could also double-buffer, but alternative is to simply ensure all
threads have read smem* before writing to it again in the next loop
iteration

* Remove unused smem from ssm_scan_f32
ggml/src/ggml-cuda/ssm-scan.cu