diff --git a/projects/rccl/src/include/transport.h b/projects/rccl/src/include/transport.h index 59f83c9a88..bc9b7779d8 100644 --- a/projects/rccl/src/include/transport.h +++ b/projects/rccl/src/include/transport.h @@ -53,7 +53,7 @@ struct ncclTransportComm { struct ncclTransport { const char name[4]; - ncclResult_t (*fillInfo)(ncclTinfo_t*, int); + ncclResult_t (*fillInfo)(ncclTinfo_t*, int, uint64_t); ncclResult_t (*canConnect)(ncclTvalue_t*, ncclTinfo_t*, ncclTinfo_t*); ncclResult_t (*getRings)(int, int*, int*, ncclTvalue_t*, int*, int*, int*, int, int*); struct ncclTransportComm send; diff --git a/projects/rccl/src/include/utils.h b/projects/rccl/src/include/utils.h index 5a6a588c43..0ed875c161 100644 --- a/projects/rccl/src/include/utils.h +++ b/projects/rccl/src/include/utils.h @@ -11,6 +11,7 @@ #include ncclResult_t getHostName(char* hostname, int maxlen); +uint64_t getnHash(const char* string, int n); uint64_t getHostHash(); uint64_t getPidHash(); diff --git a/projects/rccl/src/init.cu b/projects/rccl/src/init.cu index 7b13426844..3e9f01d397 100644 --- a/projects/rccl/src/init.cu +++ b/projects/rccl/src/init.cu @@ -292,9 +292,9 @@ static void showVersion() { } } -static ncclResult_t fillInfo(struct ncclInfo* info, int rank) { +static ncclResult_t fillInfo(struct ncclInfo* info, int rank, uint64_t commHash) { for (int t=0; ttinfo+t, rank)); + NCCLCHECK(ncclTransports[t].fillInfo(info->tinfo+t, rank, commHash)); } return ncclSuccess; } @@ -517,11 +517,13 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, ncclUniqueId* comm int rank = comm->rank; int nranks = comm->nRanks; void* commState; + uint64_t commHash = getnHash(commId->internal, NCCL_UNIQUE_ID_BYTES); + TRACE(NCCL_INIT, "comm %p, commHash %lu, rank %d nranks %d - BEGIN", comm, commHash, rank, nranks); NCCLCHECK(bootstrapInit(commId, rank, nranks, &commState)); struct ncclInfo* allInfo; NCCLCHECK(ncclCalloc(&allInfo, nranks)); - NCCLCHECK(fillInfo(allInfo+rank, rank)); + NCCLCHECK(fillInfo(allInfo+rank, rank, commHash)); NCCLCHECK(bootstrapAllGather(commState, allInfo, sizeof(struct ncclInfo))); int* connectTransport; @@ -741,7 +743,7 @@ static ncclResult_t initTransportsAll(struct ncclComm** comms, const int* devs, NCCLCHECK(ncclCalloc(&allInfo, nranks)); for (int rank=0; rankrank = rank; diff --git a/projects/rccl/src/transport/p2p.cu b/projects/rccl/src/transport/p2p.cu index 094511630e..38a5eb63f1 100644 --- a/projects/rccl/src/transport/p2p.cu +++ b/projects/rccl/src/transport/p2p.cu @@ -43,13 +43,13 @@ struct p2pConnectInfo { /* Fill information necessary to exchange between ranks to choose whether or not * to use this transport */ -ncclResult_t p2pFillInfo(ncclTinfo_t* opaqueInfo, int rank) { +ncclResult_t p2pFillInfo(ncclTinfo_t* opaqueInfo, int rank, uint64_t commHash) { struct p2pInfo* info = (struct p2pInfo*)opaqueInfo; static_assert(sizeof(struct p2pInfo) <= sizeof(ncclTinfo_t), "p2p Info too large"); info->rank = rank; CUDACHECK(hipGetDevice(&info->cudaDev)); - info->hostHash=getHostHash(); - info->pidHash=getPidHash(); + info->hostHash=getHostHash()+commHash; + info->pidHash=getPidHash()+commHash; // Get PCI Bus Id. We need to get the bus ID through CUDA first, since the // cudaDev is a CUDA runtime dev number which could be different from the diff --git a/projects/rccl/src/transport/shm.cu b/projects/rccl/src/transport/shm.cu index 557a32c86b..0ba168b2bf 100644 --- a/projects/rccl/src/transport/shm.cu +++ b/projects/rccl/src/transport/shm.cu @@ -54,13 +54,13 @@ struct shmRecvResources { /* Fill information necessary to exchange between ranks to choose whether or not * to use this transport */ -ncclResult_t shmFillInfo(ncclTinfo_t* opaqueInfo, int rank) { +ncclResult_t shmFillInfo(ncclTinfo_t* opaqueInfo, int rank, uint64_t commHash) { struct shmInfo* info = (struct shmInfo*)opaqueInfo; static_assert(sizeof(struct shmInfo) <= sizeof(ncclTinfo_t), "shm Info too large"); info->rank = rank; CUDACHECK(hipGetDevice(&info->cudaDev)); - info->hostHash=getHostHash(); - info->pidHash=getPidHash(); + info->hostHash=getHostHash()+commHash; + info->pidHash=getPidHash()+commHash; return ncclSuccess; }