From c366889911e58cef232f0068e4d12b17664a0e21 Mon Sep 17 00:00:00 2001 From: yboucher Date: Wed, 19 Aug 2026 00:28:50 -0700 Subject: [PATCH] exclude mpdlp files behind a flag --- build.sh | 8 +- cpp/CMakeLists.txt | 98 +++++++++++-------- cpp/src/pdlp/CMakeLists.txt | 15 ++- cpp/src/pdlp/pdhg.cu | 26 +++++ cpp/src/pdlp/pdlp.cu | 49 +++++++--- cpp/src/pdlp/pdlp.cuh | 15 ++- .../restart_strategy/pdlp_restart_strategy.cu | 19 +++- cpp/src/pdlp/solve.cu | 27 ++++- cpp/src/pdlp/solve.cuh | 2 + .../convergence_information.cu | 14 +++ .../convergence_information.hpp | 4 + .../termination_strategy.cu | 4 + cpp/tests/linear_programming/CMakeLists.txt | 10 +- 13 files changed, 218 insertions(+), 73 deletions(-) diff --git a/build.sh b/build.sh index edb2d696ea..d1303e9bb1 100755 --- a/build.sh +++ b/build.sh @@ -14,7 +14,7 @@ ARGS=$* REPODIR=$(cd "$(dirname "$0")"; pwd) LIBCUOPT_BUILD_DIR=${LIBCUOPT_BUILD_DIR:=${REPODIR}/cpp/build} -VALIDARGS="clean codegen libcuopt cuopt_grpc_server cuopt cuopt_server cuopt_sh_client docs deb -a -b -g -fsanitize -tsan -msan -v -l= --verbose-pdlp --build-lp-only --no-fetch-rapids --skip-c-python-adapters --skip-tests-build --skip-routing-build --skip-grpc-build --skip-fatbin-write --host-lineinfo --split-compile [--cmake-args=\\\"\\\"] [--cache-tool=] --install --allgpuarch --ci-only-arch --show_depr_warn -h --help" +VALIDARGS="clean codegen libcuopt cuopt_grpc_server cuopt cuopt_server cuopt_sh_client docs deb -a -b -g -fsanitize -tsan -msan -v -l= --verbose-pdlp --build-lp-only --no-fetch-rapids --skip-c-python-adapters --skip-tests-build --skip-routing-build --skip-grpc-build --skip-distributed-pdlp-build --skip-fatbin-write --host-lineinfo --split-compile [--cmake-args=\\\"\\\"] [--cache-tool=] --install --allgpuarch --ci-only-arch --show_depr_warn -h --help" HELP="$0 [ ...] [ ...] where is: clean - remove all existing build artifacts and configuration (start over) @@ -43,6 +43,7 @@ HELP="$0 [ ...] [ ...] --skip-tests-build - disable building of all tests --skip-routing-build - skip building routing components --skip-grpc-build - skip building gRPC and protobuf components (auto-enabled with -tsan) + --skip-distributed-pdlp-build - skip distributed PDLP, NCCL, and KaMinPar --skip-fatbin-write - skip the fatbin write --host-lineinfo - build with debug line information for host code --split-compile - opt in to nvcc split compilation; builds may be nondeterministic @@ -84,6 +85,7 @@ SKIP_C_PYTHON_ADAPTERS=0 SKIP_TESTS_BUILD=0 SKIP_ROUTING_BUILD=0 SKIP_GRPC_BUILD=0 +SKIP_DISTRIBUTED_PDLP_BUILD=0 WRITE_FATBIN=1 HOST_LINEINFO=0 CACHE_ARGS=() @@ -255,6 +257,9 @@ fi if hasArg --skip-grpc-build; then SKIP_GRPC_BUILD=1 fi +if hasArg --skip-distributed-pdlp-build; then + SKIP_DISTRIBUTED_PDLP_BUILD=1 +fi if hasArg --skip-fatbin-write; then WRITE_FATBIN=0 fi @@ -399,6 +404,7 @@ if buildAll || hasArg libcuopt || hasArg cuopt_grpc_server; then -DBUILD_TESTS=$((1 - ${SKIP_TESTS_BUILD})) \ -DSKIP_ROUTING_BUILD=${SKIP_ROUTING_BUILD} \ -DSKIP_GRPC_BUILD=${SKIP_GRPC_BUILD} \ + -DSKIP_DISTRIBUTED_PDLP_BUILD=${SKIP_DISTRIBUTED_PDLP_BUILD} \ -DWRITE_FATBIN=${WRITE_FATBIN} \ -DHOST_LINEINFO=${HOST_LINEINFO} \ -DPARALLEL_LEVEL="${PARALLEL_LEVEL}" \ diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 1bef3a15cd..c07cef3926 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -52,6 +52,7 @@ option(BUILD_LP_ONLY "Build only linear programming components, exclude routing option(SKIP_C_PYTHON_ADAPTERS "Skip building C and Python adapter files (cython_solve.cu and cuopt_c.cpp)" OFF) option(SKIP_ROUTING_BUILD "Skip building routing components" OFF) option(SKIP_GRPC_BUILD "Skip building gRPC and protobuf components" OFF) +option(SKIP_DISTRIBUTED_PDLP_BUILD "Skip distributed PDLP, NCCL, and KaMinPar" OFF) option(WRITE_FATBIN "Enable fatbin writing" ON) option(HOST_LINEINFO "Build with debug line information for host code" OFF) @@ -62,9 +63,14 @@ message(VERBOSE "cuOpt: Disable OpenMP: ${DISABLE_OPENMP}") message(VERBOSE "cuOpt: Build LP-only mode: ${BUILD_LP_ONLY}") message(VERBOSE "cuOpt: Skip C/Python adapters: ${SKIP_C_PYTHON_ADAPTERS}") message(VERBOSE "cuOpt: Skip routing build: ${SKIP_ROUTING_BUILD}") +message(VERBOSE "cuOpt: Skip distributed PDLP build: ${SKIP_DISTRIBUTED_PDLP_BUILD}") message(VERBOSE "cuOpt: Build with debug line information for host code: ${HOST_LINEINFO}") message(VERBOSE "cuOpt: fatbin: ${WRITE_FATBIN}") +if (NOT SKIP_DISTRIBUTED_PDLP_BUILD) + add_compile_definitions(CUOPT_ENABLE_DISTRIBUTED_PDLP) +endif () + # ################################################################################################## # - compiler options ------------------------------------------------------------------------------ @@ -333,39 +339,41 @@ find_package(CUDSS REQUIRED) # - NCCL (multi-GPU distributed PDLP) ------------------------------------------------------------- # NCCL is shipped via the conda env; no canonical CMake config target, so look it # up by name in the standard lib paths (plus CONDA_PREFIX as a hint). -set(NCCL_HINT_PREFIXES "") -if (DEFINED ENV{CONDA_PREFIX} AND NOT "$ENV{CONDA_PREFIX}" STREQUAL "") - list(APPEND NCCL_HINT_PREFIXES "$ENV{CONDA_PREFIX}") -endif () -find_path(NCCL_INCLUDE_DIR - NAMES nccl.h - HINTS ${NCCL_HINT_PREFIXES} - PATH_SUFFIXES include -) -find_library(NCCL_LIBRARY - NAMES nccl - HINTS ${NCCL_HINT_PREFIXES} - PATH_SUFFIXES lib lib64 -) -if (NOT NCCL_INCLUDE_DIR OR NOT NCCL_LIBRARY) - message(FATAL_ERROR "NCCL not found. Looked in ${NCCL_HINT_PREFIXES}. Install nccl-dev / libnccl-dev in the active env.") -endif () -add_library(nccl_external UNKNOWN IMPORTED GLOBAL) -set_target_properties(nccl_external PROPERTIES - IMPORTED_LOCATION "${NCCL_LIBRARY}" - INTERFACE_INCLUDE_DIRECTORIES "${NCCL_INCLUDE_DIR}" -) -message(STATUS "Using NCCL: ${NCCL_LIBRARY}") - -# ################################################################################################## -# - KaMinPar (multi-threaded partitioning for distributed PDLP) ------------------------------------ -# Brought in the RAPIDS way (rapids_cpm_find): uses an installed KaMinPar (deb/rpm/conda, -# discovered via its CMake config) if present, otherwise builds the pinned source via CPM. -include(cmake/thirdparty/get_kaminpar.cmake) -if (NOT TARGET KaMinPar::KaMinPar) - message(FATAL_ERROR "KaMinPar::KaMinPar was not made available by get_kaminpar.cmake") +if (NOT SKIP_DISTRIBUTED_PDLP_BUILD) + set(NCCL_HINT_PREFIXES "") + if (DEFINED ENV{CONDA_PREFIX} AND NOT "$ENV{CONDA_PREFIX}" STREQUAL "") + list(APPEND NCCL_HINT_PREFIXES "$ENV{CONDA_PREFIX}") + endif () + find_path(NCCL_INCLUDE_DIR + NAMES nccl.h + HINTS ${NCCL_HINT_PREFIXES} + PATH_SUFFIXES include + ) + find_library(NCCL_LIBRARY + NAMES nccl + HINTS ${NCCL_HINT_PREFIXES} + PATH_SUFFIXES lib lib64 + ) + if (NOT NCCL_INCLUDE_DIR OR NOT NCCL_LIBRARY) + message(FATAL_ERROR "NCCL not found. Looked in ${NCCL_HINT_PREFIXES}. Install nccl-dev / libnccl-dev in the active env.") + endif () + add_library(nccl_external UNKNOWN IMPORTED GLOBAL) + set_target_properties(nccl_external PROPERTIES + IMPORTED_LOCATION "${NCCL_LIBRARY}" + INTERFACE_INCLUDE_DIRECTORIES "${NCCL_INCLUDE_DIR}" + ) + message(STATUS "Using NCCL: ${NCCL_LIBRARY}") + + # ################################################################################################## + # - KaMinPar (multi-threaded partitioning for distributed PDLP) ------------------------------------ + # Brought in the RAPIDS way (rapids_cpm_find): uses an installed KaMinPar (deb/rpm/conda, + # discovered via its CMake config) if present, otherwise builds the pinned source via CPM. + include(cmake/thirdparty/get_kaminpar.cmake) + if (NOT TARGET KaMinPar::KaMinPar) + message(FATAL_ERROR "KaMinPar::KaMinPar was not made available by get_kaminpar.cmake") + endif () + message(STATUS "Using KaMinPar (distributed PDLP graph partitioner)") endif () -message(STATUS "Using KaMinPar (distributed PDLP graph partitioner)") # ################################################################################################## # - gRPC and Protobuf setup ----------------------------------------------------------------------- @@ -698,8 +706,6 @@ add_dependencies(cuopt_objs PSLP) # "target KaMinPar is not in any export set"). libKaMinPar.a is self-contained (the # kaminpar-common OBJECT lib is archived into it); we only need its public headers # (, which pulls in stdlib + TBB) at compile time. -target_include_directories(cuopt_objs SYSTEM PRIVATE - $) # partitioner.cpp includes . Because KaMinPar is linked by file, cuopt # does NOT inherit KaMinPar's PUBLIC compile definitions, so we must replicate the ones that # affect what compiles to: @@ -709,11 +715,15 @@ target_include_directories(cuopt_objs SYSTEM PRIVATE # with the KAMINPAR_64BIT_* options in get_kaminpar.cmake; # otherwise the public API types (e.g. copy_graph's span<> # widths) mismatch the prebuilt archive => undefined refs. -target_compile_definitions(cuopt_objs PRIVATE TBB_PREVIEW_GLOBAL_CONTROL KAMINPAR_64BIT_EDGE_IDS) # Real (non-imported) target only exists when KaMinPar is built from source via CPM; # force it to build before the objects that include its headers since it is EXCLUDE_FROM_ALL. -if (TARGET KaMinPar) +if (NOT SKIP_DISTRIBUTED_PDLP_BUILD) + target_include_directories(cuopt_objs SYSTEM PRIVATE + $) + target_compile_definitions(cuopt_objs PRIVATE TBB_PREVIEW_GLOBAL_CONTROL KAMINPAR_64BIT_EDGE_IDS) + if (TARGET KaMinPar) add_dependencies(cuopt_objs KaMinPar) + endif () endif () # ################################################################################################## @@ -790,7 +800,6 @@ if (BUILD_TESTS) OpenMP::OpenMP_CXX PRIVATE ${CUOPT_PRIVATE_CUDA_LIBS} - nccl_external $<$:protobuf::libprotobuf> $<$:gRPC::grpc++> ) @@ -813,9 +822,11 @@ if (BUILD_TESTS) ) target_link_libraries(cuopt_static PRIVATE $) add_dependencies(cuopt_static PSLP) - target_link_libraries(cuopt_static PRIVATE $) - if (TARGET KaMinPar) - add_dependencies(cuopt_static KaMinPar) + if (NOT SKIP_DISTRIBUTED_PDLP_BUILD) + target_link_libraries(cuopt_static PRIVATE nccl_external $) + if (TARGET KaMinPar) + add_dependencies(cuopt_static KaMinPar) + endif () endif () add_subdirectory(tests) endif (BUILD_TESTS) @@ -848,7 +859,6 @@ target_link_libraries(cuopt ${CUDSS_LIB_FILE} PRIVATE ${CUOPT_PRIVATE_CUDA_LIBS} - nccl_external simde::simde OpenMP::OpenMP_CXX OpenMP::OpenMP_CUDA @@ -857,9 +867,11 @@ target_link_libraries(cuopt ) target_link_libraries(cuopt PRIVATE $) add_dependencies(cuopt PSLP) -target_link_libraries(cuopt PRIVATE $) -if (TARGET KaMinPar) +if (NOT SKIP_DISTRIBUTED_PDLP_BUILD) + target_link_libraries(cuopt PRIVATE nccl_external $) + if (TARGET KaMinPar) add_dependencies(cuopt KaMinPar) + endif () endif () # Propagate compile definitions that consumers need when including cuopt headers. # These were on cuopt directly before the cuopt_objs refactor; $ diff --git a/cpp/src/pdlp/CMakeLists.txt b/cpp/src/pdlp/CMakeLists.txt index 2f90f94872..d82902534a 100644 --- a/cpp/src/pdlp/CMakeLists.txt +++ b/cpp/src/pdlp/CMakeLists.txt @@ -29,13 +29,18 @@ set(LP_CORE_FILES ${CMAKE_CURRENT_SOURCE_DIR}/termination_strategy/convergence_information.cu ${CMAKE_CURRENT_SOURCE_DIR}/optimal_batch_size_handler/optimal_batch_size_handler.cu ${CMAKE_CURRENT_SOURCE_DIR}/utilities/ping_pong_graph.cu - ${CMAKE_CURRENT_SOURCE_DIR}/distributed_pdlp/shard.cu - ${CMAKE_CURRENT_SOURCE_DIR}/distributed_pdlp/multi_gpu_engine.cu - ${CMAKE_CURRENT_SOURCE_DIR}/distributed_pdlp/distributed_algorithms.cu - ${CMAKE_CURRENT_SOURCE_DIR}/distributed_pdlp/distributed_utils.cu - ${CMAKE_CURRENT_SOURCE_DIR}/distributed_pdlp/partitioner.cpp ) +if(NOT SKIP_DISTRIBUTED_PDLP_BUILD) + list(APPEND LP_CORE_FILES + ${CMAKE_CURRENT_SOURCE_DIR}/distributed_pdlp/shard.cu + ${CMAKE_CURRENT_SOURCE_DIR}/distributed_pdlp/multi_gpu_engine.cu + ${CMAKE_CURRENT_SOURCE_DIR}/distributed_pdlp/distributed_algorithms.cu + ${CMAKE_CURRENT_SOURCE_DIR}/distributed_pdlp/distributed_utils.cu + ${CMAKE_CURRENT_SOURCE_DIR}/distributed_pdlp/partitioner.cpp + ) +endif() + # C and Python adapter files set(LP_ADAPTER_FILES ${CMAKE_CURRENT_SOURCE_DIR}/utilities/cython_solve.cu diff --git a/cpp/src/pdlp/pdhg.cu b/cpp/src/pdlp/pdhg.cu index 69ac810455..381bafa0ff 100644 --- a/cpp/src/pdlp/pdhg.cu +++ b/cpp/src/pdlp/pdhg.cu @@ -4,7 +4,9 @@ * SPDX-License-Identifier: Apache-2.0 */ /* clang-format on */ +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP #include +#endif #include #include #include @@ -510,10 +512,12 @@ void pdhg_solver_t::compute_At_y() // Multi-GPU dispatch: when the master pdhg has an engine, drive halo // exchange + per-shard SpMV via the engine. +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (is_distributed_master()) { mgpu_engine_->distributed_compute_At_y(); return; } +#endif if (!batch_mode_) { if constexpr (std::is_same_v) { @@ -568,10 +572,12 @@ void pdhg_solver_t::compute_A_x() // Multi-GPU dispatch: see compute_At_y. The engine halo-updates the // reflected_primal vector (the buffer this SpMV reads) and then drives // per-shard local cusparse SpMV. +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (is_distributed_master()) { mgpu_engine_->distributed_compute_A_x(); return; } +#endif if (!batch_mode_) { if constexpr (std::is_same_v) { @@ -1229,23 +1235,29 @@ void pdhg_solver_t::compute_next_primal_dual_solution_reflected( using f_t2 = typename type_2::type; +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (is_distributed_master()) { mgpu_engine_->sync_await_shards(stream_view_); } +#endif // Compute next primal solution reflected. if (should_major) { graph_all.run(should_major, [&]() { +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP // Adds all the shards streams into the graph capture if (is_distributed_master()) { mgpu_engine_->graph_capture_fork_to_shards(stream_view_); } +#endif compute_At_y(); if (is_distributed_master()) { +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP mgpu_engine_->for_each_shard([](auto& shard) { auto& sub_pdlp = *shard.sub_pdlp; sub_pdlp.pdhg_solver_.primal_reflected_major_projection_transform( sub_pdlp.get_primal_step_size()); }); +#endif } else if (!batch_mode_) { primal_reflected_major_projection_transform(primal_step_size); } else { @@ -1307,11 +1319,13 @@ void pdhg_solver_t::compute_next_primal_dual_solution_reflected( compute_A_x(); if (is_distributed_master()) { +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP mgpu_engine_->for_each_shard([](auto& shard) { auto& sub_pdlp = *shard.sub_pdlp; sub_pdlp.pdhg_solver_.dual_reflected_major_projection_transform( sub_pdlp.get_dual_step_size()); }); +#endif } else if (!batch_mode_) { dual_reflected_major_projection_transform(dual_step_size); } else { @@ -1338,13 +1352,17 @@ void pdhg_solver_t::compute_next_primal_dual_solution_reflected( // Multi-GPU: close the fork by joining every shard stream back into // the master stream so cudaStreamEndCapture sees a single graph // spanning all streams. +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (is_distributed_master()) { mgpu_engine_->graph_capture_join_from_shards(stream_view_); } +#endif }); } else { graph_all.run(should_major, [&]() { +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP // Same reason as above, adds all the shards streams into the graph capture if (is_distributed_master()) { mgpu_engine_->graph_capture_fork_to_shards(stream_view_); } +#endif // Compute next primal compute_At_y(); @@ -1358,11 +1376,13 @@ void pdhg_solver_t::compute_next_primal_dual_solution_reflected( #endif if (is_distributed_master()) { +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP mgpu_engine_->for_each_shard([](auto& shard) { auto& sub_pdlp = *shard.sub_pdlp; sub_pdlp.pdhg_solver_.primal_reflected_projection_transform( sub_pdlp.get_primal_step_size()); }); +#endif } else if (!batch_mode_) { primal_reflected_projection_transform(primal_step_size); } else { @@ -1425,10 +1445,12 @@ void pdhg_solver_t::compute_next_primal_dual_solution_reflected( compute_A_x(); if (is_distributed_master()) { +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP mgpu_engine_->for_each_shard([](auto& shard) { auto& sub_pdlp = *shard.sub_pdlp; sub_pdlp.pdhg_solver_.dual_reflected_projection_transform(sub_pdlp.get_dual_step_size()); }); +#endif } else if (!batch_mode_) { dual_reflected_projection_transform(dual_step_size); } else { @@ -1449,12 +1471,16 @@ void pdhg_solver_t::compute_next_primal_dual_solution_reflected( print("reflected_dual_", reflected_dual_); #endif +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (is_distributed_master()) { mgpu_engine_->graph_capture_join_from_shards(stream_view_); } +#endif }); } // sync to master stream after the graph is captured +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (is_distributed_master()) { mgpu_engine_->sync_await_master(stream_view_); } +#endif } template diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index 62b50825e7..44984c047e 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -11,8 +11,10 @@ #include #include +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP #include #include +#endif #include #include #include @@ -20,7 +22,9 @@ #include #include #include "cuopt/mathematical_optimization/pdlp/solver_solution.hpp" +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP #include "distributed_pdlp/multi_gpu_engine.hpp" +#endif #include #include @@ -384,6 +388,7 @@ pdlp_solver_t::pdlp_solver_t(mip::problem_t& op_problem, } } +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP // ============================================================================ // Distributed multi-GPU ctor. // needs placeholder_problem to be a shape-0 problem @@ -560,6 +565,7 @@ pdlp_solver_t::pdlp_solver_t( current_termination_strategy_.get_convergence_information().distributed_init_l2_norms( *multi_gpu_engine); } +#endif template void pdlp_solver_t::set_initial_primal_weight(f_t initial_primal_weight) @@ -2235,17 +2241,21 @@ void pdlp_solver_t::compute_fixed_error(std::vector& has_restarte // Computing the deltas (delta = reflected - current) // TODO batch mdoe: this only works if everyone restarts +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (is_distributed_master()) { multi_gpu_engine->for_each_shard([](auto& shard) { compute_primal_dual_deltas(shard.sub_pdlp->pdhg_solver_, shard.stream.view()); }); - } else { + } else +#endif + { compute_primal_dual_deltas(pdhg_solver_, stream_view_); } auto& cusparse_view = pdhg_solver_.get_cusparse_view(); // Distributed compute_fixed_error second part +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (is_distributed_master()) { // SpMV is the first operation in compute_interaction_and_movement so we can do halo before and // call it naturally we then reduce the local dot products @@ -2281,7 +2291,9 @@ void pdlp_solver_t::compute_fixed_error(std::vector& has_restarte }); multi_gpu_engine->allreduce_sum_inplace_to_master( [](auto& sp) -> f_t* { return sp.step_size_strategy_.get_norm_squared_delta_dual().data(); }); - } else { + } else +#endif + { // Sync to make sure all previous cuSparse operations are finished before setting the // potential_next_dual_solution RAFT_CUDA_TRY(cudaStreamSynchronize(stream_view_)); @@ -2608,6 +2620,7 @@ optimization_problem_solution_t pdlp_solver_t::run_solver(co compute_initial_primal_weight(); // Distributed counterpart of the single-GPU, happens later in the single-GPU path. +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (settings_.use_distributed_pdlp) { step_size_strategy_.get_primal_and_dual_stepsizes(primal_step_size_, dual_step_size_); multi_gpu_engine->for_each_shard([&](auto& shard) { @@ -2618,6 +2631,7 @@ optimization_problem_solution_t pdlp_solver_t::run_solver(co multi_gpu_engine->sync_await_shards(stream_view_); handle_ptr_->sync_stream(stream_view_); } +#endif // Everything below (seed-from-settings, initial_k, get_primal_and_dual_stepsizes, // initial primal/dual, projection, transpose, verbose prints, log header) @@ -2929,6 +2943,7 @@ optimization_problem_solution_t pdlp_solver_t::run_solver(co initial_scaling_strategy_.unscale_solutions(pdhg_solver_.get_primal_solution(), pdhg_solver_.get_dual_solution()); } else { +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (is_distributed_master()) { // The only branch in cuPDLPx (Stable3) multi_gpu_engine->for_each_shard([&](auto& shard) { @@ -2938,7 +2953,9 @@ optimization_problem_solution_t pdlp_solver_t::run_solver(co sub.pdhg_solver_.get_potential_next_dual_solution(), sub.pdhg_solver_.get_dual_slack()); }); - } else { + } else +#endif + { initial_scaling_strategy_.unscale_solutions( pdhg_solver_.get_potential_next_primal_solution(), pdhg_solver_.get_potential_next_dual_solution(), @@ -2975,6 +2992,7 @@ optimization_problem_solution_t pdlp_solver_t::run_solver(co initial_scaling_strategy_.scale_solutions(pdhg_solver_.get_primal_solution(), pdhg_solver_.get_dual_solution()); } else { +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (is_distributed_master()) { // The only branch in cuPDLPx (Stable3) multi_gpu_engine->for_each_shard([&](auto& shard) { @@ -2984,7 +3002,9 @@ optimization_problem_solution_t pdlp_solver_t::run_solver(co sub.pdhg_solver_.get_potential_next_dual_solution(), sub.pdhg_solver_.get_dual_slack()); }); - } else { + } else +#endif + { initial_scaling_strategy_.scale_solutions( pdhg_solver_.get_potential_next_primal_solution(), pdhg_solver_.get_potential_next_dual_solution(), @@ -3094,11 +3114,13 @@ optimization_problem_solution_t pdlp_solver_t::run_solver(co ++internal_solver_iterations_; if (settings_.hyper_params.never_restart_to_average) { restart_strategy_.increment_iteration_since_last_restart(); +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (is_distributed_master()) { multi_gpu_engine->for_each_shard([&](auto& shard) { shard.sub_pdlp->restart_strategy_.increment_iteration_since_last_restart(); }); } +#endif } } return optimization_problem_solution_t{pdlp_termination_status_t::NumericalError, @@ -3163,10 +3185,12 @@ void pdlp_solver_t::halpern_update() { raft::common::nvtx::range fun_scope("halpern_update"); +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (is_distributed_master()) { multi_gpu_engine->for_each_shard([&](auto& shard) { shard.sub_pdlp->halpern_update(); }); return; } +#endif // TODO later batch mode: handle if element in the batch have different one if restart per climber const f_t weight = f_t(restart_strategy_.weighted_average_solution_.get_iterations_since_last_restart() + 1) / @@ -3248,6 +3272,7 @@ void pdlp_solver_t::scale_problem() { // Scale problem then free scratch buffers raft::common::nvtx::range fun_scope("pdlp_solver_t::scale_problem"); +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (is_distributed_master()) { multi_gpu_engine->distributed_scaling(settings_.hyper_params, primal_size_h_, inside_mip_); @@ -3257,7 +3282,9 @@ void pdlp_solver_t::scale_problem() scaling.get_iteration_variable_scaling().resize(0, shard.stream.view()); scaling.get_iteration_constraint_matrix_scaling().resize(0, shard.stream.view()); }); - } else { + } else +#endif + { initial_scaling_strategy_.scale_problem(); // Free scratch: no further scaling passes happen after this point. @@ -3270,6 +3297,7 @@ template void pdlp_solver_t::create_spmv_op_plans() { raft::common::nvtx::range fun_scope("pdlp_solver_t::create_spmv_op_plans"); +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (is_distributed_master()) { // Distributed path: fan out the same per-shard cusparse_view call the // single-GPU path would make. @@ -3281,6 +3309,7 @@ void pdlp_solver_t::create_spmv_op_plans() handle_ptr_->sync_stream(stream_view_); return; } +#endif if constexpr (std::is_same_v) { if (!batch_mode_ && !pdhg_solver_.get_cusparse_view().mixed_precision_enabled_) { pdhg_solver_.get_cusparse_view().create_spmv_op_plans( @@ -3300,6 +3329,7 @@ void pdlp_solver_t::compute_initial_step_size() constexpr int max_iterations = 5000; constexpr f_t tolerance = f_t{1e-4}; +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (is_distributed_master()) { // Distributed dispatch: everything (sigma_max, deriving primal/dual // step sizes from master's current primal_weight_, seeding master + @@ -3308,6 +3338,7 @@ void pdlp_solver_t::compute_initial_step_size() settings_.hyper_params, dual_size_h_, scaling_factor, max_iterations, tolerance); return; } +#endif if (!settings_.hyper_params.initial_step_size_max_singular_value) { // set stepsize relative to maximum absolute value of A @@ -3484,6 +3515,7 @@ void pdlp_solver_t::compute_initial_primal_weight() { raft::common::nvtx::range fun_scope("compute_initial_primal_weight"); +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (is_distributed_master()) { // Distributed dispatch: // - short-circuit -> 1 @@ -3492,6 +3524,7 @@ void pdlp_solver_t::compute_initial_primal_weight() multi_gpu_engine->distributed_compute_initial_primal_weight(settings_.hyper_params); return; } +#endif // Here we use the combined bounds of the op_problem_scaled which may or may not be scaled yet // based on pdlp config @@ -3585,12 +3618,6 @@ pdlp_restart_strategy_t& pdlp_solver_t::get_restart_strategy return restart_strategy_; } -template -bool pdlp_solver_t::is_distributed_master() const -{ - return multi_gpu_engine.has_value(); -} - #if MIP_INSTANTIATE_FLOAT || PDLP_INSTANTIATE_FLOAT template class pdlp_solver_t; diff --git a/cpp/src/pdlp/pdlp.cuh b/cpp/src/pdlp/pdlp.cuh index 60b2c1c354..04633af5ae 100644 --- a/cpp/src/pdlp/pdlp.cuh +++ b/cpp/src/pdlp/pdlp.cuh @@ -13,7 +13,9 @@ #include #include +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP #include +#endif #include #include #include @@ -67,10 +69,12 @@ class pdlp_solver_t { bool is_batch_mode = false, bool is_distributed_sub_pdlp = false); +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP // Distributed Solver Constructor pdlp_solver_t(mip::problem_t& placeholder_problem, cuopt::mathematical_optimization::io::mps_data_model_t const& mps, pdlp_solver_settings_t const& settings); +#endif optimization_problem_solution_t run_solver(const timer_t& timer); @@ -215,10 +219,12 @@ class pdlp_solver_t { pdlp::adaptive_step_size_strategy_t step_size_strategy_; public: +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP // std::optional because multi_gpu_engine_t is non-default-constructible // (collectively bootstraps NCCL, owns RMM resources). Stays nullopt in // single-GPU mode; emplaced by the multi-GPU ctor. std::optional> multi_gpu_engine; +#endif // Inner solver pdlp::pdhg_solver_t pdhg_solver_; @@ -227,7 +233,14 @@ class pdlp_solver_t { // This solver is the distributed-PDLP master orchestrator iff it owns the // multi-GPU engine. Shards (sub-solvers) leave the optional empty -> false. // Single-GPU PDLP reports false. - bool is_distributed_master() const; + bool is_distributed_master() const + { +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP + return multi_gpu_engine.has_value(); +#else + return false; +#endif + } private: void compute_fixed_error(std::vector& has_restarted); diff --git a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu index 5d3258502a..c442542fd7 100644 --- a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu +++ b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu @@ -8,7 +8,9 @@ #include #include +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP #include +#endif #include #include #include @@ -894,6 +896,7 @@ void pdlp_restart_strategy_t::cupdlpx_restart( "If any, all should be true"); // Computing the distributed deltas +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (auto* engine = pdhg_solver.get_mgpu_engine()) { engine->for_each_shard([&](auto& shard) { auto& sub = *shard.sub_pdlp; @@ -907,7 +910,9 @@ void pdlp_restart_strategy_t::cupdlpx_restart( engine->allreduce_sum_inplace_to_master([](pdlp_solver_t& sp) -> f_t* { return sp.get_restart_strategy().last_restart_duality_gap_.dual_distance_traveled_.data(); }); - } else { + } else +#endif + { primal_dual_distance_squared_moved_from_last_restart_period( pdhg_solver, primal_size_h_, dual_size_h_); } @@ -957,6 +962,7 @@ void pdlp_restart_strategy_t::cupdlpx_restart( hyper_params_.restart_k_i, hyper_params_.restart_k_d, hyper_params_.restart_i_smooth); +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (auto* engine = pdhg_solver.get_mgpu_engine()) { engine->set_scalar_on_master_and_shards( primal_weight_value, [](auto& sp) { return sp.get_primal_weight().data(); }); @@ -966,7 +972,9 @@ void pdlp_restart_strategy_t::cupdlpx_restart( dual_step_size_value, [](auto& sp) { return sp.get_dual_step_size().data(); }); engine->set_scalar_on_master_and_shards( best_primal_weight_value, [](auto& sp) { return sp.get_best_primal_weight().data(); }); - } else { + } else +#endif + { primal_weight.set_element_async(0, primal_weight_value, stream_view_); primal_step_size.set_element_async(0, primal_step_size_value, stream_view_); dual_step_size.set_element_async(0, dual_step_size_value, stream_view_); @@ -998,13 +1006,16 @@ void pdlp_restart_strategy_t::cupdlpx_restart( stream); }; +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (auto* engine = pdhg_solver.get_mgpu_engine()) { engine->for_each_shard([&](auto& shard) { auto& sub = *shard.sub_pdlp; commit_potential_next_as_last_restart( sub.get_restart_strategy(), sub.pdhg_solver_, shard.stream.view()); }); - } else { + } else +#endif + { commit_potential_next_as_last_restart(*this, pdhg_solver, stream_view_); } @@ -1022,12 +1033,14 @@ void pdlp_restart_strategy_t::cupdlpx_restart( last_trial_fixed_point_error_[i] = std::numeric_limits::infinity(); } +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (auto* engine = pdhg_solver.get_mgpu_engine()) { engine->for_each_shard([&](auto& shard) { shard.sub_pdlp->get_restart_strategy() .weighted_average_solution_.iterations_since_last_restart_ = 0; }); } +#endif } template diff --git a/cpp/src/pdlp/solve.cu b/cpp/src/pdlp/solve.cu index 80b3da2c18..fe87682abc 100644 --- a/cpp/src/pdlp/solve.cu +++ b/cpp/src/pdlp/solve.cu @@ -2399,6 +2399,13 @@ optimization_problem_solution_t solve_lp( bool problem_checking, bool use_pdlp_solver_mode) { +#ifndef CUOPT_ENABLE_DISTRIBUTED_PDLP + cuopt_expects( + !settings.use_distributed_pdlp && + !(settings.method == method_t::PDLP && (settings.num_gpus == -1 || settings.num_gpus > 1)), + error_type_t::ValidationError, + "Distributed PDLP is not available in this build"); +#else if (settings.use_distributed_pdlp) { return solve_lp_distributed_from_mps( handle_ptr, mps_data_model, settings, use_pdlp_solver_mode); @@ -2410,10 +2417,12 @@ optimization_problem_solution_t solve_lp( return solve_lp_distributed_from_mps( handle_ptr, mps_data_model, distributed_settings, use_pdlp_solver_mode); } +#endif auto op_problem = mps_data_model_to_optimization_problem(handle_ptr, mps_data_model); return solve_lp(op_problem, settings, problem_checking, use_pdlp_solver_mode, false); } +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP template optimization_problem_solution_t solve_lp_distributed_from_mps( raft::handle_t const* handle_ptr, @@ -2655,6 +2664,7 @@ optimization_problem_solution_t solve_lp_distributed_from_mps( return sol; } +#endif // ============================================================================ // CPU problem overloads (convert to GPU, solve, convert solution back) @@ -2746,6 +2756,17 @@ std::unique_ptr> solve_lp( return std::make_unique>(std::move(gpu_solution)); } +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP +#define INSTANTIATE_DISTRIBUTED_PDLP(F_TYPE) \ + template optimization_problem_solution_t solve_lp_distributed_from_mps( \ + raft::handle_t const* handle_ptr, \ + const cuopt::mathematical_optimization::io::mps_data_model_t& mps_data_model, \ + pdlp_solver_settings_t const& settings, \ + bool use_pdlp_solver_mode); +#else +#define INSTANTIATE_DISTRIBUTED_PDLP(F_TYPE) +#endif + #define INSTANTIATE(F_TYPE) \ template CUOPT_EXPORT optimization_problem_solution_t solve_lp( \ optimization_problem_t& op_problem, \ @@ -2805,11 +2826,7 @@ std::unique_ptr> solve_lp( template cuopt::mathematical_optimization::io::mps_data_model_t \ op_problem_to_mps_data_model(const optimization_problem_t& op_problem); \ \ - template optimization_problem_solution_t solve_lp_distributed_from_mps( \ - raft::handle_t const* handle_ptr, \ - const cuopt::mathematical_optimization::io::mps_data_model_t& mps_data_model, \ - pdlp_solver_settings_t const& settings, \ - bool use_pdlp_solver_mode); \ + INSTANTIATE_DISTRIBUTED_PDLP(F_TYPE) \ \ template void set_pdlp_solver_mode(pdlp_solver_settings_t& settings); diff --git a/cpp/src/pdlp/solve.cuh b/cpp/src/pdlp/solve.cuh index 5105d59672..e2e9aceb65 100644 --- a/cpp/src/pdlp/solve.cuh +++ b/cpp/src/pdlp/solve.cuh @@ -33,6 +33,7 @@ cuopt::mathematical_optimization::optimization_problem_solution_t solv const timer_t& timer, bool is_batch_mode = false); +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP /** * @brief Distributed-PDLP entry point that consumes the host MPS data model * directly, partitioning it across GPUs without ever materializing the @@ -73,6 +74,7 @@ solve_lp_distributed_from_mps( const cuopt::mathematical_optimization::io::mps_data_model_t& mps_data_model, pdlp_solver_settings_t const& settings, bool use_pdlp_solver_mode); +#endif /** * @brief Entry point for batch PDLP. Solves multiple LPs sharing the same constraint diff --git a/cpp/src/pdlp/termination_strategy/convergence_information.cu b/cpp/src/pdlp/termination_strategy/convergence_information.cu index f72524d076..5ade954ff2 100644 --- a/cpp/src/pdlp/termination_strategy/convergence_information.cu +++ b/cpp/src/pdlp/termination_strategy/convergence_information.cu @@ -5,7 +5,9 @@ */ /* clang-format on */ +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP #include +#endif #include #include #include @@ -215,6 +217,7 @@ void convergence_information_t::init_l2_norms() } } +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP template void convergence_information_t::distributed_init_l2_norms( multi_gpu_engine_t& engine) @@ -269,6 +272,7 @@ void convergence_information_t::distributed_init_l2_norms( }, [](pdlp_shard_t& s) -> i_t { return s.rank_data.owned_var_size; }); } +#endif // --------------------------------------------------------------------------- // init_reduction_storage: allocate and size the temporary buffers used by @@ -444,6 +448,7 @@ __global__ void compute_remaining_stats_kernel( raft::abs(convergence_information_view.dual_objective[idx]); } +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP template void convergence_information_t::distributed_compute_primal_residual_and_objective( multi_gpu_engine_t& engine, const pdlp_solver_settings_t& settings) @@ -515,6 +520,7 @@ void convergence_information_t::distributed_compute_dual_residual_and_ }); apply_dual_objective_scaling_and_offset(); } +#endif template void convergence_information_t::compute_convergence_information( @@ -558,8 +564,10 @@ void convergence_information_t::compute_convergence_information( #endif if (current_pdhg_solver.is_distributed_master()) { +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP distributed_compute_primal_residual_and_objective(*current_pdhg_solver.get_mgpu_engine(), settings); +#endif } else { compute_primal_residual( op_problem_cusparse_view_, current_pdhg_solver.get_dual_tmp_resource(), dual_iterate); @@ -572,6 +580,7 @@ void convergence_information_t::compute_convergence_information( // L2 Norm if (current_pdhg_solver.is_distributed_master()) { +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP // Fused per-shard L2 reduction + D2D mirror to master. current_pdhg_solver.get_mgpu_engine()->distributed_l2_norm_to_master( [](pdlp_solver_t& sp) -> rmm::device_uvector& { @@ -583,6 +592,7 @@ void convergence_information_t::compute_convergence_information( .l2_primal_residual_.data(); }, [](pdlp_shard_t& shard) -> i_t { return shard.rank_data.owned_cstr_size; }); +#endif } else if (!batch_mode_) { my_l2_norm(primal_residual_, l2_primal_residual_, handle_ptr_); } else { @@ -627,7 +637,9 @@ void convergence_information_t::compute_convergence_information( } if (current_pdhg_solver.is_distributed_master()) { +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP distributed_compute_dual_residual_and_objective(*current_pdhg_solver.get_mgpu_engine()); +#endif } else { compute_dual_residual(op_problem_cusparse_view_, current_pdhg_solver.get_primal_tmp_resource(), @@ -641,6 +653,7 @@ void convergence_information_t::compute_convergence_information( #endif if (current_pdhg_solver.is_distributed_master()) { +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP // Same pattern as the primal L2 above. current_pdhg_solver.get_mgpu_engine()->distributed_l2_norm_to_master( [](pdlp_solver_t& sp) -> rmm::device_uvector& { @@ -652,6 +665,7 @@ void convergence_information_t::compute_convergence_information( .l2_dual_residual_.data(); }, [](pdlp_shard_t& shard) -> i_t { return shard.rank_data.owned_var_size; }); +#endif } else if (!batch_mode_) { my_l2_norm(dual_residual_, l2_dual_residual_, handle_ptr_); } else { diff --git a/cpp/src/pdlp/termination_strategy/convergence_information.hpp b/cpp/src/pdlp/termination_strategy/convergence_information.hpp index 1bcb2fc0ab..9856cb28ad 100644 --- a/cpp/src/pdlp/termination_strategy/convergence_information.hpp +++ b/cpp/src/pdlp/termination_strategy/convergence_information.hpp @@ -66,6 +66,7 @@ class convergence_information_t { const rmm::device_uvector& get_l2_norm_primal_linear_objective() const; const rmm::device_uvector& get_l2_norm_primal_right_hand_side() const; +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP // Multi-GPU counterpart of init_l2_norms(). // - per-shard thrust::transform_reduce on the OWNED prefix, folded into // host scalars (blocking on each shard's stream); @@ -75,6 +76,7 @@ class convergence_information_t { // rhs_sum_of_squares_t (skips infinite bounds and degenerate ranges, // matching the single-GPU compute_sum_bounds semantics). void distributed_init_l2_norms(multi_gpu_engine_t& engine); +#endif struct view_t { i_t primal_size; @@ -174,12 +176,14 @@ class convergence_information_t { void compute_reduced_costs_dual_objective_contribution(); +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP // ----- Distributed-PDLP sub-steps of compute_convergence_information ----- // Halo exchange, per-shard primal/residual + partial (owned) primal/dual objective, allreduce and // apply scaling+offset. void distributed_compute_primal_residual_and_objective( multi_gpu_engine_t& engine, const pdlp_solver_settings_t& settings); void distributed_compute_dual_residual_and_objective(multi_gpu_engine_t& engine); +#endif // Ctor helpers — each handles both batch and non-batch internally. void init_objective_offsets(); diff --git a/cpp/src/pdlp/termination_strategy/termination_strategy.cu b/cpp/src/pdlp/termination_strategy/termination_strategy.cu index 13acee138c..5794a03db5 100644 --- a/cpp/src/pdlp/termination_strategy/termination_strategy.cu +++ b/cpp/src/pdlp/termination_strategy/termination_strategy.cu @@ -5,7 +5,9 @@ */ /* clang-format on */ +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP #include +#endif #include #include @@ -567,12 +569,14 @@ pdlp_termination_strategy_t::fill_return_problem_solution( } // In distributed PDLP, gather solutions from the shards to the master. +#ifdef CUOPT_ENABLE_DISTRIBUTED_PDLP if (auto* engine = current_pdhg_solver.get_mgpu_engine()) { const bool is_current_live_iterate = (&primal_iterate == ¤t_pdhg_solver.get_potential_next_primal_solution()) || (&primal_iterate == ¤t_pdhg_solver.get_primal_solution()); if (is_current_live_iterate) { engine->gather_potential_next_solutions_to_master(); } } +#endif typename convergence_information_t::view_t convergence_information_view = convergence_information_.view(); diff --git a/cpp/tests/linear_programming/CMakeLists.txt b/cpp/tests/linear_programming/CMakeLists.txt index 08fd9598c0..98e88c69ce 100644 --- a/cpp/tests/linear_programming/CMakeLists.txt +++ b/cpp/tests/linear_programming/CMakeLists.txt @@ -17,10 +17,12 @@ ConfigureTest(LP_UNIT_TEST # ci/test_cpp_multi_gpu.sh picks it up via its *_MG_TEST glob. Each test # calls GTEST_SKIP when fewer than 2 GPUs are visible so the binary is # harmless on single-GPU runners. STATIC_LIB: needs internal (hidden) PDLP symbols. -ConfigureTest(PDLP_MG_TEST - ${CMAKE_CURRENT_SOURCE_DIR}/pdlp_distributed_test.cu - LABELS numopt - STATIC_LIB) +if(NOT SKIP_DISTRIBUTED_PDLP_BUILD) + ConfigureTest(PDLP_MG_TEST + ${CMAKE_CURRENT_SOURCE_DIR}/pdlp_distributed_test.cu + LABELS numopt + STATIC_LIB) +endif() # ################################################################################################## # - MPS / LP parser tests --------------------------------------------------------------------------