diff --git a/projects/rccl/src/include/transport.h b/projects/rccl/src/include/transport.h index f43464dec9..01812af1a3 100644 --- a/projects/rccl/src/include/transport.h +++ b/projects/rccl/src/include/transport.h @@ -74,8 +74,10 @@ struct ncclTransport { ncclResult_t ncclTransportP2pConnect(struct ncclComm* comm, int channelId, int nrecv, int* peerRecv, int nsend, int* peerSend, int connIndex); ncclResult_t ncclTransportP2pSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, int connIndex, int* highestTransportType=NULL); +#if CUDART_VERSION >= 12010 ncclResult_t ncclNvlsSetup(struct ncclComm* comm); ncclResult_t ncclNvlsFree(struct ncclComm* comm); +#endif enum { collNetRecv=0, collNetSend=1 }; int ncclTransportCollNetSetup(struct ncclComm* comm, struct ncclTopoGraph* collNetGraph, struct ncclChannel* channel, int masterRank, int masterPeer, int collNetGraphChannelId, int type); diff --git a/projects/rccl/src/init.cc b/projects/rccl/src/init.cc index 41f3db9f64..325ede8e04 100644 --- a/projects/rccl/src/init.cc +++ b/projects/rccl/src/init.cc @@ -373,7 +373,9 @@ static ncclResult_t commFree(ncclComm_t comm) { NCCLCHECK(ncclStrongStreamDestruct(&comm->deviceStream)); } +#if CUDART_VERSION >= 12010 if (comm->nvlsSupport) NCCLCHECK(ncclNvlsFree(comm)); +#endif struct ncclDestructor* dtor = comm->destructorHead; while (dtor != nullptr) { @@ -1268,7 +1270,9 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm // Check if we can setup CollNet if (comm->collNetSupport > 0) collNetTrySetup(comm, &collNetGraph); - //NCCLCHECKGOTO(ncclNvlsSetup(comm), ret, fail); +#if CUDART_VERSION >= 12010 + NCCLCHECKGOTO(ncclNvlsSetup(comm), ret, fail); +#endif TRACE(NCCL_INIT, "rank %d nranks %d - CONNECTED %d RINGS AND TREES", rank, nranks, comm->nChannels);