]> git.djapps.eu Git - pkg/ggml/sources/llama.cpp/commit
CUDA: Fix ssm_scan_f32 data-races (#24360)
authorOliver Simons <redacted>
Wed, 10 Jun 2026 12:27:08 +0000 (14:27 +0200)
committerGitHub <redacted>
Wed, 10 Jun 2026 12:27:08 +0000 (14:27 +0200)
commitfb83cc9a072b4e85b329d4f95c59478180204c87
tree4c7f997aabef30915423f8fd3bf911c7e2938b2f
parent039e20a2db9e87b2477c76cc04905f3e1acad77f
CUDA: Fix ssm_scan_f32 data-races (#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