AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge remote-tracking branch 'remotes/nccl/master'

Browse Source
This commit is contained in:
Wenkai Du
2020-02-27 11:45:50 -08:00
parent 63af7ab5cb c38f174bd4
commit 8e73a2ad60
7 changed files with 48 additions and 21 deletions
Download Patch File Download Diff File Expand all files Collapse all files
makefiles/common.mk
+4 -3
View File
@@ -42,9 +42,10 @@ else
endif
#$(info NVCC_GENCODE is ${NVCC_GENCODE})
CXXFLAGS := -DCUDA_MAJOR=$(CUDA_MAJOR) -DCUDA_MINOR=$(CUDA_MINOR) -fPIC -fvisibility=hidden
CXXFLAGS += -Wall -Wno-unused-function -Wno-sign-compare -std=c++11 -Wvla
CXXFLAGS += -I $(CUDA_INC)
CXXFLAGS := -DCUDA_MAJOR=$(CUDA_MAJOR) -DCUDA_MINOR=$(CUDA_MINOR) -fPIC -fvisibility=hidden \
-Wall -Wno-unused-function -Wno-sign-compare -std=c++11 -Wvla \
-I $(CUDA_INC) \
$(CXXFLAGS)
# Maxrregcount needs to be set accordingly to NCCL_MAX_NTHREADS (otherwise it will cause kernel launch errors)
# 512 : 120, 640 : 96, 768 : 80, 1024 : 60
# We would not have to set this if we used __launch_bounds__, but this only works on kernels, not on functions.
makefiles/version.mk
+1 -1
View File
@@ -1,6 +1,6 @@
##### version
NCCL_MAJOR := 2
NCCL_MINOR := 5
NCCL_PATCH := 6
NCCL_PATCH := 7
NCCL_SUFFIX :=
PKG_REVISION := 1
pkg/debian/Makefile
+2 -7
View File
@@ -16,13 +16,8 @@ DEBFILES := compat copyright libnccl-dev.install rules $(DEBGEN)
DEBTARGETS := $(patsubst %, $(DEBPREPDIR)/%, $(DEBFILES))
PKG_TIMESTAMP := $(shell date -R)
ARCH := $(shell uname -m)
PKG_ARCH ?= $(shell uname -m | sed -e "s/x86_64/amd64/g" | sed -e "s/ppc64le/ppc64el/g"| sed -e "s/aarch64/arm64/g")
PKG_MULTIARCH ?= $(shell $(CXX) -print-multiarch)
ifeq ($(PKG_MULTIARCH),)
# Hardwire the PKG_MULTIARCH directory as the RHEL6 distribution agnostic compiler (gcc 4.8.3) doesn't set it
PKG_MULTIARCH := $(ARCH)-linux-gnu
endif
PKG_ARCH ?= $(shell dpkg-architecture -qDEB_HOST_ARCH)
PKG_MULTIARCH ?= $(shell dpkg-architecture -qDEB_HOST_MULTIARCH)
prep : $(DEBTARGETS)
$(MAKE) -C ../.. lic BUILDDIR=$(BUILDDIR)
src/collectives/device/primitives.h
+8 -8
View File
@@ -182,7 +182,7 @@ class ncclPrimitives {
}
template <int DIRECTRECV>
inline __device__ const T* directRecvPtr(int i, int directOffset) {
inline __device__ const T* directRecvPtr(int i, ssize_t directOffset) {
#if defined(RCCL_USE_DIRECT_BUFFER)
return DIRECTRECV && r.recvDirectBuff[i] ? r.recvDirectBuff[i]+directOffset : recvPtr(i);
#else
@@ -191,7 +191,7 @@ class ncclPrimitives {
}
template <int DIRECTSEND>
inline __device__ T* directSendPtr(int i, int directOffset) {
inline __device__ T* directSendPtr(int i, ssize_t directOffset) {
#if defined(RCCL_USE_DIRECT_BUFFER)
return DIRECTSEND && s.sendDirectBuff[i] ? s.sendDirectBuff[i]+directOffset : sendPtr(i);
#else
@@ -219,7 +219,7 @@ inline __device__ int directSendInc(int i, int directInc, int sliceInc) {
template <int DIRECTRECV, int DIRECTSEND, int RECV, int SEND, int SRC, int DST>
inline __device__ void
GenericOp(const T* srcPtr, T* dstPtr, int nelem, int directOffset) {
GenericOp(const T* srcPtr, T* dstPtr, int nelem, ssize_t directOffset) {
int offset = 0;
int sliceSize = stepSize*SLICESTEPS;
int dataSize = max(DIVUP(nelem, 16*SLICESPERCHUNK)*16, sliceSize/32);
@@ -386,7 +386,7 @@ inline __device__ int directSendInc(int i, int directInc, int sliceInc) {
GenericOp<0, 0, 0, 1, 1, 0>(src, NULL, nelem, 0);
}
__device__ __forceinline__ void
directSend(const T* src, int directOffset, int nelem) {
directSend(const T* src, ssize_t directOffset, int nelem) {
GenericOp<0, 1, 0, 1, 1, 0>(src, NULL, nelem, directOffset);
}
@@ -395,7 +395,7 @@ inline __device__ int directSendInc(int i, int directInc, int sliceInc) {
GenericOp<0, 0, 1, 0, 0, 1>(NULL, dst, nelem, 0);
}
__device__ __forceinline__ void
directRecv(T* dst, int directOffset, int nelem) {
directRecv(T* dst, ssize_t directOffset, int nelem) {
GenericOp<1, 0, 1, 0, 0, 1>(NULL, dst, nelem, directOffset);
}
@@ -404,7 +404,7 @@ inline __device__ int directSendInc(int i, int directInc, int sliceInc) {
GenericOp<0, 0, 0, 1, 1, 1>(src, dst, nelem, 0);
}
__device__ __forceinline__ void
directCopySend(const T* src, T* dst, int directOffset, int nelem) {
directCopySend(const T* src, T* dst, ssize_t directOffset, int nelem) {
GenericOp<0, 1, 0, 1, 1, 1>(src, dst, nelem, directOffset);
}
@@ -413,7 +413,7 @@ inline __device__ int directSendInc(int i, int directInc, int sliceInc) {
GenericOp<0, 0, 1, 1, 0, 1>(NULL, dst, nelem, 0);
}
__device__ __forceinline__ void
directRecvCopySend(T* dst, int directOffset, int nelem) {
directRecvCopySend(T* dst, ssize_t directOffset, int nelem) {
GenericOp<1, 1, 1, 1, 0, 1>(NULL, dst, nelem, directOffset);
}
@@ -432,7 +432,7 @@ inline __device__ int directSendInc(int i, int directInc, int sliceInc) {
GenericOp<0, 0, 1, 1, 1, 1>(src, dst, nelem, 0);
}
__device__ __forceinline__ void
directRecvReduceCopySend(const T* src, T* dst, int directOffset, int nelem) {
directRecvReduceCopySend(const T* src, T* dst, ssize_t directOffset, int nelem) {
// Direct is only for the send part
GenericOp<0, 1, 1, 1, 1, 1>(src, dst, nelem, directOffset);
}
src/graph/paths.cc
+23
View File
@@ -179,11 +179,18 @@ static ncclResult_t addCpuStep(struct ncclTopoSystem* system, int c, int t1, int
// Remove/free paths for a given type
static void ncclTopoRemovePathType(struct ncclTopoSystem* system, int nodeType) {
for (int t=0; t<NCCL_TOPO_NODE_TYPES; t++) {
// Remove links _to_ the given type
for (int n=0; n<system->nodes[t].count; n++) {
struct ncclTopoNode* node = system->nodes[t].nodes+n;
free(node->paths[nodeType]);
node->paths[nodeType] = NULL;
}
// Remove links _from_ the given type
for (int n=0; n<system->nodes[nodeType].count; n++) {
struct ncclTopoNode* node = system->nodes[nodeType].nodes+n;
free(node->paths[t]);
node->paths[t] = NULL;
}
}
}
@@ -309,6 +316,22 @@ ncclResult_t ncclTopoTrimSystem(struct ncclTopoSystem* system, struct ncclComm*
// Trim network
ncclTopoRemovePathType(system, NET);
system->nodes[NET].count = 0;
for (int t=0; t<NCCL_TOPO_NODE_TYPES; t++) {
for (int n=0; n<system->nodes[t].count; n++) {
struct ncclTopoNode* node = system->nodes[t].nodes+n;
for (int l=0; l<node->nlinks; l++) {
struct ncclTopoLink* link = &(node->links[l]);
if (link->remNode->type == NET) {
// Remove the link
for (int i=l; i<(node->nlinks-1); i++) {
memcpy(&(node->links[i]), &(node->links[i+1]), sizeof(ncclTopoLink));
}
node->nlinks--;
l--; // revisit the same value of "l" for the next iteration, since we edited the list in the middle of the loop
}
}
}
}
}
free(domains);
free(ids);
src/graph/tuning.cc
+1 -1
View File
@@ -193,7 +193,7 @@ ncclResult_t ncclSetThresholds(struct ncclComm* comm, int minCompCap, int maxCom
// Override defaults with user env
char* str = getenv("NCCL_THREAD_THRESHOLDS");
if (str) {
ssize_t t[NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { -2 };
ssize_t t[NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = {{ -2, -2, -2 }, { -2, -2, -2}};
sscanf(str, "%ld %ld %ld %ld %ld %ld", t[0], t[0]+1, t[0]+2, t[1], t[1]+1, t[1]+2);
for (int a=0; a<NCCL_NUM_ALGORITHMS; a++) {
for (int p=0; p<NCCL_NUM_PROTOCOLS; p++) {
src/init.cc
+9 -1
View File
@@ -127,7 +127,13 @@ ncclResult_t ncclGetUniqueId(ncclUniqueId* out) {
}
// Prevent compiler from optimizing out these operations
void __attribute__((optimize("O0"))) commPoison(ncclComm_t comm) {
#ifdef __clang__
#define NCCL_NO_OPTIMIZE __attribute__((optnone))
#else
#define NCCL_NO_OPTIMIZE __attribute__((optimize("O0")))
#endif
void NCCL_NO_OPTIMIZE commPoison(ncclComm_t comm) {
comm->rank = comm->cudaDev = comm->busId = comm->nRanks = -1;
}
@@ -179,6 +185,8 @@ void *ncclCommThreadMain(void *arg) {
}
#endif
#undef NCCL_NO_OPTIMIZE
static ncclResult_t commFree(ncclComm_t comm) {
if (comm == NULL)
return ncclSuccess;