AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

msccl: ensure memory coherence after data receive (#1062)

Browse Source 
[ROCm/rccl commit: f7550d83b8]
This commit is contained in:
Wenkai Du
2024-01-30 08:22:50 -08:00
committed by GitHub
parent 367ccf5688
commit 3e66d42d95
1 changed files with 1 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/collectives/device/msccl_kernel_impl.h
+1 -1
View File
@@ -372,7 +372,7 @@ __device__ __forceinline__ void mscclRunInterpreter(
return;
}
if (t->hasDependence && tid == nthreads-1)
__atomic_store_n(&mscclFlags[bid].flag, (uint64_t) COMPUTE_FLAG(workIndex, iter, step), t->type == MSCCL_REDUCE ? __ATOMIC_RELEASE : __ATOMIC_RELAXED);
__atomic_store_n(&mscclFlags[bid].flag, (uint64_t) COMPUTE_FLAG(workIndex, iter, step), (t->type == MSCCL_REDUCE || t->type == MSCCL_RECV) ? __ATOMIC_RELEASE : __ATOMIC_RELAXED);
step++;
}
}