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

Skip doneEvent inside MSCCL by default.

Browse Source 
Added a RCCL_MSCCL_ENABLE_DONE_EVENT env var, set it be 0 by default.

The env var is to control whether to use doneEvent when invoking MSCCL
kernels.

Skipping doneEvent would cause the firmware to skip L2 cache flush,
resulting in overall performance improvement.


[ROCm/rccl commit: 12dba425de]
This commit is contained in:
Wen-Heng (Jack) Chung
2023-05-17 16:49:42 +00:00
parent f7184aef3f
commit a7fbe68535
1 changed files with 10 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/misc/msccl/msccl_setup.cc
+10 -2
View File
@@ -13,6 +13,8 @@
#include "msccl/msccl_setup.h"
#include "msccl/msccl_status.h"
RCCL_PARAM(MscclEnableDoneEvent, "MSCCL_ENABLE_DONE_EVENT", 0);
ncclResult_t mscclSetupCount(struct mscclAlgo* hostAlgo, ncclComm_t comm, size_t count, ncclDataType_t dataType) {
mscclStatus& status = mscclGetStatus();
status.stepSize = comm->buffSizes[hostAlgo->protocol] / NCCL_STEPS;
@@ -260,7 +262,9 @@ ncclResult_t mscclSetupKernel(const void* sendBuff, void* recvBuff, size_t count
ncclComm_t comm, hipStream_t stream) {
mscclStatus& status = mscclGetStatus();
if (status.lastStream != stream && status.lastStream != nullptr) {
bool enableDoneEvent = (rcclParamMscclEnableDoneEvent() == 1);
if (enableDoneEvent && (status.lastStream != stream && status.lastStream != nullptr)) {
CUDACHECK(hipStreamWaitEvent(stream, comm->doneEvent, 0));
}
@@ -284,7 +288,11 @@ ncclResult_t mscclSetupKernel(const void* sendBuff, void* recvBuff, size_t count
void *args[3] = {&comm->devComm, &devAlgo, &work};
void *func = mscclKernelEntries[(opFull.op * ncclNumTypes + dataType) * NCCL_NUM_PROTOCOLS + hostAlgo->protocol];
CUDACHECK(hipExtLaunchKernel(func, grid, block, args, 0, stream, NULL, comm->doneEvent, 0));
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));
}
status.workIndex++;
status.lastStream = stream;
return ncclSuccess;