diff --git a/makefiles/common.mk b/makefiles/common.mk index 37e81be204..2e448268a6 100644 --- a/makefiles/common.mk +++ b/makefiles/common.mk @@ -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. diff --git a/makefiles/version.mk b/makefiles/version.mk index efbdee7258..05abbc75c5 100644 --- a/makefiles/version.mk +++ b/makefiles/version.mk @@ -1,6 +1,6 @@ ##### version NCCL_MAJOR := 2 NCCL_MINOR := 5 -NCCL_PATCH := 6 +NCCL_PATCH := 7 NCCL_SUFFIX := PKG_REVISION := 1 diff --git a/pkg/debian/Makefile b/pkg/debian/Makefile index 2ce4390365..0494f3e032 100644 --- a/pkg/debian/Makefile +++ b/pkg/debian/Makefile @@ -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) diff --git a/src/collectives/device/primitives.h b/src/collectives/device/primitives.h index f2a175229f..8b2ac5a3ba 100644 --- a/src/collectives/device/primitives.h +++ b/src/collectives/device/primitives.h @@ -182,7 +182,7 @@ class ncclPrimitives { } template - 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 - 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 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); } diff --git a/src/graph/paths.cc b/src/graph/paths.cc index ce1772c39a..eba1964f44 100644 --- a/src/graph/paths.cc +++ b/src/graph/paths.cc @@ -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; tnodes[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; nnodes[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; tnodes[t].count; n++) { + struct ncclTopoNode* node = system->nodes[t].nodes+n; + for (int l=0; lnlinks; 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); diff --git a/src/graph/tuning.cc b/src/graph/tuning.cc index df2a901feb..957c63e973 100644 --- a/src/graph/tuning.cc +++ b/src/graph/tuning.cc @@ -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; arank = 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;