From 8c564e9b57dc0fb913ce31de99c6ad556ae04fa7 Mon Sep 17 00:00:00 2001 From: Christian Sigg Date: Fri, 6 Dec 2019 18:14:55 +0100 Subject: [PATCH 1/8] Fix clang build (#271) Clang doesn't understand `optimize("O0")`. It has `noopt`, which GCC doesn't understand. Wrap the difference in a macro. --- src/init.cc | 10 +++++++++- 1 file changed, 9 insertions(+), 1 deletion(-) diff --git a/src/init.cc b/src/init.cc index 1ee2a73d78..15a9e7caa5 100644 --- a/src/init.cc +++ b/src/init.cc @@ -124,10 +124,18 @@ 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__((noopt))) +#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; } +#undef NCCL_NO_OPTIMIZE + static ncclResult_t commFree(ncclComm_t comm) { if (comm == NULL) return ncclSuccess; From aa15dfb29c05adc4f2c895aa72f8873e125560fb Mon Sep 17 00:00:00 2001 From: Sylvain Jeaugey Date: Fri, 6 Dec 2019 09:55:54 -0800 Subject: [PATCH 2/8] Fix clang compilation --- src/graph/tuning.cc | 2 +- src/init.cc | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/src/graph/tuning.cc b/src/graph/tuning.cc index 89a97a3945..87afb2f598 100644 --- a/src/graph/tuning.cc +++ b/src/graph/tuning.cc @@ -192,7 +192,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 Date: Thu, 5 Dec 2019 18:34:55 -0800 Subject: [PATCH 3/8] 2.5.6-2 Fix PPC64 Debian packaging --- makefiles/version.mk | 2 +- pkg/debian/Makefile | 9 ++------- 2 files changed, 3 insertions(+), 8 deletions(-) diff --git a/makefiles/version.mk b/makefiles/version.mk index efbdee7258..80f6221bfd 100644 --- a/makefiles/version.mk +++ b/makefiles/version.mk @@ -3,4 +3,4 @@ NCCL_MAJOR := 2 NCCL_MINOR := 5 NCCL_PATCH := 6 NCCL_SUFFIX := -PKG_REVISION := 1 +PKG_REVISION := 2 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) From 3899f6e0f219b0cbf58537f791b0dd104b377750 Mon Sep 17 00:00:00 2001 From: Christian Sigg Date: Mon, 9 Dec 2019 18:31:13 +0100 Subject: [PATCH 4/8] Fix clang build (#274) The attribute is called `optnone`, not `noopt`. --- src/init.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/init.cc b/src/init.cc index 9da676a246..627f6c75a2 100644 --- a/src/init.cc +++ b/src/init.cc @@ -125,7 +125,7 @@ ncclResult_t ncclGetUniqueId(ncclUniqueId* out) { // Prevent compiler from optimizing out these operations #ifdef __clang__ -#define NCCL_NO_OPTIMIZE __attribute__((noopt)) +#define NCCL_NO_OPTIMIZE __attribute__((optnone)) #else #define NCCL_NO_OPTIMIZE __attribute__((optimize("O0"))) #endif From c7ba70ff90b357b40bf571ea3366d61e1249a0be Mon Sep 17 00:00:00 2001 From: Luke Yeager Date: Tue, 7 Jan 2020 13:29:42 -0800 Subject: [PATCH 5/8] [build] Allow setting CXXFLAGS on the command line --- makefiles/common.mk | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) 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. From 7a18fe07847300fbe7fec8d5512b3b44d8bc1716 Mon Sep 17 00:00:00 2001 From: Luke Yeager Date: Tue, 7 Jan 2020 13:29:57 -0800 Subject: [PATCH 6/8] [topology] remove NET links when trimming system This fixes a memory leak. --- src/graph/paths.cc | 23 +++++++++++++++++++++++ 1 file changed, 23 insertions(+) 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); From 3701130b3c1bcdb01c14b3cb70fe52498c1e82b7 Mon Sep 17 00:00:00 2001 From: Sylvain Jeaugey Date: Thu, 16 Jan 2020 15:40:57 -0800 Subject: [PATCH 7/8] 2.5.7-1 --- makefiles/version.mk | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/makefiles/version.mk b/makefiles/version.mk index 80f6221bfd..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 := 2 +PKG_REVISION := 1 From c38f174bd436031dbc79dce19ff969f377976a8a Mon Sep 17 00:00:00 2001 From: Sylvain Jeaugey Date: Wed, 12 Feb 2020 11:04:35 -0800 Subject: [PATCH 8/8] Fix Allgather operations above 4G with multiple GPUs per process. Fixes nccl-tests#37. Direct offsets were still on 32 bits in the low-level primitives. --- src/collectives/device/primitives.h | 18 +++++++++--------- 1 file changed, 9 insertions(+), 9 deletions(-) diff --git a/src/collectives/device/primitives.h b/src/collectives/device/primitives.h index aa3d20de56..b624359e6d 100644 --- a/src/collectives/device/primitives.h +++ b/src/collectives/device/primitives.h @@ -1,5 +1,5 @@ /************************************************************************* - * Copyright (c) 2016-2019, NVIDIA CORPORATION. All rights reserved. + * Copyright (c) 2016-2020, NVIDIA CORPORATION. All rights reserved. * * See LICENSE.txt for license information ************************************************************************/ @@ -143,12 +143,12 @@ class ncclPrimitives { } template - inline __device__ const T* directRecvPtr(int i, int directOffset) { + inline __device__ const T* directRecvPtr(int i, ssize_t directOffset) { return DIRECTRECV && recvDirectBuff[i] ? recvDirectBuff[i]+directOffset : recvPtr(i); } template - inline __device__ T* directSendPtr(int i, int directOffset) { + inline __device__ T* directSendPtr(int i, ssize_t directOffset) { return DIRECTSEND && sendDirectBuff[i] ? sendDirectBuff[i]+directOffset : sendPtr(i); } @@ -164,7 +164,7 @@ class ncclPrimitives { 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); @@ -310,7 +310,7 @@ class ncclPrimitives { 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); } @@ -319,7 +319,7 @@ class ncclPrimitives { 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); } @@ -328,7 +328,7 @@ class ncclPrimitives { 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); } @@ -337,7 +337,7 @@ class ncclPrimitives { 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); } @@ -356,7 +356,7 @@ class ncclPrimitives { 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); }