diff --git a/projects/rccl/src/collectives/device/msccl_kernel_impl.h b/projects/rccl/src/collectives/device/msccl_kernel_impl.h index c211c86921..897481e695 100644 --- a/projects/rccl/src/collectives/device/msccl_kernel_impl.h +++ b/projects/rccl/src/collectives/device/msccl_kernel_impl.h @@ -460,7 +460,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), __ATOMIC_RELAXED); + __atomic_store_n(&mscclFlags[bid].flag, (uint64_t) COMPUTE_FLAG(workIndex, iter, step), t->type == MSCCL_REDUCE ? __ATOMIC_RELEASE : __ATOMIC_RELAXED); step++; } }