From 20975921dde19e7c5c28df4ca562c8b0abebd405 Mon Sep 17 00:00:00 2001 From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com> Date: Tue, 2 Jul 2019 09:27:16 -0700 Subject: [PATCH] Fix share memory collision in multi-communicator case. (#93) Current SHM object name would only use pidHash and ranks as identification, which would collide each other when program runs with multiple communicators. Here we added commId info into pidHash, it makes 'pidHash'es of different communicators keeping in same process will be distincted with each other. Ported from original commit: https://github.com/lowintelligence/nccl/commits/shm [ROCm/rccl commit: 949d680e4907927976b008830c1e81179a6ba2d6] --- projects/rccl/src/include/transport.h | 2 +- projects/rccl/src/include/utils.h | 1 + projects/rccl/src/init.cu | 10 ++++++---- projects/rccl/src/misc/utils.cu | 9 +++++++++ projects/rccl/src/transport/net.cu | 2 +- projects/rccl/src/transport/p2p.cu | 6 +++--- projects/rccl/src/transport/shm.cu | 6 +++--- 7 files changed, 24 insertions(+), 12 deletions(-) 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; }