From ca4a1dfd6719e02341168f82eadc756e6ca65826 Mon Sep 17 00:00:00 2001 From: "Wen-Heng (Jack) Chung" Date: Wed, 17 May 2023 17:50:25 +0000 Subject: [PATCH] Address review feedbacks and make the flag be disabled by default. --- src/misc/msccl/msccl_setup.cc | 11 +++++++++-- 1 file changed, 9 insertions(+), 2 deletions(-) diff --git a/src/misc/msccl/msccl_setup.cc b/src/misc/msccl/msccl_setup.cc index fd1465820e..37c2aca7ba 100644 --- a/src/misc/msccl/msccl_setup.cc +++ b/src/misc/msccl/msccl_setup.cc @@ -13,7 +13,9 @@ #include "msccl/msccl_setup.h" #include "msccl/msccl_status.h" -RCCL_PARAM(MscclEnableDoneEvent, "MSCCL_ENABLE_DONE_EVENT", 0); +#ifndef HIP_EVENT_DISABLE_FENCE +RCCL_PARAM(MscclEnableDoneEvent, "MSCCL_ENABLE_DONE_EVENT", 1); +#endif ncclResult_t mscclSetupCount(struct mscclAlgo* hostAlgo, ncclComm_t comm, size_t count, ncclDataType_t dataType) { mscclStatus& status = mscclGetStatus(); @@ -262,7 +264,12 @@ ncclResult_t mscclSetupKernel(const void* sendBuff, void* recvBuff, size_t count ncclComm_t comm, hipStream_t stream) { mscclStatus& status = mscclGetStatus(); - bool enableDoneEvent = (rcclParamMscclEnableDoneEvent() == 1); + bool enableDoneEvent = +#ifndef HIP_EVENT_DISABLE_FENCE + (rcclParamMscclEnableDoneEvent() == 1); +#else + true; +#endif if (enableDoneEvent && (status.lastStream != stream && status.lastStream != nullptr)) { CUDACHECK(hipStreamWaitEvent(stream, comm->doneEvent, 0));