diff --git a/projects/rccl/CMakeLists.txt b/projects/rccl/CMakeLists.txt index b57363b1bb..61b2503862 100644 --- a/projects/rccl/CMakeLists.txt +++ b/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") diff --git a/projects/rccl/cmake/CheckSymbolExistsNoWarn.cmake b/projects/rccl/cmake/CheckSymbolExistsNoWarn.cmake new file mode 100644 index 0000000000..b478f599c1 --- /dev/null +++ b/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() diff --git a/projects/rccl/src/graph/rome_models.cc b/projects/rccl/src/graph/rome_models.cc index ac805fe518..01a7e67ae4 100644 --- a/projects/rccl/src/graph/rome_models.cc +++ b/projects/rccl/src/graph/rome_models.cc @@ -26,6 +26,7 @@ THE SOFTWARE. #include #include #include +#include #include #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 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 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 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; } diff --git a/projects/rccl/test/AllReduceTests.cpp b/projects/rccl/test/AllReduceTests.cpp index 5d7d8cdfc0..a82a9ac19d 100644 --- a/projects/rccl/test/AllReduceTests.cpp +++ b/projects/rccl/test/AllReduceTests.cpp @@ -116,7 +116,7 @@ namespace RcclUnitTesting std::vector const inPlaceList = {false}; std::vector const managedMemList = {false}; std::vector const useHipGraphList = {false, true}; - std::vector const channelList = {"84", "112"}; + std::vector const channelList = {"84", "112"}; bool const enableSweep = false; for (auto channel : channelList) { setenv("NCCL_MIN_NCHANNELS", channel, 1); diff --git a/projects/rccl/test/AllToAllTests.cpp b/projects/rccl/test/AllToAllTests.cpp index 926cf41488..4298b2a2f4 100644 --- a/projects/rccl/test/AllToAllTests.cpp +++ b/projects/rccl/test/AllToAllTests.cpp @@ -100,7 +100,7 @@ namespace RcclUnitTesting std::vector const inPlaceList = {false}; std::vector const managedMemList = {false}; std::vector const useHipGraphList = {false, true}; - std::vector const channelList = {"112"}; + std::vector const channelList = {"112"}; bool const enableSweep = false; for (auto channel : channelList) { setenv("NCCL_MIN_NCHANNELS", channel, 1); diff --git a/projects/rccl/test/StandaloneTests.cpp b/projects/rccl/test/StandaloneTests.cpp index 97249774ee..f699aeb4c9 100644 --- a/projects/rccl/test/StandaloneTests.cpp +++ b/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 comms(numRanks); + std::vector gpuInput(numRanks); + std::vector gpuOutput(numRanks); + std::vector 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(Clock::now() - start).count(); // Check results