diff --git a/projects/rccl/src/include/net.h b/projects/rccl/src/include/net.h index f7f914b7cd..70f053e6f1 100644 --- a/projects/rccl/src/include/net.h +++ b/projects/rccl/src/include/net.h @@ -30,6 +30,7 @@ extern ncclNet_t ncclNetIb; extern ncclNet_t ncclNetSocket; extern ncclResult_t rcclNetP2pPolicy(void* handle, int isP2p); +extern int64_t ncclParamDmaBufEnable(); #if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__) extern ncclNet_t rocmNetIb; diff --git a/projects/rccl/src/transport/net_ib.cc b/projects/rccl/src/transport/net_ib.cc index deac53a3cf..7908db72ca 100644 --- a/projects/rccl/src/transport/net_ib.cc +++ b/projects/rccl/src/transport/net_ib.cc @@ -1990,7 +1990,7 @@ ib_recv: } } - useDmaBuf = (ncclIbDmaBufSupport(lComm->dev) == ncclSuccess); + useDmaBuf = (ncclIbDmaBufSupport(lComm->dev) == ncclSuccess && ncclParamDmaBufEnable()); rComm->flushEnabled = ((ncclIbGdrSupport() == ncclSuccess || useDmaBuf) && (ncclParamIbGdrFlushDisable() == 0)) ? 1 : 0; for (int i = 0; i < rComm->base.vProps.ndevs; i++) {