yezhengmao/nccl
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge branch 'master' into HEAD

Browse Source
This commit is contained in:
Ke Wen 2019-12-06 18:28:11 -08:00
commit 44b5652617
2 changed files with 10 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
src/graph/tuning.cc
View File

@ -192,7 +192,7 @@ ncclResult_t ncclSetThresholds(struct ncclComm* comm, int minCompCap, int maxCom
// Override defaults with user env // Override defaults with user env
char* str = getenv("NCCL_THREAD_THRESHOLDS"); char* str = getenv("NCCL_THREAD_THRESHOLDS");
if (str) { 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); 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 a=0; a<NCCL_NUM_ALGORITHMS; a++) {
for (int p=0; p<NCCL_NUM_PROTOCOLS; p++) { for (int p=0; p<NCCL_NUM_PROTOCOLS; p++) {

10
src/init.cc
View File

@ -124,10 +124,18 @@ ncclResult_t ncclGetUniqueId(ncclUniqueId* out) {
} }
// Prevent compiler from optimizing out these operations // 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; comm->rank = comm->cudaDev = comm->busId = comm->nRanks = -1;
} }
#undef NCCL_NO_OPTIMIZE
static ncclResult_t commFree(ncclComm_t comm) { static ncclResult_t commFree(ncclComm_t comm) {
if (comm == NULL) if (comm == NULL)
return ncclSuccess; return ncclSuccess;