fix scatter_perf crash (#1493)

* fix scatter_perf crash

* Update src/misc/msccl/msccl_lifecycle.cc

Co-authored-by: corey-derochie-amd <161367113+corey-derochie-amd@users.noreply.github.com>

* Update src/misc/msccl/msccl_lifecycle.cc

Co-authored-by: corey-derochie-amd <161367113+corey-derochie-amd@users.noreply.github.com>

* More buffsRegisteredNonGraphMode spelling fixes.

---------

Co-authored-by: corey-derochie-amd <161367113+corey-derochie-amd@users.noreply.github.com>

[ROCm/rccl commit: ff130cce7a]

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));

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);
 

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<void*>(sendBuff)));
           if (!isManagedBuffer && recvBuff) CUDACHECK(hipPointerGetAttribute(&isManagedBuffer, HIP_POINTER_ATTRIBUTE_IS_MANAGED, const_cast<void*>(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<void*>(sendBuff)));
           if (!isManagedBuffer && recvBuff) CUDACHECK(hipPointerGetAttribute(&isManagedBuffer, HIP_POINTER_ATTRIBUTE_IS_MANAGED, const_cast<void*>(recvBuff)));