From 8c564e9b57dc0fb913ce31de99c6ad556ae04fa7 Mon Sep 17 00:00:00 2001 From: Christian Sigg Date: Fri, 6 Dec 2019 18:14:55 +0100 Subject: [PATCH 1/2] 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 1ee2a73..15a9e7c 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/2] 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 89a97a3..87afb2f 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