diff options
| author | Ke Wen <kwen@nvidia.com> | 2019-12-07 05:28:11 +0300 |
|---|---|---|
| committer | Ke Wen <kwen@nvidia.com> | 2019-12-07 05:28:11 +0300 |
| commit | 44b56526179c6fb0c72350d378a8ff9184774ef9 (patch) | |
| tree | ea3859bb660c32e064a036d526269adb9657f017 | |
| parent | 6bb953d4e6b7e6856683a8cc83b9626b3bf421bb (diff) | |
| parent | aa15dfb29c05adc4f2c895aa72f8873e125560fb (diff) | |
Merge branch 'master' into HEAD
| -rw-r--r-- | src/graph/tuning.cc | 2 | ||||
| -rw-r--r-- | src/init.cc | 10 |
2 files changed, 10 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<NCCL_NUM_ALGORITHMS; a++) { for (int p=0; p<NCCL_NUM_PROTOCOLS; p++) { diff --git a/src/init.cc b/src/init.cc index 1ee2a73..9da676a 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; |