Use hipGraph instead of cudaGraph (#487)

[ROCm/rccl commit: 3669e12432]

projects/rccl/src/enqueue.cc

@@ -1101,7 +1101,7 @@ void* graphHelperFunc(void *args) {
   }
 }
 
-ncclResult_t ncclGetCudaGraph(ncclComm_t comm, cudaGraph_t* graph) {
+ncclResult_t ncclGetCudaGraph(ncclComm_t comm, hipGraph_t* graph) {
   comm->usingCudaGraph = 0;
 #if CUDART_VERSION >= 11030
   hipStreamCaptureStatus captureStatus;
@@ -1285,7 +1285,7 @@ ncclResult_t ncclEnqueueCheck(struct ncclInfo* info) {
         info->datatype, info->op, info->root, info->comm, info->comm->nRanks, info->stream);
 
     // Check whether we are in cuda graph mode
-    cudaGraph_t graph;
+    hipGraph_t graph;
     ncclComm_t comm = info->comm;
     NCCLCHECKGOTO(ncclGetCudaGraph(comm, &graph), ret, end);
 

projects/rccl/src/group.cc

@@ -174,7 +174,7 @@ ncclResult_t ncclGroupEnd() {
   for (int i=0; i<ncclGroupIndex; i++) doneArray[i] = 1;
   ncclResult_t ret = ncclGroupError;
   int usingCudaGraphAll = -1;
-  cudaGraph_t* graphs = NULL;
+  hipGraph_t* graphs = NULL;
   if (ret != ncclSuccess) goto group_cleanup;
 
   /* Launch async ncclCommInitRank */

projects/rccl/src/include/comm.h

@@ -15,8 +15,6 @@
 // [/RCCL]
 
 #if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
-  typedef void *cudaGraph_t;
-  typedef void *cudaGraphNode_t;
   #define HIPRT_CB
 #else
 #include "collectives.h"
@@ -216,7 +214,7 @@ struct ncclComm {
   struct ncclQueueInfo* enqueueInfo;
   int nQueueInfoCreated;
   int nQueueInfoDestroyed;
-  cudaGraphNode_t lastSetupNode;
+  hipGraphNode_t lastSetupNode;
   unsigned long long lastCudaGraphId;
   int driverVersion;
   pfn_cuMemGetAddressRange_t pfnCuMemGetAddressRange;

projects/rccl/src/include/enqueue.h

@@ -28,8 +28,8 @@ ncclResult_t ncclSetupP2pKernel(struct ncclInfo* info);
 ncclResult_t ncclSetupAsyncKernels(struct ncclComm* comm);
 template<int USING_CUDA_GRAPH>
 void HIPRT_CB ncclEnqueueHostSetup(void* arg);
-ncclResult_t ncclGetCudaGraph(ncclComm_t comm, cudaGraph_t* graph);
-ncclResult_t ncclCudaGraphHostSetup(ncclComm_t comm, cudaGraph_t graph);
+ncclResult_t ncclGetCudaGraph(ncclComm_t comm, hipGraph_t* graph);
+ncclResult_t ncclCudaGraphHostSetup(ncclComm_t comm, hipGraph_t graph);
 
 struct ncclBuffRegInfo {
   void* sendbuffsBase[NCCL_MAX_INTRA_RANKS];