msccl: fix data corruption with MTYPE_RW (#1014)

[ROCm/rccl commit: 81602814a7]
This commit is contained in:
Wenkai Du
2023-12-11 20:33:15 -08:00
committato da GitHub
parent fe5a902f97
commit 48107b18c9
@@ -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++;
}
}