diff --git a/projects/rccl/ext-src/mem-reg.patch b/projects/rccl/ext-src/mem-reg.patch index 0e848cf6c5..57c50a5341 100644 --- a/projects/rccl/ext-src/mem-reg.patch +++ b/projects/rccl/ext-src/mem-reg.patch @@ -1,24 +1,20 @@ diff --git a/apps/nccl/include/nccl.h b/apps/nccl/include/nccl.h -index bfdb226..7fd07a8 100644 +index bfdb226..70d15cf 100644 --- a/apps/nccl/include/nccl.h +++ b/apps/nccl/include/nccl.h -@@ -167,6 +167,14 @@ ncclResult_t pncclCommCuDevice(const ncclComm_t comm, int* device); - ncclResult_t ncclCommUserRank(const ncclComm_t comm, int* rank); - ncclResult_t pncclCommUserRank(const ncclComm_t comm, int* rank); +@@ -370,6 +370,10 @@ ncclResult_t ncclSend(const void* sendbuff, size_t count, ncclDataType_t datatyp + ncclResult_t pncclSend(const void* sendbuff, size_t count, ncclDataType_t datatype, int peer, ncclComm_t comm, + cudaStream_t stream); -+/* -+ * Register/Deregister -+ */ -+ncclResult_t ncclCommRegister(ncclComm_t comm, void* buff, size_t size, void** handle); +ncclResult_t ncclCommDeregister(ncclComm_t comm, void* handle); -+bool mscclpp_BuffIsRegistered(ncclComm_t comm, const void* buff, size_t count); ++bool mscclpp_BuffIsRegistered(ncclComm_t comm, const void* buff); +size_t mscclpp_BufferSize(ncclComm_t comm, void* handle); + - /* Reduction operation selector */ - typedef enum { ncclNumOps_dummy = 5 } ncclRedOp_dummy_t; - typedef enum { + /* + * Receive + * diff --git a/apps/nccl/src/nccl.cu b/apps/nccl/src/nccl.cu -index 022d398..2a39643 100644 +index 022d398..468fcf2 100644 --- a/apps/nccl/src/nccl.cu +++ b/apps/nccl/src/nccl.cu @@ -85,6 +85,7 @@ struct ncclComm { @@ -41,7 +37,7 @@ index 022d398..2a39643 100644 static size_t ncclTypeSize(ncclDataType_t type) { switch (type) { case ncclInt8: -@@ -561,6 +567,105 @@ NCCL_API ncclResult_t ncclRedOpDestroy(ncclRedOp_t, ncclComm_t) { +@@ -561,6 +567,107 @@ NCCL_API ncclResult_t ncclRedOpDestroy(ncclRedOp_t, ncclComm_t) { return ncclInternalError; } @@ -126,7 +122,9 @@ index 022d398..2a39643 100644 + return ncclSuccess; +} + -+bool mscclpp_BuffIsRegistered(ncclComm_t comm, const void* buff, size_t count){ ++bool mscclpp_BuffIsRegistered(ncclComm_t comm, const void* buff){ ++ if(buff == nullptr) ++ return false; + size_t buffBytes; + CUdeviceptr buffBasePtr; + MSCCLPP_CUTHROW(cuMemGetAddressRange(&buffBasePtr, &buffBytes, (CUdeviceptr)buff)); diff --git a/projects/rccl/src/include/mscclpp/mscclpp_nccl.h b/projects/rccl/src/include/mscclpp/mscclpp_nccl.h index 1bd3d8168d..bfe5c8e4c4 100644 --- a/projects/rccl/src/include/mscclpp/mscclpp_nccl.h +++ b/projects/rccl/src/include/mscclpp/mscclpp_nccl.h @@ -43,7 +43,7 @@ extern "C" { ncclResult_t mscclpp_ncclCommDeregister(mscclppComm_t comm, void* handle); - bool mscclpp_BuffIsRegistered(mscclppComm_t comm, const void* buff, size_t count); + bool mscclpp_BuffIsRegistered(mscclppComm_t comm, const void* buff); size_t mscclpp_BufferSize(mscclppComm_t comm, void* handle); diff --git a/projects/rccl/src/misc/msccl/msccl_lifecycle.cc b/projects/rccl/src/misc/msccl/msccl_lifecycle.cc index 66048767a2..e6a2560a27 100644 --- a/projects/rccl/src/misc/msccl/msccl_lifecycle.cc +++ b/projects/rccl/src/misc/msccl/msccl_lifecycle.cc @@ -524,13 +524,13 @@ ncclResult_t mscclEnqueueCheck( NCCLCHECK(mscclGetCaptureStatus(comm->rank, stream)); } - const bool sendBuffRegistered = mscclpp_BuffIsRegistered(comm->mscclpp_comm, sendBuff, count); - const bool recvBuffRegistered = mscclpp_BuffIsRegistered(comm->mscclpp_comm, recvBuff, count); + const bool sendBuffRegistered = mscclpp_BuffIsRegistered(comm->mscclpp_comm, sendBuff); + const bool recvBuffRegistered = mscclpp_BuffIsRegistered(comm->mscclpp_comm, recvBuff); const bool graphMode = threadLocalStatus.captureStatus != mscclNoCapture; - const bool buffsRegistedNonGraphMode = !graphMode && sendBuffRegistered && recvBuffRegistered; + const bool buffsRegisteredNonGraphMode = !graphMode && sendBuffRegistered && recvBuffRegistered; /* check if one rank per GPU and graph mode is enabled */ - if ((graphMode || buffsRegistedNonGraphMode) && comm->mscclCompatible && nBytes > 0 && (nBytes & 31) == 0) { + if ((graphMode || buffsRegisteredNonGraphMode) && comm->mscclCompatible && nBytes > 0 && (nBytes & 31) == 0) { bool isManagedBuffer = false; if (sendBuff) CUDACHECK(hipPointerGetAttribute(&isManagedBuffer, HIP_POINTER_ATTRIBUTE_IS_MANAGED, const_cast(sendBuff))); if (!isManagedBuffer && recvBuff) CUDACHECK(hipPointerGetAttribute(&isManagedBuffer, HIP_POINTER_ATTRIBUTE_IS_MANAGED, const_cast(recvBuff))); @@ -570,13 +570,13 @@ ncclResult_t mscclEnqueueCheck( NCCLCHECK(mscclGetCaptureStatus(comm->rank, stream)); } - const bool sendBuffRegistered = mscclpp_BuffIsRegistered(comm->mscclpp_comm, sendBuff, count); - const bool recvBuffRegistered = mscclpp_BuffIsRegistered(comm->mscclpp_comm, recvBuff, count); + const bool sendBuffRegistered = mscclpp_BuffIsRegistered(comm->mscclpp_comm, sendBuff); + const bool recvBuffRegistered = mscclpp_BuffIsRegistered(comm->mscclpp_comm, recvBuff); const bool graphMode = threadLocalStatus.captureStatus != mscclNoCapture; - const bool buffsRegistedNonGraphMode = !graphMode && sendBuffRegistered && recvBuffRegistered; + const bool buffsRegisteredNonGraphMode = !graphMode && sendBuffRegistered && recvBuffRegistered; /* check if one rank per GPU and graph mode is enabled */ - if ((graphMode || buffsRegistedNonGraphMode) && comm->mscclCompatible && nBytes > 0 && (nBytes & 31) == 0) { + if ((graphMode || buffsRegisteredNonGraphMode) && comm->mscclCompatible && nBytes > 0 && (nBytes & 31) == 0) { bool isManagedBuffer = false; if (sendBuff) CUDACHECK(hipPointerGetAttribute(&isManagedBuffer, HIP_POINTER_ATTRIBUTE_IS_MANAGED, const_cast(sendBuff))); if (!isManagedBuffer && recvBuff) CUDACHECK(hipPointerGetAttribute(&isManagedBuffer, HIP_POINTER_ATTRIBUTE_IS_MANAGED, const_cast(recvBuff)));