diff --git a/src/clique/CliqueManager.cc b/src/clique/CliqueManager.cc index 11b910740c..722c7ac1e1 100644 --- a/src/clique/CliqueManager.cc +++ b/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++) { diff --git a/src/include/alloc.h b/src/include/alloc.h index 622fd42f1a..0b3c11dd6a 100644 --- a/src/include/alloc.h +++ b/src/include/alloc.h @@ -109,16 +109,17 @@ finish: #define ncclCudaMalloc(...) ncclCudaMallocDebug( __FILE__, __LINE__, __VA_ARGS__) template -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) { diff --git a/src/include/comm.h b/src/include/comm.h index a71e697859..844eea5bcd 100644 --- a/src/include/comm.h +++ b/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; diff --git a/src/init.cc b/src/init.cc index 7ebe13c74d..a32f47033d 100644 --- a/src/init.cc +++ b/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)); diff --git a/src/transport.cc b/src/transport.cc index 52f39a04c4..82c0c86a26 100644 --- a/src/transport.cc +++ b/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; inRanks; 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; diff --git a/src/transport/coll_net.cc b/src/transport/coll_net.cc index cdde6eb908..2d4c12ca85 100644 --- a/src/transport/coll_net.cc +++ b/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)); diff --git a/src/transport/net.cc b/src/transport/net.cc index 1f5da03cb4..f818f2d4be 100644 --- a/src/transport/net.cc +++ b/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; } } diff --git a/src/transport/p2p.cc b/src/transport/p2p.cc index ec94bfd6d2..d6bfaf29c8 100644 --- a/src/transport/p2p.cc +++ b/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) { diff --git a/src/transport/shm.cc b/src/transport/shm.cc index 6bdfaf4878..6c28ec65d9 100644 --- a/src/transport/shm.cc +++ b/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; idevFifo, 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 +#include +#include +#include +#include +#include + +#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>(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>(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; +} diff --git a/tools/GraphBench/Makefile b/tools/GraphBench/Makefile new file mode 100644 index 0000000000..4500c0a691 --- /dev/null +++ b/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) diff --git a/tools/ib-test/ib_test.cpp b/tools/ib-test/ib_test.cpp index 7a5e09f7a5..640b4f4328 100755 --- a/tools/ib-test/ib_test.cpp +++ b/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); }