From 65d78e9a1d37b6833b5b95805af23641478669a0 Mon Sep 17 00:00:00 2001 From: gilbertlee-amd <44450918+gilbertlee-amd@users.noreply.github.com> Date: Fri, 9 Sep 2022 12:12:25 -0600 Subject: [PATCH 1/2] GraphBench (#613) Adding simple GraphBench tool for comparing RCCL hipGraph performance --- tools/GraphBench/GraphBench.cpp | 202 ++++++++++++++++++++++++++++++++ tools/GraphBench/Makefile | 26 ++++ 2 files changed, 228 insertions(+) create mode 100644 tools/GraphBench/GraphBench.cpp create mode 100644 tools/GraphBench/Makefile diff --git a/tools/GraphBench/GraphBench.cpp b/tools/GraphBench/GraphBench.cpp new file mode 100644 index 0000000000..2f29285538 --- /dev/null +++ b/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 +#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) From dd56135a9a23f59ea5f056dfbc271bed836c575c Mon Sep 17 00:00:00 2001 From: gilbertlee-amd <44450918+gilbertlee-amd@users.noreply.github.com> Date: Fri, 9 Sep 2022 16:30:15 -0600 Subject: [PATCH 2/2] Updating stream caching (#614) - Adding non-captured hipStream for use in setup --- src/channel.cc | 4 ++-- src/clique/CliqueManager.cc | 4 ++-- src/include/alloc.h | 20 +++++++------------- src/include/comm.h | 1 + src/init.cc | 9 +++++++-- src/transport.cc | 10 +++------- src/transport/coll_net.cc | 2 +- src/transport/net.cc | 6 +++--- src/transport/p2p.cc | 4 ++-- tools/ib-test/ib_test.cpp | 4 ++-- 10 files changed, 30 insertions(+), 34 deletions(-) diff --git a/src/channel.cc b/src/channel.cc index e9cfa6664f..5018a39ae5 100644 --- a/src/channel.cc +++ b/src/channel.cc @@ -18,11 +18,11 @@ ncclResult_t initChannel(struct ncclComm* comm, int channelid) { channel->id = channelid; // Ring index to user rank table. - NCCLCHECK(ncclCudaCalloc(&channel->ring.devUserRanks, comm->nRanks)); + NCCLCHECK(ncclCudaCalloc(&channel->ring.devUserRanks, comm->nRanks, comm->sideStream)); NCCLCHECK(ncclCalloc(&channel->ring.userRanks, comm->nRanks)); // Communication structures with peers. - NCCLCHECK(ncclCudaCalloc(&channel->devPeers, comm->nRanks+1)); // The extra one rank is for collnet root (i.e. network) + NCCLCHECK(ncclCudaCalloc(&channel->devPeers, comm->nRanks+1, comm->sideStream)); // The extra one rank is for collnet root (i.e. network) NCCLCHECK(ncclCalloc(&channel->peers, comm->nRanks+1)); for (size_t i=0; inRanks+1; ++i) { for (int b=0; b -static ncclResult_t ncclCudaCallocDebug(const char *filefunc, int line, T** ptr, size_t nelem, bool isFineGrain = false) { - - // Need async stream for P2P pre-connect + CUDA Graph - static bool streamCreated = false; - static hipStream_t stream; - if (rcclParamEnableHipGraph() && !streamCreated) - { - // Create stream only once to avoid performance penalty - CUDACHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); - streamCreated = true; - } +static ncclResult_t ncclCudaCallocDebug(const char *filefunc, int line, T** ptr, size_t nelem, hipStream_t sideStream = nullptr, bool isFineGrain = false) { if (isFineGrain) CUDACHECK(hipExtMallocWithFlags((void**)ptr, nelem*sizeof(T), hipDeviceMallocFinegrained)); @@ -96,10 +86,14 @@ static ncclResult_t ncclCudaCallocDebug(const char *filefunc, int line, T** ptr, CUDACHECK(hipMalloc(ptr, nelem*sizeof(T))); if (rcclParamEnableHipGraph()) { + hipStream_t stream = sideStream; + if (stream == nullptr) + CUDACHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); + CUDACHECK(hipMemsetAsync(*ptr, 0, nelem*sizeof(T), stream)); CUDACHECK(hipStreamSynchronize(stream)); - // NOTE: Currently the re-used stream is not destroyed - //CUDACHECK(hipStreamDestroy(stream)); + if (sideStream == nullptr) + CUDACHECK(hipStreamDestroy(stream)); } else { CUDACHECK(hipMemset(*ptr, 0, nelem*sizeof(T))); CUDACHECK(hipStreamSynchronize(NULL)); diff --git a/src/include/comm.h b/src/include/comm.h index c7f04d44a7..41ac83922f 100644 --- a/src/include/comm.h +++ b/src/include/comm.h @@ -220,6 +220,7 @@ struct ncclComm { int p2pRecvCount; // [RCCL] + hipStream_t sideStream; //CliqueManager* cliqueManager; // CliqueManager handles pointer collection / distribution for clique-based kernels //int rootPid; // Process ID of root // [/RCCL] diff --git a/src/init.cc b/src/init.cc index 31cdf595b2..1f921d41f2 100644 --- a/src/init.cc +++ b/src/init.cc @@ -303,6 +303,8 @@ static ncclResult_t commFree(ncclComm_t comm) { CUDACHECK(hipStreamDestroy(comm->groupStream)); } + CUDACHECK(hipStreamDestroy(comm->sideStream)); + // Last rank frees shared resources between threads int isLast; NCCLCHECK(ncclCpuBarrierIn(comm, &isLast)); @@ -365,6 +367,9 @@ 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->doneEvent = doneEvent; comm->intDoneEvent = intDoneEvent; comm->checkPointers = ncclParamCheckPointers() == 1 ? true : false; @@ -385,7 +390,7 @@ static ncclResult_t commAlloc(ncclComm_t* comret, int ndev, int rank, int virtua comm->argsptrs[0] = &comm->devComm; #ifdef ENABLE_PROFILING - NCCLCHECK(ncclCudaCalloc(&comm->hostDevComm.devProf, MAXCHANNELS*PROFILE_NUM_LAUNCHES)); + NCCLCHECK(ncclCudaCalloc(&comm->hostDevComm.devProf, MAXCHANNELS*PROFILE_NUM_LAUNCHES, comm->sideStream)); #endif #ifdef ENABLE_COLLTRACE @@ -445,7 +450,7 @@ static ncclResult_t commAlloc(ncclComm_t* comret, int ndev, int rank, int virtua static ncclResult_t devCommSetup(ncclComm_t comm) { ncclDevCommAndChannels *devCommAndChans; - NCCLCHECK(ncclCudaCalloc(&devCommAndChans, 1)); + NCCLCHECK(ncclCudaCalloc(&devCommAndChans, 1, comm->sideStream)); comm->devComm = &devCommAndChans->comm; comm->hostDevComm.channels = devCommAndChans->channels; diff --git a/src/transport.cc b/src/transport.cc index 6b279cbdbb..da8fe8d2da 100644 --- a/src/transport.cc +++ b/src/transport.cc @@ -81,9 +81,6 @@ 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 - hipStream_t transportSetupStream; - CUDACHECK(hipStreamCreateWithFlags(&transportSetupStream, hipStreamNonBlocking)); int highestType = TRANSPORT_P2P; // track highest transport type struct ncclConnect data[2*MAXCHANNELS]; @@ -137,7 +134,7 @@ 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, sizeof(struct ncclConnector), hipMemcpyHostToDevice, transportSetupStream)); + CUDACHECK(hipMemcpyAsync(comm->channels[c].devPeers[sendPeer].send+connIndex, conn, sizeof(struct ncclConnector), hipMemcpyHostToDevice, comm->sideStream)); } } TIME_STOP(3); @@ -147,14 +144,13 @@ 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, sizeof(struct ncclConnector), hipMemcpyHostToDevice, transportSetupStream)); + CUDACHECK(hipMemcpyAsync(comm->channels[c].devPeers[recvPeer].recv+connIndex, conn, sizeof(struct ncclConnector), hipMemcpyHostToDevice, comm->sideStream)); } } TIME_STOP(4); comm->connectRecv[recvPeer+comm->nRanks*connIndex] = comm->connectSend[sendPeer+comm->nRanks*connIndex] = 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 e7bec6e70b..cb0ae2a482 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 10572d3af4..56d5748bff 100644 --- a/src/transport/net.cc +++ b/src/transport/net.cc @@ -396,7 +396,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)); } @@ -575,7 +575,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) { @@ -689,7 +689,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 5101eae3bc..481b86319c 100644 --- a/src/transport/p2p.cc +++ b/src/transport/p2p.cc @@ -115,7 +115,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; @@ -347,7 +347,7 @@ static ncclResult_t p2pProxySetup(struct ncclProxyConnection* connection, struct 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/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); }