Merge branch 'master' into HEAD

[ROCm/rccl commit: 44b5652617]

projects/rccl/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++) {

projects/rccl/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;