Public/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-09-19 17:24:57 +02:00
Code Issues Packages Projects Releases Wiki Activity

use __threadfence_system explicitly (not in ggml_cuda_ar_signal_set)

Browse Source
This commit is contained in:
Scott Cutler
2026-04-22 22:22:35 -07:00
parent 372d40830d
commit 5028250b52
1 changed files with 2 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/allreduce.cu
+2 -2
View File
@@ -30,9 +30,7 @@
// ---------------------------------------------------------------------------
static __device__ __forceinline__ void ggml_cuda_ar_signal_set(int * p) {
__threadfence_system(); // ensure all prior writes (D2H data) are globally visible
*(volatile int *)p = 1;
__threadfence_system(); // ensure the signal itself is globally visible
}
static __device__ __forceinline__ int ggml_cuda_ar_signal_get(const int * p) {
return *(const volatile int *)p;
@@ -130,6 +128,8 @@ static __global__ void ggml_cuda_ar_f32_kernel(
if (tid == 0) {
ggml_cuda_ar_signal_set(arrival_mine);
__threadfence_system(); // ensure the signal itself is visible across all GPUs
#if GGML_CUDA_AR_WATCHDOG
int writeback = ggml_cuda_ar_signal_get(arrival_mine);
int spin = 0;