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

Remove hipEventDisableSystemFence (#1122)

Browse Source 
There is no indication that disabling system fence has any latency improvement.
Removing it per recommendation from HIP.
This commit is contained in:
Wenkai Du
2024-03-25 08:01:57 -07:00
committed by GitHub
parent c2fc1d6809
commit 5976f757dd
3 changed files with 2 additions and 28 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/misc/msccl/msccl_setup.cc
+2 -17
View File
@@ -16,10 +16,6 @@
#include "msccl/msccl_setup.h"
#include "msccl/msccl_status.h"
#ifndef HIP_EVENT_DISABLE_FENCE
RCCL_PARAM(MscclEnableDoneEvent, "MSCCL_ENABLE_DONE_EVENT", 1);
#endif
RCCL_PARAM(MscclWorkFifoDepth, "MSCCL_WORK_FIFO_DEPTH", 256<<10);
static inline size_t computeSizeNeeded(size_t nBytes, int nScratchChunks, int nChunksPerLoop) {
@@ -417,14 +413,7 @@ ncclResult_t mscclSetupKernel(const void* sendBuff, void* recvBuff, size_t count
mscclStatus& status = mscclGetStatus();
mscclThreadLocalStatus& threadLocalStatus = mscclGetThreadLocalStatus();
bool enableDoneEvent =
#ifndef HIP_EVENT_DISABLE_FENCE
(rcclParamMscclEnableDoneEvent() == 1);
#else
true;
#endif
if (enableDoneEvent && (status.lastStream != stream && status.lastStream != nullptr)) {
if (status.lastStream != stream && status.lastStream != nullptr) {
CUDACHECK(hipStreamWaitEvent(stream, comm->doneEvent, 0));
}
@@ -526,11 +515,7 @@ ncclResult_t mscclSetupKernel(const void* sendBuff, void* recvBuff, size_t count
void *args[3] = {&comm->devComm, &devAlgo, &workPtr};
void *func = mscclKernelEntries[fnIndex];
if (enableDoneEvent) {
CUDACHECK(hipExtLaunchKernel(func, grid, block, args, 0, stream, NULL, comm->doneEvent, 0));
} else {
CUDACHECK(hipExtLaunchKernel(func, grid, block, args, 0, stream, NULL, NULL, 0));
}
CUDACHECK(hipExtLaunchKernel(func, grid, block, args, 0, stream, NULL, comm->doneEvent, 0));
status.workIndex++;
status.lastStream = stream;
return ncclSuccess;