diff --git a/cpp/include/cuopt/mathematical_optimization/constants.h b/cpp/include/cuopt/mathematical_optimization/constants.h index 3656791a98..66c8aa3755 100644 --- a/cpp/include/cuopt/mathematical_optimization/constants.h +++ b/cpp/include/cuopt/mathematical_optimization/constants.h @@ -36,6 +36,7 @@ #define CUOPT_NODE_LIMIT "node_limit" #define CUOPT_PDLP_SOLVER_MODE "pdlp_solver_mode" #define CUOPT_METHOD "method" +#define CUOPT_CONCURRENT_NNZ_CUTOFF "concurrent_nnz_cutoff" #define CUOPT_PER_CONSTRAINT_RESIDUAL "per_constraint_residual" #define CUOPT_SAVE_BEST_PRIMAL_SO_FAR "save_best_primal_so_far" #define CUOPT_FIRST_PRIMAL_FEASIBLE "first_primal_feasible" diff --git a/cpp/include/cuopt/mathematical_optimization/mip/solver_settings.hpp b/cpp/include/cuopt/mathematical_optimization/mip/solver_settings.hpp index 7ed45f1f9b..f3dc5fe340 100644 --- a/cpp/include/cuopt/mathematical_optimization/mip/solver_settings.hpp +++ b/cpp/include/cuopt/mathematical_optimization/mip/solver_settings.hpp @@ -147,6 +147,7 @@ class mip_solver_settings_t { i_t strong_branching_simplex_iteration_limit = -1; i_t num_gpus = 1; method_t method{method_t::Concurrent}; + i_t concurrent_nnz_cutoff{50'000'000}; bool log_to_console = true; std::string log_file; diff --git a/cpp/include/cuopt/mathematical_optimization/pdlp/solver_settings.hpp b/cpp/include/cuopt/mathematical_optimization/pdlp/solver_settings.hpp index fd6c497850..da772aee98 100644 --- a/cpp/include/cuopt/mathematical_optimization/pdlp/solver_settings.hpp +++ b/cpp/include/cuopt/mathematical_optimization/pdlp/solver_settings.hpp @@ -353,6 +353,9 @@ class pdlp_solver_settings_t { // distributed_pdlp_partitioner_t for the meaning of each value. distributed_pdlp_partitioner_t distributed_pdlp_partitioner{distributed_pdlp_partitioner_t::Auto}; method_t method{method_t::Concurrent}; + // TODO: Remove this cutoff once concurrent CPU solver memory usage and cuDSS long running kernels + // are resolved. -1 disables the cutoff regardless of the reduced problem's NNZ. + i_t concurrent_nnz_cutoff{50'000'000}; bool inside_mip{false}; // For concurrent termination std::atomic* concurrent_halt{nullptr}; diff --git a/cpp/src/math_optimization/solver_settings.cu b/cpp/src/math_optimization/solver_settings.cu index 5a4ab72c32..5d655c94c5 100644 --- a/cpp/src/math_optimization/solver_settings.cu +++ b/cpp/src/math_optimization/solver_settings.cu @@ -174,6 +174,8 @@ solver_settings_t::solver_settings_t() : pdlp_settings(), mip_settings {CUOPT_PDLP_SOLVER_MODE, reinterpret_cast(&pdlp_settings.pdlp_solver_mode), CUOPT_PDLP_SOLVER_MODE_STABLE1, CUOPT_PDLP_SOLVER_MODE_STABLE3, CUOPT_PDLP_SOLVER_MODE_STABLE3}, {CUOPT_METHOD, reinterpret_cast(&pdlp_settings.method), CUOPT_METHOD_CONCURRENT, CUOPT_METHOD_BARRIER, CUOPT_METHOD_CONCURRENT}, {CUOPT_METHOD, reinterpret_cast(&mip_settings.method), CUOPT_METHOD_CONCURRENT, CUOPT_METHOD_BARRIER, CUOPT_METHOD_CONCURRENT}, + {CUOPT_CONCURRENT_NNZ_CUTOFF, &pdlp_settings.concurrent_nnz_cutoff, -1, std::numeric_limits::max(), 50'000'000, "skip Barrier and dual simplex in concurrent solves at this reduced NNZ; -1 disables the cutoff"}, + {CUOPT_CONCURRENT_NNZ_CUTOFF, &mip_settings.concurrent_nnz_cutoff, -1, std::numeric_limits::max(), 50'000'000, "skip Barrier and dual simplex in concurrent solves at this reduced NNZ; -1 disables the cutoff"}, {CUOPT_NUM_CPU_THREADS, &mip_settings.num_cpu_threads, -1, std::numeric_limits::max(), -1}, {CUOPT_AUGMENTED, &pdlp_settings.augmented, -1, 1, -1}, {CUOPT_FOLDING, &pdlp_settings.folding, -1, 1, -1}, diff --git a/cpp/src/mip_heuristics/diversity/diversity_manager.cu b/cpp/src/mip_heuristics/diversity/diversity_manager.cu index 96165e6064..55fb1dd29b 100644 --- a/cpp/src/mip_heuristics/diversity/diversity_manager.cu +++ b/cpp/src/mip_heuristics/diversity/diversity_manager.cu @@ -577,6 +577,7 @@ solution_t diversity_manager_t::run_solver() pdlp_settings.first_primal_feasible = false; pdlp_settings.concurrent_halt = &global_concurrent_halt; pdlp_settings.method = context.settings.method; + pdlp_settings.concurrent_nnz_cutoff = context.settings.concurrent_nnz_cutoff; pdlp_settings.inside_mip = true; pdlp_settings.pdlp_solver_mode = pdlp_solver_mode_t::Stable2; pdlp_settings.num_gpus = context.settings.num_gpus; diff --git a/cpp/src/mip_heuristics/mip_constants.hpp b/cpp/src/mip_heuristics/mip_constants.hpp index 9719377da1..11a67aef3f 100644 --- a/cpp/src/mip_heuristics/mip_constants.hpp +++ b/cpp/src/mip_heuristics/mip_constants.hpp @@ -29,10 +29,6 @@ #define CUOPT_MIP_BATCH_PDLP_REQUIRED_THREAD_COUNT 3 #define CUOPT_MIP_CLIQUE_CUTS_REQUIRED_THREAD_COUNT 3 -// MIP-only gate: skip the concurrent barrier when fewer threads are available than this -// (1 PDLP + 1 dual simplex + 1 barrier). Stand-alone LP always runs all three. -#define CUOPT_CONCURRENT_LP_BARRIER_REQUIRED_THREAD_COUNT 3 - /* @brief Priority classes for the omp tasks. Highest value = higher priority. * Note that this only gives a hint to the runtime, such that the high priority * is not guarantee to be executed before a low priority one (i.e., do not rely on diff --git a/cpp/src/mip_heuristics/solver.cu b/cpp/src/mip_heuristics/solver.cu index 166afd832b..058955f95e 100644 --- a/cpp/src/mip_heuristics/solver.cu +++ b/cpp/src/mip_heuristics/solver.cu @@ -262,10 +262,11 @@ solution_t mip_solver_t::run_solver() if (run_presolve && context.problem_ptr->n_integer_vars == 0) { CUOPT_LOG_INFO("Problem reduced to a LP, running concurrent LP"); pdlp_solver_settings_t settings{}; - settings.time_limit = timer_.remaining_time(); - auto lp_timer = timer_t(settings.time_limit); - settings.method = method_t::Concurrent; - settings.presolver = presolver_t::None; + settings.time_limit = timer_.remaining_time(); + auto lp_timer = timer_t(settings.time_limit); + settings.method = method_t::Concurrent; + settings.concurrent_nnz_cutoff = context.settings.concurrent_nnz_cutoff; + settings.presolver = presolver_t::None; auto opt_sol = solve_lp_with_method(*context.problem_ptr, settings, lp_timer); diff --git a/cpp/src/pdlp/pdlp_constants.hpp b/cpp/src/pdlp/pdlp_constants.hpp index 1651fd4212..7391c0f723 100644 --- a/cpp/src/pdlp/pdlp_constants.hpp +++ b/cpp/src/pdlp/pdlp_constants.hpp @@ -9,6 +9,8 @@ #include +#include + namespace cuopt::mathematical_optimization::pdlp { inline constexpr int block_size = 128; @@ -38,6 +40,29 @@ inline constexpr bool deterministic_batch_pdlp = true; inline constexpr bool enable_batch_resizing = true; +inline constexpr int concurrent_barrier_required_thread_count = 3; + +inline constexpr bool should_skip_concurrent_cpu_solvers(std::size_t num_nonzeros, int nnz_cutoff) +{ + return nnz_cutoff >= 0 && num_nonzeros >= static_cast(nnz_cutoff); +} + +inline constexpr bool should_enable_concurrent_barrier(std::size_t num_nonzeros, + int nnz_cutoff, + bool inside_mip, + int available_threads) +{ + return !should_skip_concurrent_cpu_solvers(num_nonzeros, nnz_cutoff) && + (!inside_mip || available_threads >= concurrent_barrier_required_thread_count); +} + +inline constexpr bool should_enable_concurrent_dual_simplex(std::size_t num_nonzeros, + int nnz_cutoff, + bool inside_mip) +{ + return !inside_mip && !should_skip_concurrent_cpu_solvers(num_nonzeros, nnz_cutoff); +} + // Value used to determine what we see as too small (the value) or too large (1/value) values when // computing the new primal weight during the restart. template diff --git a/cpp/src/pdlp/solve.cu b/cpp/src/pdlp/solve.cu index 38202c6b51..d9892ab55f 100644 --- a/cpp/src/pdlp/solve.cu +++ b/cpp/src/pdlp/solve.cu @@ -1585,11 +1585,15 @@ optimization_problem_solution_t run_concurrent( // Make sure allocations are done on the original stream problem.handle_ptr->sync_stream(); - // Stand-alone LP always runs all three concurrently. MIP gates the barrier so we don't - // overshoot num_cpu_threads (need 1 PDLP + 1 dual simplex + 1 barrier). + // Keep the concurrent solver thread count correct when CPU solvers are skipped. + const auto num_nonzeros = problem.coefficients.size(); const int available_threads = omp_in_parallel() ? omp_get_num_threads() : omp_get_max_threads(); - const bool enable_barrier = - !settings.inside_mip || available_threads >= CUOPT_CONCURRENT_LP_BARRIER_REQUIRED_THREAD_COUNT; + const bool skip_cpu_solvers = + pdlp::should_skip_concurrent_cpu_solvers(num_nonzeros, settings.concurrent_nnz_cutoff); + const bool enable_barrier = pdlp::should_enable_concurrent_barrier( + num_nonzeros, settings.concurrent_nnz_cutoff, settings.inside_mip, available_threads); + const bool enable_dual_simplex = pdlp::should_enable_concurrent_dual_simplex( + num_nonzeros, settings.concurrent_nnz_cutoff, settings.inside_mip); if (settings.num_gpus > 1) { int device_count = raft::device_setter::get_device_count(); @@ -1601,11 +1605,13 @@ optimization_problem_solution_t run_concurrent( device_count > 1, error_type_t::RuntimeError, "Multi-GPU mode requires at least 2 GPUs"); } - // Initialize the dual simplex structures before we run PDLP. - // Otherwise, CUDA API calls to the problem stream may occur in both threads and throw graph - // capture off - simplex::user_problem_t dual_simplex_problem = - cuopt_problem_to_user_problem(problem.handle_ptr, problem, false); + // Initialize the shared CPU solver structures before we run PDLP. Otherwise, CUDA API calls to + // the problem stream may occur in multiple threads and throw graph capture off. + std::unique_ptr> concurrent_cpu_problem; + if (enable_barrier || enable_dual_simplex) { + concurrent_cpu_problem = std::make_unique>( + cuopt_problem_to_user_problem(problem.handle_ptr, problem, false)); + } // Dual simplex / barrier results — written by tasks, read after the taskgroup barrier. std::unique_ptr, simplex::lp_status_t, f_t, f_t, f_t>> sol_dual_simplex_ptr; @@ -1620,10 +1626,21 @@ optimization_problem_solution_t run_concurrent( // library init is now recovered by manual_cuda_graph_t::run, so the previous main-thread // preflight (eager handle construction + cuDSS warmup) is no longer needed. std::unique_ptr barrier_handle_ptr; - if (!enable_barrier) { - CUOPT_LOG_DEBUG("MIP: skipping concurrent barrier, %d threads available < %d required.", + if (skip_cpu_solvers) { + CUOPT_LOG_CONDITIONAL_INFO( + !settings.inside_mip, + "Skipping concurrent barrier and dual simplex: reduced problem has %zu nonzeros " + "(CONCURRENT_NNZ_CUTOFF: %d).", + num_nonzeros, + settings.concurrent_nnz_cutoff); + CUOPT_LOG_DEBUG( + "Skipping concurrent CPU solvers: reduced problem has %zu nonzeros (cutoff: %d).", + num_nonzeros, + settings.concurrent_nnz_cutoff); + } else if (!enable_barrier) { + CUOPT_LOG_DEBUG("MIP: skipping concurrent Barrier, %d threads available < %d required.", available_threads, - CUOPT_CONCURRENT_LP_BARRIER_REQUIRED_THREAD_COUNT); + pdlp::concurrent_barrier_required_thread_count); } // Dispatch barrier + dual simplex as OMP tasks (not std::threads) so they consume slots from @@ -1639,7 +1656,7 @@ optimization_problem_solution_t run_concurrent( auto dispatch_concurrent_solvers = [&]() { #pragma omp taskgroup { - // Barrier task — always on for stand-alone LP, gated on enable_barrier for MIP. + // Barrier task — gated by the reduced-problem size and, for MIP, available CPU threads. if (enable_barrier) { #pragma omp task default(shared) { @@ -1647,7 +1664,7 @@ optimization_problem_solution_t run_concurrent( auto call_barrier_thread = [&]() { cuda::stream_ref barrier_stream = cuda::stream_ref{cudaStreamPerThread}; barrier_handle_ptr = std::make_unique(barrier_stream); - run_barrier_thread(dual_simplex_problem, + run_barrier_thread(*concurrent_cpu_problem, settings_pdlp, sol_barrier_ptr, timer, @@ -1673,13 +1690,13 @@ optimization_problem_solution_t run_concurrent( } } - // Dual simplex task — skipped from MIP (B&B already drives it separately). - if (!settings.inside_mip) { + // Dual simplex task — skipped for large LPs and from MIP (B&B drives it separately). + if (enable_dual_simplex) { #pragma omp task default(shared) { try { run_dual_simplex_thread( - dual_simplex_problem, settings_pdlp, sol_dual_simplex_ptr, timer); + *concurrent_cpu_problem, settings_pdlp, sol_dual_simplex_ptr, timer); } catch (const std::exception& e) { CUOPT_LOG_ERROR("Exception in concurrent dual simplex LP: %s", e.what()); dual_simplex_exception = std::current_exception(); @@ -1717,7 +1734,7 @@ optimization_problem_solution_t run_concurrent( dispatch_concurrent_solvers(); } else { // Stand-alone LP: stand up a local team sized for 1 dispatcher + 1 per spawned task. - const int num_workers = 1 + (settings.inside_mip ? 0 : 1) + (enable_barrier ? 1 : 0); + const int num_workers = 1 + (enable_dual_simplex ? 1 : 0) + (enable_barrier ? 1 : 0); #pragma omp parallel num_threads(num_workers) default(shared) { #pragma omp single @@ -1735,13 +1752,13 @@ optimization_problem_solution_t run_concurrent( if (dual_simplex_exception) { std::rethrow_exception(dual_simplex_exception); } if (barrier_exception) { std::rethrow_exception(barrier_exception); } - // Both CPU solvers have joined, so release their shared host model before converting outputs. - dual_simplex_problem = simplex::user_problem_t(problem.handle_ptr); + // The CPU solver tasks have joined, so release their shared host model before converting outputs. + concurrent_cpu_problem.reset(); f_t end_time = timer.elapsed_time(); CUOPT_LOG_CONDITIONAL_INFO(!settings.inside_mip, "Concurrent time: %.3fs", end_time); - const auto dual_simplex_status = (!settings.inside_mip && sol_dual_simplex_ptr != nullptr) + const auto dual_simplex_status = (enable_dual_simplex && sol_dual_simplex_ptr != nullptr) ? std::get<1>(*sol_dual_simplex_ptr) : simplex::lp_status_t::CONCURRENT_LIMIT; const auto barrier_status = (enable_barrier && sol_barrier_ptr != nullptr) diff --git a/cpp/tests/linear_programming/pdlp_test.cu b/cpp/tests/linear_programming/pdlp_test.cu index 715d43b7a4..61b56fb62c 100644 --- a/cpp/tests/linear_programming/pdlp_test.cu +++ b/cpp/tests/linear_programming/pdlp_test.cu @@ -177,6 +177,54 @@ TEST(pdlp_class, concurrent_pdlp_exception_joins_worker_threads) testing::HasSubstr("all_primal_feasible only applies in batch mode")); } +TEST(pdlp_class, concurrent_cpu_solvers_gate_on_reduced_problem_size_and_mip_threads) +{ + using cuopt::mathematical_optimization::pdlp::concurrent_barrier_required_thread_count; + using cuopt::mathematical_optimization::pdlp::should_enable_concurrent_barrier; + using cuopt::mathematical_optimization::pdlp::should_enable_concurrent_dual_simplex; + using cuopt::mathematical_optimization::pdlp::should_skip_concurrent_cpu_solvers; + + constexpr int nnz_cutoff = 50'000'000; + EXPECT_FALSE(should_skip_concurrent_cpu_solvers(nnz_cutoff - 1, nnz_cutoff)); + EXPECT_TRUE(should_skip_concurrent_cpu_solvers(nnz_cutoff, nnz_cutoff)); + EXPECT_FALSE(should_skip_concurrent_cpu_solvers(nnz_cutoff, -1)); + + EXPECT_TRUE(should_enable_concurrent_barrier(nnz_cutoff - 1, nnz_cutoff, false, 1)); + EXPECT_FALSE(should_enable_concurrent_barrier(nnz_cutoff, nnz_cutoff, false, 32)); + EXPECT_TRUE(should_enable_concurrent_barrier(nnz_cutoff, -1, false, 32)); + EXPECT_TRUE(should_enable_concurrent_barrier( + nnz_cutoff - 1, nnz_cutoff, true, concurrent_barrier_required_thread_count)); + EXPECT_FALSE(should_enable_concurrent_barrier( + nnz_cutoff - 1, nnz_cutoff, true, concurrent_barrier_required_thread_count - 1)); + + EXPECT_TRUE(should_enable_concurrent_dual_simplex(nnz_cutoff - 1, nnz_cutoff, false)); + EXPECT_FALSE(should_enable_concurrent_dual_simplex(nnz_cutoff, nnz_cutoff, false)); + EXPECT_TRUE(should_enable_concurrent_dual_simplex(nnz_cutoff, -1, false)); + EXPECT_FALSE(should_enable_concurrent_dual_simplex(nnz_cutoff - 1, nnz_cutoff, true)); +} + +TEST(pdlp_class, concurrent_cutoff_runs_pdlp_without_cpu_solvers) +{ + const raft::handle_t handle_{}; + + auto path = make_path_absolute("linear_programming/afiro_original.mps"); + cuopt::mathematical_optimization::io::mps_data_model_t op_problem = + cuopt::mathematical_optimization::io::read_mps(path, true); + + auto settings = pdlp_solver_settings_t{}; + settings.method = cuopt::mathematical_optimization::method_t::Concurrent; + settings.presolver = cuopt::mathematical_optimization::presolver_t::None; + settings.concurrent_nnz_cutoff = 0; + + testing::internal::CaptureStdout(); + optimization_problem_solution_t solution = solve_lp(&handle_, op_problem, settings); + const auto logs = testing::internal::GetCapturedStdout(); + + EXPECT_THAT(logs, testing::HasSubstr("Skipping concurrent barrier and dual simplex")); + EXPECT_THAT(logs, testing::HasSubstr("(CONCURRENT_NNZ_CUTOFF: 0)")); + EXPECT_EQ((int)solution.get_termination_status(), CUOPT_TERMINATION_STATUS_OPTIMAL); +} + TEST(pdlp_class, concurrent_null_solver_ptrs_inside_mip) { const raft::handle_t handle_{}; @@ -191,7 +239,7 @@ TEST(pdlp_class, concurrent_null_solver_ptrs_inside_mip) settings.inside_mip = true; // inside_mip skips dual simplex. Setting threads to 1 ensures barrier is also disabled - // (< CUOPT_CONCURRENT_LP_BARRIER_REQUIRED_THREAD_COUNT), leaving both sol_dual_simplex_ptr + // (< concurrent_barrier_required_thread_count), leaving both sol_dual_simplex_ptr // and sol_barrier_ptr null. const int prev_threads = omp_get_max_threads(); omp_set_num_threads(1); diff --git a/cpp/tests/linear_programming/unit_tests/solver_settings_test.cu b/cpp/tests/linear_programming/unit_tests/solver_settings_test.cu index 2374a8f2c7..c3f5e9bf67 100644 --- a/cpp/tests/linear_programming/unit_tests/solver_settings_test.cu +++ b/cpp/tests/linear_programming/unit_tests/solver_settings_test.cu @@ -91,6 +91,17 @@ TEST(SolverSettingsTest, SequenceSolveParameter) EXPECT_TRUE(settings.get_pdlp_settings().sequence_solve); } +TEST(SolverSettingsTest, ConcurrentNnzCutoffParameter) +{ + solver_settings_t settings; + EXPECT_EQ(settings.get_pdlp_settings().concurrent_nnz_cutoff, 50'000'000); + EXPECT_EQ(settings.get_mip_settings().concurrent_nnz_cutoff, 50'000'000); + + settings.set_parameter_from_string(CUOPT_CONCURRENT_NNZ_CUTOFF, "-1"); + EXPECT_EQ(settings.get_pdlp_settings().concurrent_nnz_cutoff, -1); + EXPECT_EQ(settings.get_mip_settings().concurrent_nnz_cutoff, -1); +} + TEST(SolverSettingsTest, warm_start_smaller_vector) { const raft::handle_t handle_{};