Hide or fix all build warnings (#1331)

* Changing C-strings to be const.

* Changed variable-length arrays to std::vector to avoid warnings. VLA is a compiler extension.

* Changed `#define` inside functions into `constexpr int` to preserve scoping and avoid macro redefinition warnings.

* Disabled warnings for modifying `CMAKE_CXX_FLAGS` caused by `check_symbol_exists`, which temporarily modifies the flag to do a compile check.

* Fixed VLA in rccl UT.

[ROCm/rccl commit: 1c45962273]

projects/rccl/CMakeLists.txt

@@ -55,6 +55,7 @@ set(DEFAULT_GPUS
 include(CheckIncludeFiles)
 include(CheckSymbolExists)
 include(cmake/Dependencies.cmake) # GTest, rocm-cmake, rocm_local_targets
+include(cmake/CheckSymbolExistsNoWarn.cmake)
 
 list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
 

projects/rccl/cmake/CheckSymbolExistsNoWarn.cmake

@@ -0,0 +1,40 @@
+# MIT License
+#
+# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+# copies of the Software, and to permit persons to whom the Software is
+# furnished to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in all
+# copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+# SOFTWARE.
+
+# These overrides are due to CMake CHECK_SYMBOL_EXISTS modifying CMAKE_CXX_FLAGS to do a test compile,
+# while ROCMChecks gives a warning if this variable is modified manually without a target.
+
+# We now choose to disable ROCMChecks for this one case.
+
+set(DISABLE_ROCM_CHECK OFF)
+
+function(rocm_check_toolchain_var var access value list_file)
+  if(NOT DISABLE_ROCM_CHECK)
+    _rocm_check_toolchain_var("${var}" "${access}" "${value}" "${list_file}")
+  endif()
+endfunction()
+
+macro(CHECK_SYMBOL_EXISTS)
+  set(DISABLE_ROCM_CHECK ON)
+  _check_symbol_exists(${ARGN})
+  set(DISABLE_ROCM_CHECK OFF)
+endmacro()

projects/rccl/src/graph/rome_models.cc

@@ -26,6 +26,7 @@ THE SOFTWARE.
 #include <math.h>
 #include <sys/time.h>
 #include <algorithm>
+#include <vector>
 #include <string.h>
 #include "rome_models.h"
 
@@ -1746,17 +1747,17 @@ int checkAlltoallWidth(struct rcclRomeModel *romeTopo) {
 }
 
 ncclResult_t parseA2a8P(struct ncclTopoSystem* system, struct ncclTopoGraph* graph, const char *ringBase) {
-  #define NUMA_CPUS 2
-  #define NUMA_GPUS 4
-  #define NUMA_PERMUTE_COUNT 24
-  #define TOTAL_PERMUTE_COUNT (NUMA_PERMUTE_COUNT*NUMA_PERMUTE_COUNT)
+  constexpr int NUMA_CPUS = 2;
+  constexpr int NUMA_GPUS = 4;
+  constexpr int NUMA_PERMUTE_COUNT = 24;
+  constexpr int TOTAL_PERMUTE_COUNT = (NUMA_PERMUTE_COUNT*NUMA_PERMUTE_COUNT);
 
   static char ringRemap[256];
   int i;
 
-  int ngpus = system->nodes[GPU].count;
-  int ncpus = system->nodes[CPU].count;
-  int nnets = system->nodes[NET].count;
+  const int ngpus = system->nodes[GPU].count;
+  const int ncpus = system->nodes[CPU].count;
+  const int nnets = system->nodes[NET].count;
 
   // number of GPUs and NICs on each numa node is used as first screening pattern
   struct rcclRomeModel romeTopo;
@@ -1775,11 +1776,12 @@ ncclResult_t parseA2a8P(struct ncclTopoSystem* system, struct ncclTopoGraph* gra
   int *all_gpu_permutations = (int *)malloc(TOTAL_PERMUTE_COUNT*NUMA_CPUS*NUMA_GPUS*sizeof(int));
   struct timeval tvs, tve;
   gettimeofday(&tvs, NULL);
+  std::vector<int> r(ngpus), g(ngpus);
   for (i = 0; i < sizeof(romeTopoModels)/sizeof(romeTopoModels[0]); i++) {
     if (romeTopo.nCpus != romeTopoModels[i].nCpus || romeTopo.nGpus != romeTopoModels[i].nGpus ||
       romeTopo.nNics != romeTopoModels[i].nNics || romeTopo.nLinks != romeTopoModels[i].nLinks) continue;
     if (strcmp(romeTopoModels[i].pattern, pattern)) continue;
-    int j, r[ngpus], g[ngpus];
+    int j;
     int numa_gpu_permutations[NUMA_CPUS][NUMA_PERMUTE_COUNT][NUMA_GPUS];
     if (isAlltoall != checkAlltoallWidth(romeTopoModels+i))
       continue;
@@ -1798,12 +1800,12 @@ ncclResult_t parseA2a8P(struct ncclTopoSystem* system, struct ncclTopoGraph* gra
         if (romeTopo.gpuNuma[k] != j) continue;
         g[(2+cnt++)%ngpusPerNuma] = k;
       }
-      std::sort(g, g+ngpusPerNuma);
+      std::sort(g.begin(), g.begin()+ngpusPerNuma);
       do {
         for (int n = 0; n < ngpusPerNuma; n++)
           numa_gpu_permutations[j][npermute][n] = g[n];
         npermute++;
-      } while (std::next_permutation(g, g+ngpusPerNuma));
+      } while (std::next_permutation(g.begin(), g.begin()+ngpusPerNuma));
       if (npermute != NUMA_PERMUTE_COUNT) break;
     }
     if (j < ncpus) continue;
@@ -2108,17 +2110,17 @@ ncclResult_t parseRome4P2H(struct ncclTopoSystem* system, struct ncclTopoGraph*
 }
 
 ncclResult_t parse1H16P(struct ncclTopoSystem* system, struct ncclTopoGraph* graph) {
-  #define NUMA_CPUS 4
-  #define NUMA_GPUS 4
-  #define NUMA_PERMUTE_COUNT 24
-  #define TOTAL_PERMUTE_COUNT (NUMA_PERMUTE_COUNT*NUMA_PERMUTE_COUNT*NUMA_PERMUTE_COUNT*NUMA_PERMUTE_COUNT)
+  constexpr int NUMA_CPUS = 4;
+  constexpr int NUMA_GPUS = 4;
+  constexpr int NUMA_PERMUTE_COUNT = 24;
+  constexpr int TOTAL_PERMUTE_COUNT = (NUMA_PERMUTE_COUNT*NUMA_PERMUTE_COUNT*NUMA_PERMUTE_COUNT*NUMA_PERMUTE_COUNT);
 
   static char ringRemap[256];
   int i;
 
-  int ngpus = system->nodes[GPU].count;
-  int ncpus = system->nodes[CPU].count;
-  int nnets = system->nodes[NET].count;
+  const int ngpus = system->nodes[GPU].count;
+  const int ncpus = system->nodes[CPU].count;
+  const int nnets = system->nodes[NET].count;
 
   // only valid on Rome
   int arch, vendor, model;
@@ -2139,11 +2141,12 @@ ncclResult_t parse1H16P(struct ncclTopoSystem* system, struct ncclTopoGraph* gra
   int *all_gpu_permutations = (int *)malloc(TOTAL_PERMUTE_COUNT*NUMA_CPUS*NUMA_GPUS*sizeof(int));
   struct timeval tvs, tve;
   gettimeofday(&tvs, NULL);
+  std::vector<int> r(ngpus), g(ngpus);
   for (i = 0; i < sizeof(romeTopoModels)/sizeof(romeTopoModels[0]); i++) {
     if (romeTopo.nCpus != romeTopoModels[i].nCpus || romeTopo.nGpus != romeTopoModels[i].nGpus ||
       romeTopo.nNics != romeTopoModels[i].nNics || romeTopo.nLinks != romeTopoModels[i].nLinks) continue;
     if (strcmp(romeTopoModels[i].pattern, pattern)) continue;
-    int j, r[ngpus], g[ngpus];
+    int j;
     int numa_gpu_permutations[NUMA_CPUS][NUMA_PERMUTE_COUNT][NUMA_GPUS];
     // permute GPUs for each CPU NUMA nodes
     for (j = 0; j < ncpus; j++) {
@@ -2160,12 +2163,12 @@ ncclResult_t parse1H16P(struct ncclTopoSystem* system, struct ncclTopoGraph* gra
         if (romeTopo.gpuNuma[k] != j) continue;
         g[(2+cnt++)%ngpusPerNuma] = k;
       }
-      std::sort(g, g+ngpusPerNuma);
+      std::sort(g.begin(), g.begin()+ngpusPerNuma);
       do {
         for (int n = 0; n < ngpusPerNuma; n++)
           numa_gpu_permutations[j][npermute][n] = g[n];
         npermute++;
-      } while (std::next_permutation(g, g+ngpusPerNuma));
+      } while (std::next_permutation(g.begin(), g.begin()+ngpusPerNuma));
       if (npermute != NUMA_PERMUTE_COUNT) break;
     }
     if (j < ncpus) continue;
@@ -2245,13 +2248,13 @@ ncclResult_t parse1H16P(struct ncclTopoSystem* system, struct ncclTopoGraph* gra
 }
 
 ncclResult_t parse4H4P(struct ncclTopoSystem* system, struct ncclTopoGraph* graph) {
-  #define NUM_HIVES 4
-  #define HIVE_GPUS 4
+  constexpr int NUM_HIVES = 4;
+  constexpr int HIVE_GPUS = 4;
 
   static char ringRemap[256];
 
-  int ngpus = system->nodes[GPU].count;
-  int nnets = system->nodes[NET].count;
+  const int ngpus = system->nodes[GPU].count;
+  const int nnets = system->nodes[NET].count;
 
   // only valid on Rome
   int arch, vendor, model;
@@ -2267,7 +2270,7 @@ ncclResult_t parse4H4P(struct ncclTopoSystem* system, struct ncclTopoGraph* grap
   // only match for system with 16 GPUs
   if (ngpus != NUM_HIVES*HIVE_GPUS || nnets != NUM_HIVES*HIVE_GPUS) return ncclSuccess;
 
-  int g_hives[ngpus], n_hives[nnets];
+  std::vector<int> g_hives(ngpus), n_hives(nnets);
   int ng_hives[NUM_HIVES];
 
   // try to sort GPUs into hives
@@ -2338,6 +2341,6 @@ ncclResult_t parse4H4P(struct ncclTopoSystem* system, struct ncclTopoGraph* grap
     system->type |= RCCL_TOPO_4P2H_ROME;
   parseOptions(system, rome_model_68.options);
   // create 4P4H based on reference and remapped ids
-  NCCLCHECK(parseGraph(rome_model_68.ringBase, system, graph, g_hives, n_hives, false));
+  NCCLCHECK(parseGraph(rome_model_68.ringBase, system, graph, g_hives.data(), n_hives.data(), false));
   return ncclSuccess;
 }

projects/rccl/test/AllReduceTests.cpp

@@ -116,7 +116,7 @@ namespace RcclUnitTesting
         std::vector<bool>           const inPlaceList     = {false};
         std::vector<bool>           const managedMemList  = {false};
         std::vector<bool>           const useHipGraphList = {false, true};
-        std::vector<char *>         const channelList     = {"84", "112"};
+        std::vector<const char *>   const channelList     = {"84", "112"};
         bool                        const enableSweep     = false; 
         for (auto channel : channelList) {
           setenv("NCCL_MIN_NCHANNELS", channel, 1);

projects/rccl/test/AllToAllTests.cpp

@@ -100,7 +100,7 @@ namespace RcclUnitTesting
         std::vector<bool>           const inPlaceList     = {false};
         std::vector<bool>           const managedMemList  = {false};
         std::vector<bool>           const useHipGraphList = {false, true};
-        std::vector<char *>         const channelList     = {"112"}; 
+        std::vector<const char *>   const channelList     = {"112"}; 
         bool                        const enableSweep     = false;
         for (auto channel : channelList) {
           setenv("NCCL_MIN_NCHANNELS", channel, 1);

projects/rccl/test/StandaloneTests.cpp

@@ -171,8 +171,11 @@ namespace RcclUnitTesting
     hipDeviceProp_t devProp;
     HIPCALL(hipGetDeviceProperties(&devProp, 0));
     // Initialize RCCL
-    int numRanks = 2;
+    constexpr int numRanks = 2;
     std::vector<ncclComm_t> comms(numRanks);
+    std::vector<int*> gpuInput(numRanks);
+    std::vector<int*> gpuOutput(numRanks);
+    std::vector<hipStream_t> stream(numRanks);
 
     char *proto = std::getenv("NCCL_PROTO");
     const char* protocolList[3] = {"LL", "LL128", "Simple"};
@@ -198,10 +201,6 @@ namespace RcclUnitTesting
       }
 
       // Prepare GPU data arrays
-      int* gpuInput[numRanks];
-      int* gpuOutput[numRanks];
-      hipStream_t stream[numRanks];
-
       for (int rank = 0; rank < numRanks; rank++) {
         HIPCALL(hipSetDevice(rank));
         HIPCALL(hipStreamCreate(&stream[rank]));
@@ -218,7 +217,7 @@ namespace RcclUnitTesting
           HIPCALL(hipSetDevice(rank));
           HIPCALL(hipMemset(gpuOutput[rank], 0, N * sizeof(int)));
           HIPCALL(hipDeviceSynchronize());
-	}
+        }
 
         // Initiate the allreduce
         NCCLCHECK(ncclGroupStart());
@@ -233,7 +232,7 @@ namespace RcclUnitTesting
           HIPCALL(hipStreamSynchronize(stream[rank]));
         }
 
-	if (iter >= 0)
+        if (iter >= 0)
           usElapsed += duration_cast<microseconds>(Clock::now() - start).count();
 
         // Check results