From 8145873d12d20dbb3248eeb68e926d5f24c84b6b Mon Sep 17 00:00:00 2001 From: romerojosh Date: Thu, 16 Jul 2026 18:07:30 +0000 Subject: [PATCH 1/2] Add autotuning candidate controls Signed-off-by: romerojosh --- docs/api/f_api.rst | 1 + docs/autotuning.rst | 9 +- docs/env_vars.rst | 40 ++++ include/cudecomp.h | 3 +- include/internal/autotune.h | 12 +- src/autotune.cc | 324 ++++++++++++++++++++++--------- src/cudecomp.cc | 40 ++-- src/cudecomp_m.cuf | 3 +- tests/ctest/api_tests.cc | 185 ++++++++++++++++++ tests/ctest/fortran_api_test.f90 | 4 + 10 files changed, 508 insertions(+), 113 deletions(-) diff --git a/docs/api/f_api.rst b/docs/api/f_api.rst index 6395c7de..54875a71 100644 --- a/docs/api/f_api.rst +++ b/docs/api/f_api.rst @@ -70,6 +70,7 @@ _________________________________ :f cudecompAutotuneGridMode grid_mode: which communication (transpose/halo) to use to autotune process grid (default: CUDECOMP_AUTOTUNE_GRID_TRANSPOSE) :f cudecompDataType dtype: datatype to use during autotuning (default: CUDECOMP_DOUBLE) :f logical allow_uneven_distributions: flag to control whether autotuning allows process grids that result in uneven distributions of elements across processes (default: true) + :f logical disable_mpi_backends: flag to disable MPI backend options during autotuning (default: false) :f logical disable_nccl_backends: flag to disable NCCL backend options during autotuning (default: false) :f logical disable_nvshmem_backends: flag to disable NVSHMEM backend options during autotuning (default: false) :f real(c_double) skip_threshold: threshold used to skip testing slow configurations; skip configuration if :code:`skip_threshold * t > t_best`, where :code:`t` is the duration of the first timed trial for the configuration and :code:`t_best` is the average trial time of the current best configuration (default: 0.0) diff --git a/docs/autotuning.rst b/docs/autotuning.rst index 0dfd7571..5e571daa 100644 --- a/docs/autotuning.rst +++ b/docs/autotuning.rst @@ -117,19 +117,22 @@ The :code:`dtype` entry in the options struct controls which data type cuDecomp options%dtype = CUDECOMP_DOUBLE -The :code:`disable_nccl_backends` and :code:`disable_nvshmem_backends` entries are boolean flags controlling whether -the autotuner will test transpose and halo communication backends using the NCCL or NVSHMEM libraries respectively. -By default, these flags are set to false and NCCL and NVSHMEM backends are enabled. +The :code:`disable_mpi_backends`, :code:`disable_nccl_backends`, and :code:`disable_nvshmem_backends` entries are +boolean flags controlling whether the autotuner will test transpose and halo communication backends using the MPI, +NCCL, or NVSHMEM libraries respectively. These exclusions are applied after any backend environment-variable filter. +By default, these flags are set to false and all three backend families are enabled. .. tabs:: .. code-tab:: c++ + options.disable_mpi_backends = false; options.disable_nccl_backends = false; options.disable_nvshmem_backends = false; .. code-tab:: fortran + options%disable_mpi_backends = .false. options%disable_nccl_backends = .false. options%disable_nvshmem_backends = .false. diff --git a/docs/env_vars.rst b/docs/env_vars.rst index e4fae790..023d9bb6 100644 --- a/docs/env_vars.rst +++ b/docs/env_vars.rst @@ -104,3 +104,43 @@ CUDECOMP_USE_COL_MAJOR_RANK_ORDER When :code:`rank_order` is left at :code:`CUDECOMP_RANK_ORDER_DEFAULT`, this environment variable is still honored for backward compatibility. Setting this variable to :code:`1` enables column-major rank assignment; otherwise rank assignment is row-major. Explicit :code:`rank_order` settings take precedence over this environment variable. + +CUDECOMP_AUTOTUNE_TRANSPOSE_BACKENDS +------------------------------------ +(since v0.7.0) + +:code:`CUDECOMP_AUTOTUNE_TRANSPOSE_BACKENDS` restricts the candidates used when transpose communication backends are +autotuned. Its value is a comma-separated inclusion list of uppercase backend names. Prefixing the complete value with +:code:`^` changes it to an exclusion list. Accepted names are :code:`MPI_P2P`, :code:`MPI_P2P_PL`, :code:`MPI_A2A`, +:code:`NCCL`, :code:`NCCL_PL`, :code:`NVSHMEM`, :code:`NVSHMEM_PL`, and :code:`NVSHMEM_SM`. + +Default setting is unset (no additional backend filtering). + +CUDECOMP_AUTOTUNE_HALO_BACKENDS +------------------------------- +(since v0.7.0) + +:code:`CUDECOMP_AUTOTUNE_HALO_BACKENDS` restricts the candidates used when halo communication backends are autotuned. +Its value is a comma-separated inclusion list of uppercase backend names. Prefixing the complete value with :code:`^` +changes it to an exclusion list. Accepted names are :code:`MPI`, :code:`MPI_BLOCKING`, :code:`NCCL`, :code:`NVSHMEM`, +and :code:`NVSHMEM_BLOCKING`. + +Default setting is unset (no additional backend filtering). + +CUDECOMP_AUTOTUNE_P_ROWS_RANGE +------------------------------ +(since v0.7.0) + +:code:`CUDECOMP_AUTOTUNE_P_ROWS_RANGE` restricts process-grid autotuning to row dimensions in the inclusive range +:code:`,`. + +Default setting is unset (no row-dimension restriction). + +CUDECOMP_AUTOTUNE_P_COLS_RANGE +------------------------------ +(since v0.7.0) + +:code:`CUDECOMP_AUTOTUNE_P_COLS_RANGE` restricts process-grid autotuning to column dimensions in the inclusive range +:code:`,`. + +Default setting is unset (no column-dimension restriction). diff --git a/include/cudecomp.h b/include/cudecomp.h index 009ae958..af6a8584 100644 --- a/include/cudecomp.h +++ b/include/cudecomp.h @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 * * Licensed under the Apache License, Version 2.0 (the "License"); @@ -158,6 +158,7 @@ typedef struct { cudecompDataType_t dtype; ///< datatype to use during autotuning (default: CUDECOMP_DOUBLE) bool allow_uneven_decompositions; ///< flag to control whether autotuning allows process grids that result in uneven ///< distributions of elements across processes (default: true) + bool disable_mpi_backends; ///< flag to disable MPI backend options during autotuning (default: false) bool disable_nccl_backends; ///< flag to disable NCCL backend options during autotuning (default: false) bool disable_nvshmem_backends; ///< flag to disable NVSHMEM backend options during autotuning (default: false) double skip_threshold; ///< threshold used to skip testing slow configurations; skip configuration diff --git a/include/internal/autotune.h b/include/internal/autotune.h index 5c8bcc14..ae55c77a 100644 --- a/include/internal/autotune.h +++ b/include/internal/autotune.h @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 * * Licensed under the Apache License, Version 2.0 (the "License"); @@ -18,9 +18,19 @@ #ifndef CUDECOMP_AUTOTUNE_H #define CUDECOMP_AUTOTUNE_H +#include +#include +#include + #include "cudecomp.h" namespace cudecomp { +std::vector +getAutotuneTransposeBackendCandidates(const cudecompGridDescAutotuneOptions_t* options); +std::vector +getAutotuneHaloBackendCandidates(const cudecompGridDescAutotuneOptions_t* options); +std::vector> getAutotunePdimCandidates(int nranks, cudecompRankOrder_t rank_order); + void autotuneTransposeBackend(cudecompHandle_t handle, cudecompGridDesc_t grid_desc, const cudecompGridDescAutotuneOptions_t* options); void autotuneHaloBackend(cudecompHandle_t handle, cudecompGridDesc_t grid_desc, diff --git a/src/autotune.cc b/src/autotune.cc index 1ad66153..8c9ab40e 100644 --- a/src/autotune.cc +++ b/src/autotune.cc @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 * * Licensed under the Apache License, Version 2.0 (the "License"); @@ -17,14 +17,18 @@ #include #include +#include #include #include #include #include +#include #include #include #include +#include #include +#include #include #include @@ -38,11 +42,14 @@ #include "cudecomp.h" #include "internal/checks.h" #include "internal/common.h" +#include "internal/exceptions.h" #include "internal/halo.h" #include "internal/performance.h" #include "internal/transpose.h" namespace cudecomp { +using namespace std::string_view_literals; + namespace { struct cudaBufferGuardDeleter { @@ -98,6 +105,65 @@ static std::vector> getPdimCandidates(int nranks, cudecom return pdim_list; } +template +static std::vector +filterBackendCandidates(const char* env_name, const std::array, N>& supported, + std::vector available) { + const char* env_value = std::getenv(env_name); + if (!env_value) return available; + + std::string_view value(env_value); + bool exclude = !value.empty() && value.front() == '^'; + if (exclude) value.remove_prefix(1); + + std::vector listed; + size_t start = 0; + while (start <= value.size()) { + size_t end = value.find(',', start); + if (end == std::string_view::npos) end = value.size(); + auto name = value.substr(start, end - start); + auto match = + std::find_if(supported.begin(), supported.end(), [name](const auto& entry) { return entry.first == name; }); + if (match == supported.end()) { + std::string message = + std::string(env_name) + " contains unknown or empty backend name '" + std::string(name) + "'"; + THROW_INVALID_USAGE(message.c_str()); + } + listed.push_back(match->second); + if (end == value.size()) break; + start = end + 1; + } + + available.erase(std::remove_if(available.begin(), available.end(), + [&](Backend backend) { + bool is_listed = std::find(listed.begin(), listed.end(), backend) != listed.end(); + return exclude ? is_listed : !is_listed; + }), + available.end()); + return available; +} + +static std::pair parsePdimRange(const char* env_name) { + std::string_view value(std::getenv(env_name)); + size_t comma = value.find(','); + int32_t min = 0; + int32_t max = 0; + bool valid = comma != std::string_view::npos && comma > 0 && comma + 1 < value.size() && + value.find(',', comma + 1) == std::string_view::npos; + if (valid) { + auto min_result = std::from_chars(value.data(), value.data() + comma, min); + auto max_result = std::from_chars(value.data() + comma + 1, value.data() + value.size(), max); + valid = min_result.ec == std::errc() && min_result.ptr == value.data() + comma && max_result.ec == std::errc() && + max_result.ptr == value.data() + value.size() && min >= 0 && max > 0 && min <= max; + } + if (!valid) { + std::string message = + std::string(env_name) + " must be comma-separated nonnegative min and positive max with min <= max"; + THROW_INVALID_USAGE(message.c_str()); + } + return {min, max}; +} + template static std::vector processTimings(cudecompHandle_t handle, std::vector times, T scale = 1) { std::sort(times.begin(), times.end()); double t_min = times[0]; @@ -123,6 +189,89 @@ template static std::vector processTimings(cudecompHandle_t hand } // namespace +std::vector +getAutotuneTransposeBackendCandidates(const cudecompGridDescAutotuneOptions_t* options) { + static constexpr std::array supported{ + std::pair{"MPI_P2P"sv, CUDECOMP_TRANSPOSE_COMM_MPI_P2P}, + std::pair{"MPI_P2P_PL"sv, CUDECOMP_TRANSPOSE_COMM_MPI_P2P_PL}, + std::pair{"MPI_A2A"sv, CUDECOMP_TRANSPOSE_COMM_MPI_A2A}, + std::pair{"NCCL"sv, CUDECOMP_TRANSPOSE_COMM_NCCL}, + std::pair{"NCCL_PL"sv, CUDECOMP_TRANSPOSE_COMM_NCCL_PL}, + std::pair{"NVSHMEM"sv, CUDECOMP_TRANSPOSE_COMM_NVSHMEM}, + std::pair{"NVSHMEM_PL"sv, CUDECOMP_TRANSPOSE_COMM_NVSHMEM_PL}, + std::pair{"NVSHMEM_SM"sv, CUDECOMP_TRANSPOSE_COMM_NVSHMEM_SM}, + }; + std::vector candidates{CUDECOMP_TRANSPOSE_COMM_MPI_P2P, + CUDECOMP_TRANSPOSE_COMM_MPI_P2P_PL, + CUDECOMP_TRANSPOSE_COMM_MPI_A2A, CUDECOMP_TRANSPOSE_COMM_NCCL, + CUDECOMP_TRANSPOSE_COMM_NCCL_PL}; +#ifdef ENABLE_NVSHMEM + candidates.insert(candidates.end(), {CUDECOMP_TRANSPOSE_COMM_NVSHMEM, CUDECOMP_TRANSPOSE_COMM_NVSHMEM_PL, + CUDECOMP_TRANSPOSE_COMM_NVSHMEM_SM}); +#endif + candidates = filterBackendCandidates("CUDECOMP_AUTOTUNE_TRANSPOSE_BACKENDS", supported, std::move(candidates)); + candidates.erase(std::remove_if(candidates.begin(), candidates.end(), + [&](auto backend) { + return (options->disable_mpi_backends && transposeBackendRequiresMpi(backend)) || + (options->disable_nccl_backends && transposeBackendRequiresNccl(backend)) || + (options->disable_nvshmem_backends && + transposeBackendRequiresNvshmem(backend)); + }), + candidates.end()); + if (candidates.empty()) { + THROW_INVALID_USAGE("Transpose backend autotuning has no usable candidates after applying filters"); + } + return candidates; +} + +std::vector +getAutotuneHaloBackendCandidates(const cudecompGridDescAutotuneOptions_t* options) { + static constexpr std::array supported{ + std::pair{"MPI"sv, CUDECOMP_HALO_COMM_MPI}, + std::pair{"MPI_BLOCKING"sv, CUDECOMP_HALO_COMM_MPI_BLOCKING}, + std::pair{"NCCL"sv, CUDECOMP_HALO_COMM_NCCL}, + std::pair{"NVSHMEM"sv, CUDECOMP_HALO_COMM_NVSHMEM}, + std::pair{"NVSHMEM_BLOCKING"sv, CUDECOMP_HALO_COMM_NVSHMEM_BLOCKING}, + }; + std::vector candidates{CUDECOMP_HALO_COMM_MPI, CUDECOMP_HALO_COMM_MPI_BLOCKING, + CUDECOMP_HALO_COMM_NCCL}; +#ifdef ENABLE_NVSHMEM + candidates.insert(candidates.end(), {CUDECOMP_HALO_COMM_NVSHMEM, CUDECOMP_HALO_COMM_NVSHMEM_BLOCKING}); +#endif + candidates = filterBackendCandidates("CUDECOMP_AUTOTUNE_HALO_BACKENDS", supported, std::move(candidates)); + candidates.erase(std::remove_if(candidates.begin(), candidates.end(), + [&](auto backend) { + return (options->disable_mpi_backends && haloBackendRequiresMpi(backend)) || + (options->disable_nccl_backends && haloBackendRequiresNccl(backend)) || + (options->disable_nvshmem_backends && haloBackendRequiresNvshmem(backend)); + }), + candidates.end()); + if (candidates.empty()) { + THROW_INVALID_USAGE("Halo backend autotuning has no usable candidates after applying filters"); + } + return candidates; +} + +std::vector> getAutotunePdimCandidates(int nranks, cudecompRankOrder_t rank_order) { + auto candidates = getPdimCandidates(nranks, rank_order); + const char* rows_value = std::getenv("CUDECOMP_AUTOTUNE_P_ROWS_RANGE"); + const char* cols_value = std::getenv("CUDECOMP_AUTOTUNE_P_COLS_RANGE"); + std::pair rows{1, INT32_MAX}; + std::pair cols{1, INT32_MAX}; + if (rows_value) rows = parsePdimRange("CUDECOMP_AUTOTUNE_P_ROWS_RANGE"); + if (cols_value) cols = parsePdimRange("CUDECOMP_AUTOTUNE_P_COLS_RANGE"); + candidates.erase(std::remove_if(candidates.begin(), candidates.end(), + [&](const auto& pdims) { + return pdims[0] < rows.first || pdims[0] > rows.second || pdims[1] < cols.first || + pdims[1] > cols.second; + }), + candidates.end()); + if (candidates.empty()) { + THROW_INVALID_USAGE("Process-grid autotuning has no usable candidates after applying filters"); + } + return candidates; +} + void autotuneTransposeBackend(cudecompHandle_t handle, cudecompGridDesc_t grid_desc, const cudecompGridDescAutotuneOptions_t* options) { if (handle->rank == 0) printf("CUDECOMP: Running transpose autotuning...\n"); @@ -139,21 +288,9 @@ void autotuneTransposeBackend(cudecompHandle_t handle, cudecompGridDesc_t grid_d bool need_nccl = false; bool need_nvshmem = false; if (autotune_comm) { - comm_backend_list = {CUDECOMP_TRANSPOSE_COMM_MPI_P2P, CUDECOMP_TRANSPOSE_COMM_MPI_P2P_PL, - CUDECOMP_TRANSPOSE_COMM_MPI_A2A}; - if (!options->disable_nccl_backends) { - comm_backend_list.push_back(CUDECOMP_TRANSPOSE_COMM_NCCL); - comm_backend_list.push_back(CUDECOMP_TRANSPOSE_COMM_NCCL_PL); - need_nccl = true; - } -#ifdef ENABLE_NVSHMEM - if (!options->disable_nvshmem_backends) { - comm_backend_list.push_back(CUDECOMP_TRANSPOSE_COMM_NVSHMEM); - comm_backend_list.push_back(CUDECOMP_TRANSPOSE_COMM_NVSHMEM_PL); - comm_backend_list.push_back(CUDECOMP_TRANSPOSE_COMM_NVSHMEM_SM); - need_nvshmem = true; - } -#endif + comm_backend_list = getAutotuneTransposeBackendCandidates(options); + need_nccl = std::any_of(comm_backend_list.begin(), comm_backend_list.end(), transposeBackendRequiresNccl); + need_nvshmem = std::any_of(comm_backend_list.begin(), comm_backend_list.end(), transposeBackendRequiresNvshmem); } else { comm_backend_list = {grid_desc->config.transpose_comm_backend}; if (transposeBackendRequiresNccl(comm_backend_list[0])) { need_nccl = true; } @@ -161,6 +298,8 @@ void autotuneTransposeBackend(cudecompHandle_t handle, cudecompGridDesc_t grid_d if (transposeBackendRequiresNvshmem(comm_backend_list[0])) { need_nvshmem = true; } #endif } + const bool need_non_nvshmem = std::any_of(comm_backend_list.begin(), comm_backend_list.end(), + [](auto backend) { return !transposeBackendRequiresNvshmem(backend); }); bool need_data2 = false; for (int i = 0; i < 4; ++i) { @@ -169,7 +308,7 @@ void autotuneTransposeBackend(cudecompHandle_t handle, cudecompGridDesc_t grid_d std::vector> pdim_list; if (autotune_pdims) { - pdim_list = getPdimCandidates(handle->nranks, grid_desc->config.rank_order); + pdim_list = getAutotunePdimCandidates(handle->nranks, grid_desc->config.rank_order); } else { pdim_list.push_back({grid_desc->config.pdims[0], grid_desc->config.pdims[1]}); } @@ -291,40 +430,42 @@ void autotuneTransposeBackend(cudecompHandle_t handle, cudecompGridDesc_t grid_d work_nvshmem_guard = transposeWorkspaceGuard(work_nvshmem, {handle, grid_desc, CUDECOMP_TRANSPOSE_COMM_NVSHMEM}); - // Check if there is enough memory for separate non-NVSHMEM allocated work buffer - const bool allow_nvshmem_workspace_fallback = !handle->cuda_cumem_enable && !handle->nccl_enable_ubr; - auto ret = cudaMalloc(&work, work_sz); - int any_oom = (ret == cudaErrorMemoryAllocation); - CHECK_MPI(MPI_Allreduce(MPI_IN_PLACE, &any_oom, 1, MPI_INT, MPI_LOR, handle->mpi_comm)); - if (any_oom) { - if (ret == cudaSuccess) { - CHECK_CUDA(cudaFree(work)); - work = nullptr; - } else if (ret == cudaErrorMemoryAllocation) { - cudaGetLastError(); // Reset CUDA error state + if (need_non_nvshmem) { + // Check if there is enough memory for separate non-NVSHMEM allocated work buffer + const bool allow_nvshmem_workspace_fallback = !handle->cuda_cumem_enable && !handle->nccl_enable_ubr; + auto ret = cudaMalloc(&work, work_sz); + int any_oom = (ret == cudaErrorMemoryAllocation); + CHECK_MPI(MPI_Allreduce(MPI_IN_PLACE, &any_oom, 1, MPI_INT, MPI_LOR, handle->mpi_comm)); + if (any_oom) { + if (ret == cudaSuccess) { + CHECK_CUDA(cudaFree(work)); + work = nullptr; + } else if (ret == cudaErrorMemoryAllocation) { + cudaGetLastError(); // Reset CUDA error state + } else { + CHECK_CUDA(ret); + } + if (!allow_nvshmem_workspace_fallback) { + THROW_CUDA_ERROR( + "Cannot allocate separate non-NVSHMEM workspace during autotuning while cuMem or NCCL user buffer " + "registration is enabled."); + } + if (handle->rank == 0) { + printf("CUDECOMP:WARN: Cannot allocate separate workspace for non-NVSHMEM backends during " + "autotuning. Using NVSHMEM allocated workspace for all backends, which may cause issues " + "for some MPI implementations. See documentation for more details and suggested workarounds.\n"); + } + work = work_nvshmem; } else { CHECK_CUDA(ret); + CHECK_CUDA(cudaFree(work)); + auto backend = (need_nccl) ? CUDECOMP_TRANSPOSE_COMM_NCCL : CUDECOMP_TRANSPOSE_COMM_MPI_P2P; + tmp = grid_desc->config.transpose_comm_backend; + grid_desc->config.transpose_comm_backend = backend; + CHECK_CUDECOMP(cudecompMalloc(handle, grid_desc, reinterpret_cast(&work), work_sz)); + grid_desc->config.transpose_comm_backend = tmp; + work_guard = transposeWorkspaceGuard(work, {handle, grid_desc, backend}); } - if (!allow_nvshmem_workspace_fallback) { - THROW_CUDA_ERROR( - "Cannot allocate separate non-NVSHMEM workspace during autotuning while cuMem or NCCL user buffer " - "registration is enabled."); - } - if (handle->rank == 0) { - printf("CUDECOMP:WARN: Cannot allocate separate workspace for non-NVSHMEM backends during " - "autotuning. Using NVSHMEM allocated workspace for all backends, which may cause issues " - "for some MPI implementations. See documentation for more details and suggested workarounds.\n"); - } - work = work_nvshmem; - } else { - CHECK_CUDA(ret); - CHECK_CUDA(cudaFree(work)); - auto backend = (need_nccl) ? CUDECOMP_TRANSPOSE_COMM_NCCL : CUDECOMP_TRANSPOSE_COMM_MPI_P2P; - tmp = grid_desc->config.transpose_comm_backend; - grid_desc->config.transpose_comm_backend = backend; - CHECK_CUDECOMP(cudecompMalloc(handle, grid_desc, reinterpret_cast(&work), work_sz)); - grid_desc->config.transpose_comm_backend = tmp; - work_guard = transposeWorkspaceGuard(work, {handle, grid_desc, backend}); } #endif } else { @@ -554,7 +695,7 @@ void autotuneTransposeBackend(cudecompHandle_t handle, cudecompGridDesc_t grid_d // Free test data and workspace if (need_nvshmem) { - if (work != work_nvshmem) { + if (need_non_nvshmem && work != work_nvshmem) { auto tmp = grid_desc->config.transpose_comm_backend; grid_desc->config.transpose_comm_backend = (need_nccl) ? CUDECOMP_TRANSPOSE_COMM_NCCL : CUDECOMP_TRANSPOSE_COMM_MPI_P2P; @@ -646,18 +787,9 @@ void autotuneHaloBackend(cudecompHandle_t handle, cudecompGridDesc_t grid_desc, bool need_nccl = false; bool need_nvshmem = false; if (autotune_comm) { - comm_backend_list = {CUDECOMP_HALO_COMM_MPI, CUDECOMP_HALO_COMM_MPI_BLOCKING}; - if (!options->disable_nccl_backends) { - comm_backend_list.push_back(CUDECOMP_HALO_COMM_NCCL); - need_nccl = true; - } -#ifdef ENABLE_NVSHMEM - if (!options->disable_nvshmem_backends) { - comm_backend_list.push_back(CUDECOMP_HALO_COMM_NVSHMEM); - comm_backend_list.push_back(CUDECOMP_HALO_COMM_NVSHMEM_BLOCKING); - need_nvshmem = true; - } -#endif + comm_backend_list = getAutotuneHaloBackendCandidates(options); + need_nccl = std::any_of(comm_backend_list.begin(), comm_backend_list.end(), haloBackendRequiresNccl); + need_nvshmem = std::any_of(comm_backend_list.begin(), comm_backend_list.end(), haloBackendRequiresNvshmem); } else { comm_backend_list = {grid_desc->config.halo_comm_backend}; if (haloBackendRequiresNccl(comm_backend_list[0])) { need_nccl = true; } @@ -665,10 +797,12 @@ void autotuneHaloBackend(cudecompHandle_t handle, cudecompGridDesc_t grid_desc, if (haloBackendRequiresNvshmem(comm_backend_list[0])) { need_nvshmem = true; } #endif } + const bool need_non_nvshmem = std::any_of(comm_backend_list.begin(), comm_backend_list.end(), + [](auto backend) { return !haloBackendRequiresNvshmem(backend); }); std::vector> pdim_list; if (autotune_pdims) { - pdim_list = getPdimCandidates(handle->nranks, grid_desc->config.rank_order); + pdim_list = getAutotunePdimCandidates(handle->nranks, grid_desc->config.rank_order); } else { pdim_list.push_back({grid_desc->config.pdims[0], grid_desc->config.pdims[1]}); } @@ -761,40 +895,42 @@ void autotuneHaloBackend(cudecompHandle_t handle, cudecompGridDesc_t grid_desc, grid_desc->config.halo_comm_backend = tmp; work_nvshmem_guard = haloWorkspaceGuard(work_nvshmem, {handle, grid_desc, CUDECOMP_HALO_COMM_NVSHMEM}); - // Check if there is enough memory for separate non-NVSHMEM allocated work buffer - const bool allow_nvshmem_workspace_fallback = !handle->cuda_cumem_enable && !handle->nccl_enable_ubr; - auto ret = cudaMalloc(&work, work_sz); - int any_oom = (ret == cudaErrorMemoryAllocation); - CHECK_MPI(MPI_Allreduce(MPI_IN_PLACE, &any_oom, 1, MPI_INT, MPI_LOR, handle->mpi_comm)); - if (any_oom) { - if (ret == cudaSuccess) { - CHECK_CUDA(cudaFree(work)); - work = nullptr; - } else if (ret == cudaErrorMemoryAllocation) { - cudaGetLastError(); // Reset CUDA error state + if (need_non_nvshmem) { + // Check if there is enough memory for separate non-NVSHMEM allocated work buffer + const bool allow_nvshmem_workspace_fallback = !handle->cuda_cumem_enable && !handle->nccl_enable_ubr; + auto ret = cudaMalloc(&work, work_sz); + int any_oom = (ret == cudaErrorMemoryAllocation); + CHECK_MPI(MPI_Allreduce(MPI_IN_PLACE, &any_oom, 1, MPI_INT, MPI_LOR, handle->mpi_comm)); + if (any_oom) { + if (ret == cudaSuccess) { + CHECK_CUDA(cudaFree(work)); + work = nullptr; + } else if (ret == cudaErrorMemoryAllocation) { + cudaGetLastError(); // Reset CUDA error state + } else { + CHECK_CUDA(ret); + } + if (!allow_nvshmem_workspace_fallback) { + THROW_CUDA_ERROR( + "Cannot allocate separate non-NVSHMEM workspace during autotuning while cuMem or NCCL user buffer " + "registration is enabled."); + } + if (handle->rank == 0) { + printf("CUDECOMP:WARN: Cannot allocate separate workspace for non-NVSHMEM backends during " + "autotuning. Using NVSHMEM allocated workspace for all backends, which may cause issues " + "for some MPI implementations. See documentation for more details and suggested workarounds.\n"); + } + work = work_nvshmem; } else { CHECK_CUDA(ret); + CHECK_CUDA(cudaFree(work)); + auto backend = (need_nccl) ? CUDECOMP_HALO_COMM_NCCL : CUDECOMP_HALO_COMM_MPI; + tmp = grid_desc->config.halo_comm_backend; + grid_desc->config.halo_comm_backend = backend; + CHECK_CUDECOMP(cudecompMalloc(handle, grid_desc, reinterpret_cast(&work), work_sz)); + grid_desc->config.halo_comm_backend = tmp; + work_guard = haloWorkspaceGuard(work, {handle, grid_desc, backend}); } - if (!allow_nvshmem_workspace_fallback) { - THROW_CUDA_ERROR( - "Cannot allocate separate non-NVSHMEM workspace during autotuning while cuMem or NCCL user buffer " - "registration is enabled."); - } - if (handle->rank == 0) { - printf("CUDECOMP:WARN: Cannot allocate separate workspace for non-NVSHMEM backends during " - "autotuning. Using NVSHMEM allocated workspace for all backends, which may cause issues " - "for some MPI implementations. See documentation for more details and suggested workarounds.\n"); - } - work = work_nvshmem; - } else { - CHECK_CUDA(ret); - CHECK_CUDA(cudaFree(work)); - auto backend = (need_nccl) ? CUDECOMP_HALO_COMM_NCCL : CUDECOMP_HALO_COMM_MPI; - tmp = grid_desc->config.halo_comm_backend; - grid_desc->config.halo_comm_backend = backend; - CHECK_CUDECOMP(cudecompMalloc(handle, grid_desc, reinterpret_cast(&work), work_sz)); - grid_desc->config.halo_comm_backend = tmp; - work_guard = haloWorkspaceGuard(work, {handle, grid_desc, backend}); } #endif } else { @@ -938,7 +1074,7 @@ void autotuneHaloBackend(cudecompHandle_t handle, cudecompGridDesc_t grid_desc, // Free test data and workspace if (need_nvshmem) { - if (work != work_nvshmem) { + if (need_non_nvshmem && work != work_nvshmem) { auto tmp = grid_desc->config.halo_comm_backend; grid_desc->config.halo_comm_backend = (need_nccl) ? CUDECOMP_HALO_COMM_NCCL : CUDECOMP_HALO_COMM_MPI; work_guard.release(); diff --git a/src/cudecomp.cc b/src/cudecomp.cc index 1f3be866..84de63d5 100644 --- a/src/cudecomp.cc +++ b/src/cudecomp.cc @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2022-2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-FileCopyrightText: Copyright (c) 2022-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 * * Licensed under the Apache License, Version 2.0 (the "License"); @@ -846,13 +846,9 @@ cudecompResult_t cudecompGridDescCreate(cudecompHandle_t handle, cudecompGridDes // Check some autotuning options bool autotune_transpose_backend = false; bool autotune_halo_backend = false; - bool autotune_disable_nccl_backends = false; - bool autotune_disable_nvshmem_backends = false; if (options) { autotune_transpose_backend = options->autotune_transpose_backend; autotune_halo_backend = options->autotune_halo_backend; - autotune_disable_nccl_backends = options->disable_nccl_backends; - autotune_disable_nvshmem_backends = options->disable_nvshmem_backends; } checkConfig(handle, config, autotune_transpose_backend, autotune_halo_backend); @@ -868,6 +864,25 @@ cudecompResult_t cudecompGridDescCreate(cudecompHandle_t handle, cudecompGridDes auto comm_backend = grid_desc->config.transpose_comm_backend; auto halo_comm_backend = grid_desc->config.halo_comm_backend; + std::vector autotune_transpose_candidates; + std::vector autotune_halo_candidates; + if (autotune_transpose_backend) { autotune_transpose_candidates = getAutotuneTransposeBackendCandidates(options); } + if (autotune_halo_backend) { autotune_halo_candidates = getAutotuneHaloBackendCandidates(options); } + if (autotune_pdims) { (void)getAutotunePdimCandidates(handle->nranks, grid_desc->config.rank_order); } + + bool need_nccl = + (!autotune_transpose_backend && transposeBackendRequiresNccl(comm_backend)) || + (!autotune_halo_backend && haloBackendRequiresNccl(halo_comm_backend)) || + std::any_of(autotune_transpose_candidates.begin(), autotune_transpose_candidates.end(), + transposeBackendRequiresNccl) || + std::any_of(autotune_halo_candidates.begin(), autotune_halo_candidates.end(), haloBackendRequiresNccl); + bool need_nvshmem = + (!autotune_transpose_backend && transposeBackendRequiresNvshmem(comm_backend)) || + (!autotune_halo_backend && haloBackendRequiresNvshmem(halo_comm_backend)) || + std::any_of(autotune_transpose_candidates.begin(), autotune_transpose_candidates.end(), + transposeBackendRequiresNvshmem) || + std::any_of(autotune_halo_candidates.begin(), autotune_halo_candidates.end(), haloBackendRequiresNvshmem); + // If transpose_mem_order not used, set based on transpose_axis_contiguous settings) grid_desc->transpose_mem_order_set = (config->transpose_mem_order[0][0] >= 0); @@ -901,8 +916,7 @@ cudecompResult_t cudecompGridDescCreate(cudecompHandle_t handle, cudecompGridDes } // Initialize NCCL communicator if needed - if (transposeBackendRequiresNccl(comm_backend) || haloBackendRequiresNccl(halo_comm_backend) || - ((autotune_transpose_backend || autotune_halo_backend) && !autotune_disable_nccl_backends)) { + if (need_nccl) { if (!handle->nccl_comm) { handle->nccl_comm = ncclCommFromMPIComm(handle->mpi_comm); } if (!handle->nccl_local_comm) { if (grid_desc->config.pdims[0] > 0 && grid_desc->config.pdims[1] > 0) { @@ -939,8 +953,7 @@ cudecompResult_t cudecompGridDescCreate(cudecompHandle_t handle, cudecompGridDes } // Initialize NVSHMEM if needed - if (transposeBackendRequiresNvshmem(comm_backend) || haloBackendRequiresNvshmem(halo_comm_backend) || - ((autotune_transpose_backend || autotune_halo_backend) && !autotune_disable_nvshmem_backends)) { + if (need_nvshmem) { #ifdef ENABLE_NVSHMEM nvshmem_runtime = acquireNvshmemRuntime(handle); grid_desc->nvshmem_runtime = nvshmem_runtime; @@ -969,11 +982,11 @@ cudecompResult_t cudecompGridDescCreate(cudecompHandle_t handle, cudecompGridDes } // Setup final row and column communicators - bool need_nvshmem = transposeBackendRequiresNvshmem(grid_desc->config.transpose_comm_backend) || - haloBackendRequiresNvshmem(grid_desc->config.halo_comm_backend); - createCommInfo(handle, grid_desc, need_nvshmem); + bool final_need_nvshmem = transposeBackendRequiresNvshmem(grid_desc->config.transpose_comm_backend) || + haloBackendRequiresNvshmem(grid_desc->config.halo_comm_backend); + createCommInfo(handle, grid_desc, final_need_nvshmem); #ifdef ENABLE_NVSHMEM - if (need_nvshmem) { + if (final_need_nvshmem) { if (!nvshmem_runtime || !nvshmem_runtime->initialized) { nvshmem_runtime = acquireNvshmemRuntime(handle); } CHECK_CUDA(cudaMalloc(&grid_desc->nvshmem_block_counters, handle->nranks * sizeof(int))); CHECK_CUDA(cudaMemset(grid_desc->nvshmem_block_counters, 0, handle->nranks * sizeof(int))); @@ -1086,6 +1099,7 @@ cudecompResult_t cudecompGridDescAutotuneOptionsSetDefaults(cudecompGridDescAuto options->grid_mode = CUDECOMP_AUTOTUNE_GRID_TRANSPOSE; options->dtype = CUDECOMP_DOUBLE; options->allow_uneven_decompositions = true; + options->disable_mpi_backends = false; options->disable_nccl_backends = false; options->disable_nvshmem_backends = false; options->skip_threshold = 0.0; diff --git a/src/cudecomp_m.cuf b/src/cudecomp_m.cuf index 8e63d35f..42560d56 100644 --- a/src/cudecomp_m.cuf +++ b/src/cudecomp_m.cuf @@ -1,4 +1,4 @@ -! SPDX-FileCopyrightText: Copyright (c) 2022-2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +! SPDX-FileCopyrightText: Copyright (c) 2022-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. ! SPDX-License-Identifier: Apache-2.0 ! ! Licensed under the Apache License, Version 2.0 (the "License"); @@ -115,6 +115,7 @@ module cudecomp integer(c_int32_t) :: grid_mode ! which communication (transpose/halo) to use to autotune process grid integer(c_int32_t) :: dtype ! datatype to use during autotuning logical(c_bool) :: allow_uneven_decompositions ! flag to control whether autotuning allows uneven decompositions (based on gdims_dist if provided, gdims otherwise) + logical(c_bool) :: disable_mpi_backends ! flag to disable MPI backend options during autotuning logical(c_bool) :: disable_nccl_backends ! flag to disable NCCL backend options during autotuning logical(c_bool) :: disable_nvshmem_backends ! flag to disable NVSHMEM backend options during autotuning real(c_double) :: skip_threshold ! threshold used to skip testing slow configurations diff --git a/tests/ctest/api_tests.cc b/tests/ctest/api_tests.cc index a56ba67a..24327a2f 100644 --- a/tests/ctest/api_tests.cc +++ b/tests/ctest/api_tests.cc @@ -3,10 +3,15 @@ * SPDX-License-Identifier: Apache-2.0 */ +#include #include #include +#include +#include #include #include +#include +#include #include @@ -17,6 +22,8 @@ #endif #include "cudecomp.h" +#include "internal/autotune.h" +#include "internal/exceptions.h" #include "gpu_test_utils.h" #include "mpi_test_utils.h" @@ -24,6 +31,44 @@ namespace { +class ScopedEnvVar { +public: + ScopedEnvVar(const char* name, const char* value) : name_(name) { + const char* previous = std::getenv(name); + if (previous) { + had_previous_ = true; + previous_ = previous; + } + setenv(name, value, 1); + } + + ~ScopedEnvVar() { + if (had_previous_) { + setenv(name_.c_str(), previous_.c_str(), 1); + } else { + unsetenv(name_.c_str()); + } + } + +private: + std::string name_; + bool had_previous_ = false; + std::string previous_; +}; + +template std::vector withoutCandidates(std::vector candidates, std::initializer_list excluded) { + candidates.erase(std::remove_if(candidates.begin(), candidates.end(), + [&](T candidate) { + return std::find(excluded.begin(), excluded.end(), candidate) != excluded.end(); + }), + candidates.end()); + return candidates; +} + +template void expectSameCandidates(const std::vector& actual, const std::vector& expected) { + EXPECT_TRUE(std::is_permutation(actual.begin(), actual.end(), expected.begin(), expected.end())); +} + constexpr int kApiTestRanks = 4; constexpr std::array kGdims{9, 10, 11}; constexpr std::array kGdimsDist{8, 9, 10}; @@ -238,6 +283,7 @@ TEST(ApiGridDescAutotuneOptionsSetDefaultsTest, SetsDocumentedDefaults) { EXPECT_EQ(CUDECOMP_AUTOTUNE_GRID_TRANSPOSE, options.grid_mode); EXPECT_EQ(CUDECOMP_DOUBLE, options.dtype); EXPECT_TRUE(options.allow_uneven_decompositions); + EXPECT_FALSE(options.disable_mpi_backends); EXPECT_FALSE(options.disable_nccl_backends); EXPECT_FALSE(options.disable_nvshmem_backends); EXPECT_EQ(0.0, options.skip_threshold); @@ -263,6 +309,132 @@ TEST(ApiGridDescAutotuneOptionsSetDefaultsTest, SetsDocumentedDefaults) { } } +TEST(AutotuneCandidateFilterTest, AppliesTransposeInclusionList) { + cudecompGridDescAutotuneOptions_t options; + ASSERT_EQ(CUDECOMP_RESULT_SUCCESS, cudecompGridDescAutotuneOptionsSetDefaults(&options)); + ScopedEnvVar filter("CUDECOMP_AUTOTUNE_TRANSPOSE_BACKENDS", "NCCL,MPI_A2A,MPI_P2P"); + + auto candidates = cudecomp::getAutotuneTransposeBackendCandidates(&options); + std::vector expected{CUDECOMP_TRANSPOSE_COMM_MPI_P2P, CUDECOMP_TRANSPOSE_COMM_MPI_A2A, + CUDECOMP_TRANSPOSE_COMM_NCCL}; + expectSameCandidates(candidates, expected); +} + +TEST(AutotuneCandidateFilterTest, AppliesTransposeExclusionList) { + cudecompGridDescAutotuneOptions_t options; + ASSERT_EQ(CUDECOMP_RESULT_SUCCESS, cudecompGridDescAutotuneOptionsSetDefaults(&options)); + auto expected = + withoutCandidates(cudecomp::getAutotuneTransposeBackendCandidates(&options), {CUDECOMP_TRANSPOSE_COMM_MPI_P2P}); + ScopedEnvVar filter("CUDECOMP_AUTOTUNE_TRANSPOSE_BACKENDS", "^MPI_P2P"); + + auto candidates = cudecomp::getAutotuneTransposeBackendCandidates(&options); + expectSameCandidates(candidates, expected); +} + +TEST(AutotuneCandidateFilterTest, AppliesHaloInclusionList) { + cudecompGridDescAutotuneOptions_t options; + ASSERT_EQ(CUDECOMP_RESULT_SUCCESS, cudecompGridDescAutotuneOptionsSetDefaults(&options)); + ScopedEnvVar halo_filter("CUDECOMP_AUTOTUNE_HALO_BACKENDS", "NCCL,MPI_BLOCKING"); + + auto candidates = cudecomp::getAutotuneHaloBackendCandidates(&options); + std::vector expected{CUDECOMP_HALO_COMM_MPI_BLOCKING, CUDECOMP_HALO_COMM_NCCL}; + expectSameCandidates(candidates, expected); +} + +TEST(AutotuneCandidateFilterTest, AppliesTransposeBackendFamilyDisables) { + cudecompGridDescAutotuneOptions_t options; + ASSERT_EQ(CUDECOMP_RESULT_SUCCESS, cudecompGridDescAutotuneOptionsSetDefaults(&options)); + auto all_candidates = cudecomp::getAutotuneTransposeBackendCandidates(&options); + + options.disable_mpi_backends = true; + auto candidates = cudecomp::getAutotuneTransposeBackendCandidates(&options); + auto expected = + withoutCandidates(all_candidates, {CUDECOMP_TRANSPOSE_COMM_MPI_P2P, CUDECOMP_TRANSPOSE_COMM_MPI_P2P_PL, + CUDECOMP_TRANSPOSE_COMM_MPI_A2A}); + expectSameCandidates(candidates, expected); + + options.disable_mpi_backends = false; + options.disable_nccl_backends = true; + candidates = cudecomp::getAutotuneTransposeBackendCandidates(&options); + expected = withoutCandidates(all_candidates, {CUDECOMP_TRANSPOSE_COMM_NCCL, CUDECOMP_TRANSPOSE_COMM_NCCL_PL}); + expectSameCandidates(candidates, expected); + + options.disable_nccl_backends = false; + options.disable_nvshmem_backends = true; + candidates = cudecomp::getAutotuneTransposeBackendCandidates(&options); + expected = withoutCandidates(all_candidates, {CUDECOMP_TRANSPOSE_COMM_NVSHMEM, CUDECOMP_TRANSPOSE_COMM_NVSHMEM_PL, + CUDECOMP_TRANSPOSE_COMM_NVSHMEM_SM}); + expectSameCandidates(candidates, expected); +} + +TEST(AutotuneCandidateFilterTest, AppliesHaloBackendFamilyDisables) { + cudecompGridDescAutotuneOptions_t options; + ASSERT_EQ(CUDECOMP_RESULT_SUCCESS, cudecompGridDescAutotuneOptionsSetDefaults(&options)); + auto all_candidates = cudecomp::getAutotuneHaloBackendCandidates(&options); + + options.disable_mpi_backends = true; + auto candidates = cudecomp::getAutotuneHaloBackendCandidates(&options); + auto expected = withoutCandidates(all_candidates, {CUDECOMP_HALO_COMM_MPI, CUDECOMP_HALO_COMM_MPI_BLOCKING}); + expectSameCandidates(candidates, expected); + + options.disable_mpi_backends = false; + options.disable_nccl_backends = true; + candidates = cudecomp::getAutotuneHaloBackendCandidates(&options); + expected = withoutCandidates(all_candidates, {CUDECOMP_HALO_COMM_NCCL}); + expectSameCandidates(candidates, expected); + + options.disable_nccl_backends = false; + options.disable_nvshmem_backends = true; + candidates = cudecomp::getAutotuneHaloBackendCandidates(&options); + expected = withoutCandidates(all_candidates, {CUDECOMP_HALO_COMM_NVSHMEM, CUDECOMP_HALO_COMM_NVSHMEM_BLOCKING}); + expectSameCandidates(candidates, expected); +} + +TEST(AutotuneCandidateFilterTest, RejectsMalformedAndEmptyBackendSets) { + cudecompGridDescAutotuneOptions_t options; + ASSERT_EQ(CUDECOMP_RESULT_SUCCESS, cudecompGridDescAutotuneOptionsSetDefaults(&options)); + { + ScopedEnvVar filter("CUDECOMP_AUTOTUNE_HALO_BACKENDS", "MPI,,NCCL"); + EXPECT_THROW(cudecomp::getAutotuneHaloBackendCandidates(&options), cudecomp::InvalidUsage); + } + { + ScopedEnvVar filter("CUDECOMP_AUTOTUNE_HALO_BACKENDS", "^MPI,MPI_BLOCKING,NCCL,NVSHMEM,NVSHMEM_BLOCKING"); + EXPECT_THROW(cudecomp::getAutotuneHaloBackendCandidates(&options), cudecomp::InvalidUsage); + } +} + +TEST(AutotuneCandidateFilterTest, FiltersProcessGridRanges) { + { + ScopedEnvVar rows("CUDECOMP_AUTOTUNE_P_ROWS_RANGE", "2,4"); + auto candidates = cudecomp::getAutotunePdimCandidates(4, CUDECOMP_RANK_ORDER_ROW_MAJOR); + std::vector> expected{{4, 1}, {2, 2}}; + expectSameCandidates(candidates, expected); + } + { + ScopedEnvVar cols("CUDECOMP_AUTOTUNE_P_COLS_RANGE", "2,4"); + auto candidates = cudecomp::getAutotunePdimCandidates(4, CUDECOMP_RANK_ORDER_ROW_MAJOR); + std::vector> expected{{2, 2}, {1, 4}}; + expectSameCandidates(candidates, expected); + } + { + ScopedEnvVar rows("CUDECOMP_AUTOTUNE_P_ROWS_RANGE", "2,2"); + ScopedEnvVar cols("CUDECOMP_AUTOTUNE_P_COLS_RANGE", "2,2"); + EXPECT_EQ((std::vector>{{2, 2}}), + cudecomp::getAutotunePdimCandidates(4, CUDECOMP_RANK_ORDER_ROW_MAJOR)); + } +} + +TEST(AutotuneCandidateFilterTest, RejectsMalformedAndEmptyProcessGridRanges) { + { + ScopedEnvVar rows("CUDECOMP_AUTOTUNE_P_ROWS_RANGE", "2,1"); + EXPECT_THROW(cudecomp::getAutotunePdimCandidates(4, CUDECOMP_RANK_ORDER_ROW_MAJOR), cudecomp::InvalidUsage); + } + { + ScopedEnvVar cols("CUDECOMP_AUTOTUNE_P_COLS_RANGE", "3,3"); + EXPECT_THROW(cudecomp::getAutotunePdimCandidates(4, CUDECOMP_RANK_ORDER_ROW_MAJOR), cudecomp::InvalidUsage); + } +} + TEST(ApiGridDescAutotuneOptionsSetDefaultsTest, RejectsInvalidArguments) { EXPECT_EQ(CUDECOMP_RESULT_INVALID_USAGE, cudecompGridDescAutotuneOptionsSetDefaults(nullptr)); } @@ -918,6 +1090,19 @@ TEST_F(ApiGridDescCreateTest, RejectsInvalidAutotuneInputs) { EXPECT_EQ(CUDECOMP_RESULT_INVALID_USAGE, cudecompGridDescCreate(handle_, &unused_grid_desc, &config, &options)); } +TEST_F(ApiGridDescCreateTest, IgnoresAutotuneEnvironmentFiltersForFixedSelections) { + auto config = distributedConfig(); + auto options = fastAutotuneOptions(); + ScopedEnvVar transpose_backends("CUDECOMP_AUTOTUNE_TRANSPOSE_BACKENDS", "invalid"); + ScopedEnvVar halo_backends("CUDECOMP_AUTOTUNE_HALO_BACKENDS", "invalid"); + ScopedEnvVar rows("CUDECOMP_AUTOTUNE_P_ROWS_RANGE", "invalid"); + ScopedEnvVar cols("CUDECOMP_AUTOTUNE_P_COLS_RANGE", "invalid"); + + cudecompGridDesc_t grid_desc = nullptr; + CHECK_CUDECOMP_GLOBAL(active_comm_, cudecompGridDescCreate(handle_, &grid_desc, &config, &options)); + cudecomp_test::gridDescGuard grid_desc_guard(handle_, grid_desc); +} + TEST_F(ApiGridDescDestroyTest, RejectsInvalidArguments) { EXPECT_EQ(CUDECOMP_RESULT_INVALID_USAGE, cudecompGridDescDestroy(handle_, nullptr)); } diff --git a/tests/ctest/fortran_api_test.f90 b/tests/ctest/fortran_api_test.f90 index 37ed0a84..6111d5b1 100644 --- a/tests/ctest/fortran_api_test.f90 +++ b/tests/ctest/fortran_api_test.f90 @@ -109,6 +109,9 @@ subroutine test_default_values() call expect_equal_int(options%dtype, CUDECOMP_DOUBLE, "default autotune dtype") call expect_equal_int(options%halo_axis, 1, "default Fortran halo axis") if (.not. options%allow_uneven_decompositions) call record_failure("default uneven decompositions flag is false") + if (options%disable_mpi_backends) call record_failure("default disable MPI backends flag is true") + if (options%disable_nccl_backends) call record_failure("default disable NCCL backends flag is true") + if (options%disable_nvshmem_backends) call record_failure("default disable NVSHMEM backends flag is true") if (options%autotune_transpose_backend) call record_failure("default transpose backend autotune flag is true") if (options%autotune_halo_backend) call record_failure("default halo backend autotune flag is true") end subroutine test_default_values @@ -141,6 +144,7 @@ subroutine test_grid_descriptor_contracts(handle) options%n_warmup_trials = 0 options%n_trials = 1 options%dtype = CUDECOMP_FLOAT + options%disable_mpi_backends = .true. options%disable_nccl_backends = .true. options%disable_nvshmem_backends = .true. options%halo_axis = 3 From dbe1b0ff6f7716fa9e20a69e8e4829e077c99427 Mon Sep 17 00:00:00 2001 From: romerojosh Date: Thu, 16 Jul 2026 18:19:53 +0000 Subject: [PATCH 2/2] Use singular process-grid environment names Signed-off-by: romerojosh --- docs/env_vars.rst | 8 ++++---- src/autotune.cc | 8 ++++---- tests/ctest/api_tests.cc | 16 ++++++++-------- 3 files changed, 16 insertions(+), 16 deletions(-) diff --git a/docs/env_vars.rst b/docs/env_vars.rst index 023d9bb6..15426e83 100644 --- a/docs/env_vars.rst +++ b/docs/env_vars.rst @@ -127,20 +127,20 @@ and :code:`NVSHMEM_BLOCKING`. Default setting is unset (no additional backend filtering). -CUDECOMP_AUTOTUNE_P_ROWS_RANGE +CUDECOMP_AUTOTUNE_P_ROW_RANGE ------------------------------ (since v0.7.0) -:code:`CUDECOMP_AUTOTUNE_P_ROWS_RANGE` restricts process-grid autotuning to row dimensions in the inclusive range +:code:`CUDECOMP_AUTOTUNE_P_ROW_RANGE` restricts process-grid autotuning to row dimensions in the inclusive range :code:`,`. Default setting is unset (no row-dimension restriction). -CUDECOMP_AUTOTUNE_P_COLS_RANGE +CUDECOMP_AUTOTUNE_P_COL_RANGE ------------------------------ (since v0.7.0) -:code:`CUDECOMP_AUTOTUNE_P_COLS_RANGE` restricts process-grid autotuning to column dimensions in the inclusive range +:code:`CUDECOMP_AUTOTUNE_P_COL_RANGE` restricts process-grid autotuning to column dimensions in the inclusive range :code:`,`. Default setting is unset (no column-dimension restriction). diff --git a/src/autotune.cc b/src/autotune.cc index 8c9ab40e..80b92bff 100644 --- a/src/autotune.cc +++ b/src/autotune.cc @@ -254,12 +254,12 @@ getAutotuneHaloBackendCandidates(const cudecompGridDescAutotuneOptions_t* option std::vector> getAutotunePdimCandidates(int nranks, cudecompRankOrder_t rank_order) { auto candidates = getPdimCandidates(nranks, rank_order); - const char* rows_value = std::getenv("CUDECOMP_AUTOTUNE_P_ROWS_RANGE"); - const char* cols_value = std::getenv("CUDECOMP_AUTOTUNE_P_COLS_RANGE"); + const char* rows_value = std::getenv("CUDECOMP_AUTOTUNE_P_ROW_RANGE"); + const char* cols_value = std::getenv("CUDECOMP_AUTOTUNE_P_COL_RANGE"); std::pair rows{1, INT32_MAX}; std::pair cols{1, INT32_MAX}; - if (rows_value) rows = parsePdimRange("CUDECOMP_AUTOTUNE_P_ROWS_RANGE"); - if (cols_value) cols = parsePdimRange("CUDECOMP_AUTOTUNE_P_COLS_RANGE"); + if (rows_value) rows = parsePdimRange("CUDECOMP_AUTOTUNE_P_ROW_RANGE"); + if (cols_value) cols = parsePdimRange("CUDECOMP_AUTOTUNE_P_COL_RANGE"); candidates.erase(std::remove_if(candidates.begin(), candidates.end(), [&](const auto& pdims) { return pdims[0] < rows.first || pdims[0] > rows.second || pdims[1] < cols.first || diff --git a/tests/ctest/api_tests.cc b/tests/ctest/api_tests.cc index 24327a2f..fb975dff 100644 --- a/tests/ctest/api_tests.cc +++ b/tests/ctest/api_tests.cc @@ -405,20 +405,20 @@ TEST(AutotuneCandidateFilterTest, RejectsMalformedAndEmptyBackendSets) { TEST(AutotuneCandidateFilterTest, FiltersProcessGridRanges) { { - ScopedEnvVar rows("CUDECOMP_AUTOTUNE_P_ROWS_RANGE", "2,4"); + ScopedEnvVar rows("CUDECOMP_AUTOTUNE_P_ROW_RANGE", "2,4"); auto candidates = cudecomp::getAutotunePdimCandidates(4, CUDECOMP_RANK_ORDER_ROW_MAJOR); std::vector> expected{{4, 1}, {2, 2}}; expectSameCandidates(candidates, expected); } { - ScopedEnvVar cols("CUDECOMP_AUTOTUNE_P_COLS_RANGE", "2,4"); + ScopedEnvVar cols("CUDECOMP_AUTOTUNE_P_COL_RANGE", "2,4"); auto candidates = cudecomp::getAutotunePdimCandidates(4, CUDECOMP_RANK_ORDER_ROW_MAJOR); std::vector> expected{{2, 2}, {1, 4}}; expectSameCandidates(candidates, expected); } { - ScopedEnvVar rows("CUDECOMP_AUTOTUNE_P_ROWS_RANGE", "2,2"); - ScopedEnvVar cols("CUDECOMP_AUTOTUNE_P_COLS_RANGE", "2,2"); + ScopedEnvVar rows("CUDECOMP_AUTOTUNE_P_ROW_RANGE", "2,2"); + ScopedEnvVar cols("CUDECOMP_AUTOTUNE_P_COL_RANGE", "2,2"); EXPECT_EQ((std::vector>{{2, 2}}), cudecomp::getAutotunePdimCandidates(4, CUDECOMP_RANK_ORDER_ROW_MAJOR)); } @@ -426,11 +426,11 @@ TEST(AutotuneCandidateFilterTest, FiltersProcessGridRanges) { TEST(AutotuneCandidateFilterTest, RejectsMalformedAndEmptyProcessGridRanges) { { - ScopedEnvVar rows("CUDECOMP_AUTOTUNE_P_ROWS_RANGE", "2,1"); + ScopedEnvVar rows("CUDECOMP_AUTOTUNE_P_ROW_RANGE", "2,1"); EXPECT_THROW(cudecomp::getAutotunePdimCandidates(4, CUDECOMP_RANK_ORDER_ROW_MAJOR), cudecomp::InvalidUsage); } { - ScopedEnvVar cols("CUDECOMP_AUTOTUNE_P_COLS_RANGE", "3,3"); + ScopedEnvVar cols("CUDECOMP_AUTOTUNE_P_COL_RANGE", "3,3"); EXPECT_THROW(cudecomp::getAutotunePdimCandidates(4, CUDECOMP_RANK_ORDER_ROW_MAJOR), cudecomp::InvalidUsage); } } @@ -1095,8 +1095,8 @@ TEST_F(ApiGridDescCreateTest, IgnoresAutotuneEnvironmentFiltersForFixedSelection auto options = fastAutotuneOptions(); ScopedEnvVar transpose_backends("CUDECOMP_AUTOTUNE_TRANSPOSE_BACKENDS", "invalid"); ScopedEnvVar halo_backends("CUDECOMP_AUTOTUNE_HALO_BACKENDS", "invalid"); - ScopedEnvVar rows("CUDECOMP_AUTOTUNE_P_ROWS_RANGE", "invalid"); - ScopedEnvVar cols("CUDECOMP_AUTOTUNE_P_COLS_RANGE", "invalid"); + ScopedEnvVar rows("CUDECOMP_AUTOTUNE_P_ROW_RANGE", "invalid"); + ScopedEnvVar cols("CUDECOMP_AUTOTUNE_P_COL_RANGE", "invalid"); cudecompGridDesc_t grid_desc = nullptr; CHECK_CUDECOMP_GLOBAL(active_comm_, cudecompGridDescCreate(handle_, &grid_desc, &config, &options));