Merge branch 'develop' into 2.13.4

src/clique/CliqueManager.cc

@@ -190,7 +190,7 @@ ncclResult_t CliqueManager::Init(ncclUniqueId const* commId, int suffix)
       hipIpcMemHandle_t handle;
       // Allocate fine-grained device memory on rank 0 and get IPC handle for it
       // Re-usable barrier consists of (globalCount / globalSense) pair of integers
-      NCCLCHECKGOTO(ncclCudaCalloc(&m_fineGrainBarrierMem, NCCL_MAX_OPS * 2 * sizeof(int), true), res, dropback);
+      NCCLCHECKGOTO(ncclCudaCalloc(&m_fineGrainBarrierMem, NCCL_MAX_OPS * 2 * sizeof(int), nullptr, true), res, dropback);
       if (hipIpcGetMemHandle(&handle, m_fineGrainBarrierMem) != hipSuccess)
       {
         WARN("Unable to get IPC handle for barrier memory");
@@ -228,7 +228,7 @@ ncclResult_t CliqueManager::Init(ncclUniqueId const* commId, int suffix)
     // First rank prepares fine-grained memory shared across ranks used for the two barrier variables
     if (m_rank == 0)
     {
-      NCCLCHECKGOTO(ncclCudaCalloc(&m_staticGpuBarrierMem, NCCL_MAX_OPS * 2 * sizeof(int), true), res, dropback);
+      NCCLCHECKGOTO(ncclCudaCalloc(&m_staticGpuBarrierMem, NCCL_MAX_OPS * 2 * sizeof(int), nullptr, true), res, dropback);
       // Prepare all barriers
       for (int opIndex = 0; opIndex < NCCL_MAX_OPS; opIndex++)
       {

src/include/alloc.h

@@ -109,16 +109,17 @@ finish:
 #define ncclCudaMalloc(...) ncclCudaMallocDebug( __FILE__, __LINE__, __VA_ARGS__)
 
 template <typename T>
-ncclResult_t ncclCudaCallocDebug(const char *filefunc, int line, T** ptr, size_t nelem, bool isFineGrain = false) {
+ncclResult_t ncclCudaCallocDebug(const char *filefunc, int line, T** ptr, size_t nelem, hipStream_t sideStream = nullptr, bool isFineGrain = false) {
   ncclResult_t result = ncclSuccess;
   uint64_t time0=0, time1=0, time2=0;
   hipStreamCaptureMode mode = hipStreamCaptureModeRelaxed;
   *ptr = nullptr;
   CUDACHECK(hipThreadExchangeStreamCaptureMode(&mode));
   // Need a side stream so as not to interfere with graph capture.
-  hipStream_t stream;
+  hipStream_t stream = sideStream;
   time0 = clockNano();
-  CUDACHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
+  if (stream == nullptr)
+    CUDACHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
   time1 = clockNano();
   if (isFineGrain)
     CUDACHECKGOTO(hipExtMallocWithFlags((void**)ptr, nelem*sizeof(T), hipDeviceMallocFinegrained), result, finish);
@@ -127,7 +128,8 @@ ncclResult_t ncclCudaCallocDebug(const char *filefunc, int line, T** ptr, size_t
   time2 = clockNano();
   CUDACHECKGOTO(hipMemsetAsync(*ptr, 0, nelem*sizeof(T), stream), result, finish);
   CUDACHECKGOTO(hipStreamSynchronize(stream), result, finish);
-  CUDACHECKGOTO(hipStreamDestroy(stream), result, finish);
+  if (sideStream == nullptr)
+    CUDACHECKGOTO(hipStreamDestroy(stream), result, finish);
   int dev;
   CUDACHECK(hipGetDevice(&dev));
   if (dev < MAX_ALLOC_TRACK_NGPU) {

src/include/comm.h

@@ -277,6 +277,8 @@ struct ncclComm {
   int persistentRefs; // number of persistent plan-lists capturing this comm
   struct ncclTasks tasks;
 
+  hipStream_t sideStream; // [RCCL] Cached non-captured stream
+
   // user-created reduction ops
   int userRedOpCapacity, userRedOpFreeHead;
   ncclUserRedOp *userRedOps;

src/init.cc

@@ -396,6 +396,7 @@ static ncclResult_t commFree(ncclComm_t comm) {
     NCCLCHECK(dtor->fn(dtor));
     dtor = dtor->next;
   }
+  CUDACHECK(hipStreamDestroy(comm->sideStream));
 
   commPoison(comm); // Important that this does not interfere with anything used below.
 
@@ -480,6 +481,8 @@ static ncclResult_t commAlloc(ncclComm_t* comret, int ndev, int rank, int virtua
   NCCLCHECK(getBusId(comm->cudaDev, &comm->busId));
   TRACE(NCCL_INIT,"comm %p rank %d nranks %d cudaDev %d busId %lx", comm, rank, ndev, comm->cudaDev, comm->busId);
 
+  // RCCL: create persistent stream for calloc
+  CUDACHECK(hipStreamCreateWithFlags(&comm->sideStream, hipStreamNonBlocking));
   comm->checkPointers = ncclParamCheckPointers() == 1 ? true : false;
   comm->dmaBufSupport = (dmaBufSupported(comm) == ncclSuccess) ? true : false;
   comm->fatalError = ncclSuccess;
@@ -523,6 +526,7 @@ static ncclResult_t commAlloc(ncclComm_t* comret, int ndev, int rank, int virtua
 }
 
 static ncclResult_t devCommSetup(ncclComm_t comm) {
+
   NCCLCHECK(ncclStrongStreamAcquireUncaptured(&comm->deviceStream));
 
   int nRanks = comm->nRanks;
@@ -590,7 +594,7 @@ static ncclResult_t devCommSetup(ncclComm_t comm) {
 #endif
 
 #ifdef ENABLE_PROFILING
-  NCCLCHECK(ncclCudaCalloc(&tmpCommAndChans.comm.devProf, MAXCHANNELS*PROFILE_NUM_LAUNCHES));
+  NCCLCHECK(ncclCudaCalloc(&tmpCommAndChans.comm.devProf, MAXCHANNELS*PROFILE_NUM_LAUNCHES), comm->sideStream);
 #endif
 
   NCCLCHECK(ncclCudaMemcpyAsync(devCommAndChans, &tmpCommAndChans, 1, comm->deviceStream.stream));

src/transport.cc

@@ -78,12 +78,8 @@ void dumpData(struct ncclConnect* data, int ndata) {
 }
 
 ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, int connIndex, int* highestTransportType/*=NULL*/) {
-  // Stream used during transport setup; need for P2P pre-connect + CUDA Graph
   int highestType = TRANSPORT_P2P;  // track highest transport type
 
-  hipStream_t transportSetupStream;
-  CUDACHECK(hipStreamCreateWithFlags(&transportSetupStream, hipStreamNonBlocking));
-
   struct ncclConnect data[2*MAXCHANNELS];
   for (int i=1; i<comm->nRanks; i++) {
     int bootstrapTag = (i<<8) + (graph ? graph->id+1 : 0);
@@ -135,8 +131,9 @@ ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph*
         struct ncclConnector* conn = comm->channels[c].peers[sendPeer].send + connIndex;
         NCCLCHECK(conn->transportComm->connect(comm, sendData++, 1, comm->rank, conn));
         conn->connected = 1;
-        CUDACHECK(hipMemcpyAsync(&comm->channels[c].devPeers[sendPeer].send[connIndex], &conn->conn, sizeof(struct ncclConnInfo), hipMemcpyHostToDevice, transportSetupStream));
-        CUDACHECK(hipMemcpyAsync(&comm->channels[c].devPeers[sendPeer].send[connIndex], &conn->conn, sizeof(struct ncclConnInfo), hipMemcpyHostToDevice, transportSetupStream));
+
+        CUDACHECK(hipMemcpyAsync(&comm->channels[c].devPeers[sendPeer].send[connIndex], &conn->conn, sizeof(struct ncclConnInfo), hipMemcpyHostToDevice, comm->sideStream));
+        CUDACHECK(hipMemcpyAsync(&comm->channels[c].devPeers[sendPeer].send[connIndex], &conn->conn, sizeof(struct ncclConnInfo), hipMemcpyHostToDevice, comm->sideStream));
       }
     }
     TIME_STOP(3);
@@ -146,14 +143,14 @@ ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph*
         struct ncclConnector* conn = comm->channels[c].peers[recvPeer].recv + connIndex;
         NCCLCHECK(conn->transportComm->connect(comm, recvData++, 1, comm->rank, conn));
         conn->connected = 1;
-        CUDACHECK(hipMemcpyAsync(&comm->channels[c].devPeers[recvPeer].recv[connIndex], &conn->conn, sizeof(struct ncclConnInfo), hipMemcpyHostToDevice, transportSetupStream));
+
+        CUDACHECK(hipMemcpyAsync(&comm->channels[c].devPeers[recvPeer].recv[connIndex], &conn->conn, sizeof(struct ncclConnInfo), hipMemcpyHostToDevice, comm->sideStream));
       }
     }
     TIME_STOP(4);
     comm->connectRecv[recvPeer+comm->nRanks*(connIndex == NCCL_CONN_IDX_P2P_NET ? NCCL_CONN_IDX_P2P_NET : 0)] = comm->connectSend[sendPeer+comm->nRanks*(connIndex == NCCL_CONN_IDX_P2P_NET ? NCCL_CONN_IDX_P2P_NET : 0)] = 0;
   }
-  CUDACHECK(hipStreamSynchronize(transportSetupStream));
-  CUDACHECK(hipStreamDestroy(transportSetupStream));
+  CUDACHECK(hipStreamSynchronize(comm->sideStream));
   if (highestTransportType != NULL) *highestTransportType = highestType;
   TIME_PRINT("P2P Setup/Connect");
   return ncclSuccess;

src/transport/coll_net.cc

@@ -351,7 +351,7 @@ static ncclResult_t sharedBuffersInit(struct ncclComm* comm, int cuda, char** gp
   *size = state->size;
 
   if (cuda && state->cudaBuff == NULL) {
-    NCCLCHECK(ncclCudaCalloc(&state->cudaBuff, *size, cuda));
+    NCCLCHECK(ncclCudaCalloc(&state->cudaBuff, *size, comm->sideStream, cuda));
   }
   if (!cuda && state->hostBuff == NULL) {
     NCCLCHECK(ncclCudaHostCalloc(&state->hostBuff, *size));

src/transport/net.cc

@@ -387,7 +387,7 @@ static ncclResult_t sharedBuffersInit(struct ncclComm* comm, int cuda, int local
   if (size) *size = state->size;
 
   if (cuda && state->cudaBuff == NULL) {
-    NCCLCHECK(ncclCudaCalloc(&state->cudaBuff, state->size, cuda));
+    NCCLCHECK(ncclCudaCalloc(&state->cudaBuff, state->size, comm->sideStream, cuda));
     if (sameProcess == 0) {
       CUDACHECK(hipIpcGetMemHandle(&state->ipc, state->cudaBuff));
     }
@@ -566,7 +566,7 @@ static ncclResult_t sendProxyConnect(struct ncclProxyConnection* connection, str
       if (!map->sameProcess) {
         ALIGN_SIZE(map->mems[NCCL_NET_MAP_DEVMEM].size, CUDA_IPC_MIN);
       }
-      NCCLCHECK(ncclCudaCalloc(&map->mems[NCCL_NET_MAP_DEVMEM].gpuPtr, map->mems[NCCL_NET_MAP_DEVMEM].size, resources->useGdr));
+      NCCLCHECK(ncclCudaCalloc(&map->mems[NCCL_NET_MAP_DEVMEM].gpuPtr, map->mems[NCCL_NET_MAP_DEVMEM].size, comm->sideStream, resources->useGdr));
       map->mems[NCCL_NET_MAP_DEVMEM].cpuPtr = map->mems[NCCL_NET_MAP_DEVMEM].gpuPtr;
     }
     if (!map->sameProcess) {
@@ -704,7 +704,7 @@ static ncclResult_t recvProxyConnect(struct ncclProxyConnection* connection, str
 
   if (map->mems[NCCL_NET_MAP_DEVMEM].size) {
     if (resources->shared == 0) {
-      NCCLCHECK(ncclCudaCalloc(&map->mems[NCCL_NET_MAP_DEVMEM].gpuPtr, map->mems[NCCL_NET_MAP_DEVMEM].size, resources->useGdr));
+      NCCLCHECK(ncclCudaCalloc(&map->mems[NCCL_NET_MAP_DEVMEM].gpuPtr, map->mems[NCCL_NET_MAP_DEVMEM].size, comm->sideStream, resources->useGdr));
       map->mems[NCCL_NET_MAP_DEVMEM].cpuPtr = map->mems[NCCL_NET_MAP_DEVMEM].gpuPtr;
     }
   }

src/transport/p2p.cc

@@ -156,7 +156,7 @@ ncclResult_t p2pCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTop
     // Check that legacy IPC support is available (WSL WAR)
     char *dummy;
     hipIpcMemHandle_t ipc;
-    NCCLCHECK(ncclCudaCalloc(&dummy, CUDA_IPC_MIN));
+    CUDACHECK(hipMalloc(&dummy, CUDA_IPC_MIN));
     if (hipIpcGetMemHandle(&ipc, dummy) != hipSuccess) {
       INFO(NCCL_INIT|NCCL_P2P,"Legacy IPC not supported");
       *ret = 0;
@@ -425,7 +425,7 @@ static ncclResult_t p2pSendProxySetup(struct ncclProxyConnection* connection, st
     NCCLCHECK(ncclCalloc(&proxyInfo, 1));
     connection->transportResources = proxyInfo;
 
-    NCCLCHECK(ncclCudaCalloc(&proxyInfo->ceDevBuff, comm->buffSizes[NCCL_PROTO_SIMPLE], true));
+    NCCLCHECK(ncclCudaCalloc(&proxyInfo->ceDevBuff, comm->buffSizes[NCCL_PROTO_SIMPLE], comm->sideStream, true));
 
     char shmPath[PATH_MAX];
     shmPath[0] = '\0';
@@ -443,7 +443,7 @@ static ncclResult_t p2pSendProxySetup(struct ncclProxyConnection* connection, st
     int size = *((int*)reqBuff);
     if (respSize != sizeof(struct ncclP2pBuff)) return ncclInternalError;
     struct ncclP2pBuff* p2pBuff = (struct ncclP2pBuff*)respBuff;
-    NCCLCHECK(ncclCudaCalloc((char**)&p2pBuff->directPtr, size, true));
+    NCCLCHECK(ncclCudaCalloc((char**)&p2pBuff->directPtr, size, comm->sideStream, true));
     connection->transportResources = p2pBuff->directPtr;
     hipError_t res = hipIpcGetMemHandle(&p2pBuff->devIpc, p2pBuff->directPtr);
     if (res != hipSuccess) {
@@ -462,7 +462,7 @@ static ncclResult_t p2pRecvProxySetup(struct ncclProxyConnection* connection, st
   int size = *((int*)reqBuff);
   if (respSize != sizeof(struct ncclP2pBuff)) return ncclInternalError;
   struct ncclP2pBuff* p2pBuff = (struct ncclP2pBuff*)respBuff;
-  NCCLCHECK(ncclCudaCalloc((char**)&p2pBuff->directPtr, size, true));
+  NCCLCHECK(ncclCudaCalloc((char**)&p2pBuff->directPtr, size, comm->sideStream, true));
   connection->transportResources = p2pBuff->directPtr;
   hipError_t res = hipIpcGetMemHandle(&p2pBuff->devIpc, p2pBuff->directPtr);
   if (res != hipSuccess) {

src/transport/shm.cc

@@ -210,7 +210,7 @@ static ncclResult_t shmSendProxyConnect(struct ncclProxyConnection* connection,
   NCCLCHECK(ncclCalloc(&proxyInfo, 1));
   if (reqSize != sizeof(struct shmProxyInfo)) return ncclInternalError;
   memcpy(proxyInfo, reqBuff, reqSize);
-  NCCLCHECK(ncclCudaCalloc(&proxyInfo->devFifo, comm->buffSizes[NCCL_PROTO_SIMPLE]));
+  NCCLCHECK(ncclCudaCalloc(&proxyInfo->devFifo, comm->buffSizes[NCCL_PROTO_SIMPLE], comm->sideStream));
   NCCLCHECK(ncclCudaHostCalloc(&proxyInfo->ceRecvMem, 1));
   CUDACHECK(hipStreamCreateWithFlags(&proxyInfo->stream, hipStreamNonBlocking));
   for (int i=0; i<NCCL_STEPS; i++) {
@@ -228,7 +228,7 @@ static ncclResult_t shmRecvProxyConnect(struct ncclProxyConnection* connection,
   NCCLCHECK(ncclCalloc(&proxyInfo, 1));
   if (reqSize != sizeof(struct shmProxyInfo)) return ncclInternalError;
   memcpy(proxyInfo, reqBuff, reqSize);
-  NCCLCHECK(ncclCudaCalloc(&proxyInfo->devFifo, comm->buffSizes[NCCL_PROTO_SIMPLE]));
+  NCCLCHECK(ncclCudaCalloc(&proxyInfo->devFifo, comm->buffSizes[NCCL_PROTO_SIMPLE], comm->sideStream));
   NCCLCHECK(ncclCudaHostCalloc(&proxyInfo->ceRecvMem, 1));
   CUDACHECK(hipStreamCreateWithFlags(&proxyInfo->stream, hipStreamNonBlocking));
   for (int i=0; i<NCCL_STEPS; i++) {

tools/GraphBench/GraphBench.cpp

@@ -0,0 +1,202 @@
+/*
+Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <iostream>
+#include <cstdio>
+#include <string>
+#include <chrono>
+#include <hip/hip_runtime.h>
+#include <rccl/rccl.h>
+
+#define HIP_CALL(cmd)                                                 \
+  do {                                                                \
+    hipError_t error = (cmd);                                         \
+    if (error != hipSuccess)                                          \
+    {                                                                   \
+      std::cout << "Encountered HIP error (" << hipGetErrorString(error) << ") at line " \
+                << __LINE__ << " in file " << __FILE__ << "\n";         \
+      exit(-1);                                                         \
+    }                                                                   \
+  } while (0)
+
+#define NCCL_CALL(cmd) \
+  do { \
+    ncclResult_t error = (cmd);                 \
+    if (error != ncclSuccess)                   \
+    {                                           \
+      std::cout << "Encountered NCCL error (" << ncclGetErrorString(error) << ") at line " \
+                << __LINE__ << " in file " << __FILE__ << "\n";         \
+      exit(-1);                                                         \
+    }                                                                   \
+  } while (0)
+
+int main(int argc, char **argv)
+{
+  int nranks;
+  HIP_CALL(hipGetDeviceCount(&nranks));
+
+  // Initialize communicators for each rank
+  ncclComm_t comm[nranks];
+  NCCL_CALL(ncclCommInitAll(comm, nranks, NULL));
+
+  // Allocate GPU resources
+  hipStream_t stream[nranks];
+  int* iputCpu[nranks];
+  int* iputGpu[nranks];
+  int* oputGpu[nranks];
+  int* oputCpu[nranks];
+  int* pattern;
+  int* expected;
+
+  int maxN = (1<<24);
+
+  expected = (int*)calloc(maxN, sizeof(int));
+  for (int r = 0; r < nranks; r++)
+  {
+    HIP_CALL(hipSetDevice(r));
+    HIP_CALL(hipStreamCreate(&stream[r]));
+    HIP_CALL(hipMalloc((void **)&iputGpu[r], maxN * sizeof(int)));
+    HIP_CALL(hipMalloc((void **)&oputGpu[r], maxN * sizeof(int)));
+
+    iputCpu[r] = (int*) malloc(maxN * sizeof(int));
+    oputCpu[r] = (int*) malloc(maxN * sizeof(int));
+    pattern    = (int*) malloc(maxN * sizeof(int));
+
+    for (int i = 0; i < maxN; i++)
+    {
+      iputCpu[r][i] = (r * 235 + i) % 2057;
+      oputCpu[r][i] = 0;
+      expected[i] += iputCpu[r][i];
+
+      pattern[i] = -1 - i;
+    }
+
+    HIP_CALL(hipMemcpy(iputGpu[r], iputCpu[r], maxN * sizeof(int), hipMemcpyHostToDevice));
+    HIP_CALL(hipMemcpy(oputGpu[r], oputCpu[r], maxN * sizeof(int), hipMemcpyHostToDevice));
+  }
+
+  int numWarmups    = 3;
+  int numIterations = 5;
+
+  hipGraph_t graphs[nranks];
+  hipGraphExec_t graphExec[nranks];
+
+  printf("%12s", "NumBytes");
+  for (int usingGraphs = 0; usingGraphs <= 1; usingGraphs++)
+  {
+    printf("%12s", "Setup");
+    for (int i = 1; i <= numIterations; ++i)
+      printf("%11s%d", usingGraphs ? "Graph" : "NoGraph", i);
+    printf("%12s", "Avg");
+  }
+  printf("%12s\n", "Speedup");
+
+  for (int N = 1; N <= maxN; N *= 2)
+  {
+    printf("%12lu", N * sizeof(int));
+
+    double average[2] = {};
+    for (int usingGraphs = 0; usingGraphs <= 1; usingGraphs++)
+    {
+      auto setupStart = std::chrono::high_resolution_clock::now();
+      if (usingGraphs)
+      {
+        for (int r = 0; r < nranks; ++r)
+          HIP_CALL(hipStreamBeginCapture(stream[r], hipStreamCaptureModeThreadLocal));
+
+        NCCL_CALL(ncclGroupStart());
+        for (int r = 0; r < nranks; ++r)
+        {
+          HIP_CALL(hipSetDevice(r));
+          NCCL_CALL(ncclAllReduce(iputGpu[r], oputGpu[r], N, ncclInt, ncclSum, comm[r], stream[r]));
+        }
+        NCCL_CALL(ncclGroupEnd());
+
+        for (int r = 0; r < nranks; ++r)
+          HIP_CALL(hipStreamEndCapture(stream[r], &graphs[r]));
+
+        // Instantiating graphs
+        for (int r = 0; r < nranks; ++r)
+          HIP_CALL(hipGraphInstantiate(&graphExec[r], graphs[r], NULL, NULL, 0));
+      }
+      auto setupDelta = std::chrono::high_resolution_clock::now() - setupStart;
+      double setupTime = std::chrono::duration_cast<std::chrono::duration<double, std::milli>>(setupDelta).count();
+      printf("%12.3f", setupTime);
+
+      // Perform iterations
+      average[usingGraphs] = 0;
+      for (int iteration = -numWarmups; iteration < numIterations; ++iteration)
+      {
+        auto cpuStart = std::chrono::high_resolution_clock::now();
+        if (usingGraphs)
+        {
+          for (int r = 0; r < nranks; r++)
+            HIP_CALL(hipGraphLaunch(graphExec[r], stream[r]));
+        }
+        else
+        {
+          NCCL_CALL(ncclGroupStart());
+          for (int r = 0; r < nranks; ++r)
+          {
+            HIP_CALL(hipSetDevice(r));
+            NCCL_CALL(ncclAllReduce(iputGpu[r], oputGpu[r], N, ncclInt, ncclSum, comm[r], stream[r]));
+          }
+          NCCL_CALL(ncclGroupEnd());
+        }
+        for (int r = 0; r < nranks; r++)
+          HIP_CALL(hipStreamSynchronize(stream[r]));
+
+        auto cpuDelta = std::chrono::high_resolution_clock::now() - cpuStart;
+        double iterationTime = std::chrono::duration_cast<std::chrono::duration<double, std::milli>>(cpuDelta).count();
+
+        // Check result and reset
+        bool isCorrect = true;
+        for (int r = 0; r < nranks && isCorrect; r++)
+        {
+          HIP_CALL(hipMemcpy(oputCpu[r], oputGpu[r], N * sizeof(int), hipMemcpyDeviceToHost));
+          for (int i = 0; i < N; i++)
+          {
+            if (oputCpu[r][i] != expected[i])
+            {
+              isCorrect = false;
+              printf("ERROR: Expected: %d Output %d at Index %d\n", expected[i], oputCpu[r][i], i);
+              exit(1);
+            }
+          }
+          // Fill output with input for testing reasons
+          HIP_CALL(hipMemcpy(oputGpu[r], pattern, N * sizeof(int), hipMemcpyHostToDevice));
+        }
+
+        if (iteration >= 0)
+        {
+          printf("%12.3f", iterationTime); fflush(stdout);
+          average[usingGraphs] += iterationTime;
+        }
+      }
+      average[usingGraphs] /= numIterations;
+      printf("%12.3f", average[usingGraphs]);
+    }
+    printf("%12.3f\n", average[0] / average[1]);
+    fflush(stdout);
+  }
+  return 0;
+}

tools/GraphBench/Makefile

@@ -0,0 +1,26 @@
+# Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+
+# Set to where RCCL is installed
+RCCL_INSTALL=../../build/release
+
+HIP_PATH?= $(wildcard /opt/rocm)
+ifeq (,$(HIP_PATH))
+HIP_PATH=../../..
+endif
+HIPCC=$(HIP_PATH)/bin/hipcc
+
+EXE=GraphBench
+CXXFLAGS = -std=c++11 -O3 -I../../src/include -I$(RCCL_INSTALL)/include -L$(RCCL_INSTALL) -lrccl
+
+all: $(EXE)
+
+$(EXE): $(EXE).cpp $(shell find -regex ".*\.\hpp")
+	$(HIPCC) $(CXXFLAGS) $< -o $@
+
+test: $(EXE)
+	LD_LIBRARY_PATH=$(RCCL_INSTALL) RCCL_ENABLE_HIPGRAPH=1 ./$(EXE)
+
+testInfo: $(EXE)
+	NCCL_DEBUG=INFO LD_LIBRARY_PATH=$(RCCL_INSTALL) RCCL_ENABLE_HIPGRAPH=1 ./$(EXE)
+clean:
+	rm -f *.o $(EXE)

tools/ib-test/ib_test.cpp

@@ -245,7 +245,7 @@ private:
     printf("GDR Read %s\n", use_gdr_read ? "enabled" : "disabled");
 
     if (use_gdr_read) {
-      NCCLCHECK(ncclCudaCalloc(&sendDevBuffer, sendBuffSize, 1));
+      NCCLCHECK(ncclCudaCalloc(&sendDevBuffer, sendBuffSize, nullptr, 1));
       printf("Allocated sendDevBuffer %p of %d bytes, sliceSteps %d\n",
                 sendDevBuffer, sendBuffSize, sliceSteps);
     }
@@ -397,7 +397,7 @@ private:
     printf("GDR Write %s\n", use_gdr_write ? "enabled" : "disabled");
 
     if (use_gdr_write) {
-      NCCLCHECK(ncclCudaCalloc(&recvDevBuffer, recvBuffSize, 1));
+      NCCLCHECK(ncclCudaCalloc(&recvDevBuffer, recvBuffSize, nullptr, 1));
       printf("Allocated recvDevBuffer %p of %d bytes, sliceSteps %d\n",
                 recvDevBuffer, recvBuffSize, sliceSteps);
     }