diff --git a/projects/rccl/src/include/alloc.h b/projects/rccl/src/include/alloc.h index 70ac530490..0dad1eb1ad 100644 --- a/projects/rccl/src/include/alloc.h +++ b/projects/rccl/src/include/alloc.h @@ -56,16 +56,22 @@ extern struct allocationTracker allocTracker[]; template static ncclResult_t ncclCudaCallocDebug(const char *filefunc, int line, T** ptr, size_t nelem, bool isFineGrain = false) { +#if CUDART_VERSION >= 11030 // Need async stream for P2P pre-connect + CUDA Graph hipStream_t stream; CUDACHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); +#endif if (isFineGrain) CUDACHECK(hipExtMallocWithFlags((void**)ptr, nelem*sizeof(T), hipDeviceMallocFinegrained)); else CUDACHECK(hipMalloc(ptr, nelem*sizeof(T))); +#if CUDART_VERSION >= 11030 CUDACHECK(hipMemsetAsync(*ptr, 0, nelem*sizeof(T), stream)); CUDACHECK(hipStreamSynchronize(stream)); CUDACHECK(hipStreamDestroy(stream)); +#else + CUDACHECK(hipMemset(*ptr, 0, nelem*sizeof(T))); +#endif INFO(NCCL_ALLOC, "%s:%d Cuda Alloc Size %ld pointer %p", filefunc, line, nelem*sizeof(T), *ptr); int dev; CUDACHECK(hipGetDevice(&dev)); diff --git a/projects/rccl/src/transport.cc b/projects/rccl/src/transport.cc index 860b463a22..ae47f4f329 100644 --- a/projects/rccl/src/transport.cc +++ b/projects/rccl/src/transport.cc @@ -78,10 +78,11 @@ void dumpData(struct ncclConnect* data, int ndata) { } ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, int connIndex) { +#if CUDART_VERSION >= 11030 // Stream used during transport setup; need for P2P pre-connect + CUDA Graph hipStream_t transportSetupStream; CUDACHECK(hipStreamCreateWithFlags(&transportSetupStream, hipStreamNonBlocking)); - +#endif struct ncclConnect data[2*MAXCHANNELS]; for (int i=1; inRanks; i++) { int bootstrapTag = (i<<8) + (graph ? graph->id+1 : 0); @@ -123,7 +124,11 @@ 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; +#if CUDART_VERSION >= 11030 CUDACHECK(hipMemcpyAsync(comm->channels[c].devPeers[sendPeer].send+connIndex, conn, sizeof(struct ncclConnector), hipMemcpyHostToDevice, transportSetupStream)); +#else + CUDACHECK(hipMemcpy(comm->channels[c].devPeers[sendPeer].send+connIndex, conn, sizeof(struct ncclConnector), hipMemcpyHostToDevice)); +#endif } } for (int c=0; cchannels[c].peers[recvPeer].recv + connIndex; NCCLCHECK(conn->transportComm->connect(comm, recvData++, 1, comm->rank, conn)); conn->connected = 1; +#if CUDART_VERSION >= 11030 CUDACHECK(hipMemcpyAsync(comm->channels[c].devPeers[recvPeer].recv+connIndex, conn, sizeof(struct ncclConnector), hipMemcpyHostToDevice, transportSetupStream)); +#else + CUDACHECK(hipMemcpy(comm->channels[c].devPeers[recvPeer].recv+connIndex, conn, sizeof(struct ncclConnector), hipMemcpyHostToDevice)); +#endif } } comm->connectRecv[recvPeer+comm->nRanks*connIndex] = comm->connectSend[sendPeer+comm->nRanks*connIndex] = 0; } +#if CUDART_VERSION >= 11030 CUDACHECK(hipStreamSynchronize(transportSetupStream)); CUDACHECK(hipStreamDestroy(transportSetupStream)); +#endif return ncclSuccess; }