From d8884aa5cbad4ee76b21c9a139a0c5f83594917b Mon Sep 17 00:00:00 2001 From: Rajesh Gandham Date: Fri, 21 Aug 2026 16:03:53 -0700 Subject: [PATCH 1/4] First stab at refactoring --- cpp/src/branch_and_bound/CMakeLists.txt | 1 + cpp/src/branch_and_bound/branch_and_bound.cpp | 44 ++- cpp/src/branch_and_bound/branch_and_bound.hpp | 24 ++ .../diversity/diversity_manager.cu | 274 ++++++++++-------- cpp/src/mip_heuristics/solver.cu | 34 ++- 5 files changed, 237 insertions(+), 140 deletions(-) diff --git a/cpp/src/branch_and_bound/CMakeLists.txt b/cpp/src/branch_and_bound/CMakeLists.txt index 1e40c1bbf1..a3c8fa064d 100644 --- a/cpp/src/branch_and_bound/CMakeLists.txt +++ b/cpp/src/branch_and_bound/CMakeLists.txt @@ -5,6 +5,7 @@ set(BRANCH_AND_BOUND_SRC_FILES ${CMAKE_CURRENT_SOURCE_DIR}/branch_and_bound.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/concurrent_root_solver.cu ${CMAKE_CURRENT_SOURCE_DIR}/pseudo_costs.cpp ${CMAKE_CURRENT_SOURCE_DIR}/diving_heuristics.cpp ) diff --git a/cpp/src/branch_and_bound/branch_and_bound.cpp b/cpp/src/branch_and_bound/branch_and_bound.cpp index fd7987ae1c..931df3cdfb 100644 --- a/cpp/src/branch_and_bound/branch_and_bound.cpp +++ b/cpp/src/branch_and_bound/branch_and_bound.cpp @@ -6,6 +6,7 @@ /* clang-format on */ #include +#include #include #include #include @@ -2808,9 +2809,50 @@ lp_status_t branch_and_bound_t::solve_root_relaxation( root_vstatus_, edge_norms_, nullptr); + // Dual simplex has finished; stop the GPU competitors if they are still running. + gpu_root_concurrent_halt_.store(1, std::memory_order_release); } - // Wait for the root relaxation solution to be sent by the diversity manager or dual simplex + // The diversity manager prepares the GPU problem while dual simplex starts on the CPU. + // Once the GPU problem is ready, launch PDLP and barrier from here so all root-LP + // competitors are owned by this function. + while (!concurrent_root_problem_ready_.load(std::memory_order_acquire) && + *get_root_concurrent_halt() == 0) { + std::this_thread::sleep_for(std::chrono::milliseconds(1)); +#pragma omp taskyield + } + + if (*get_root_concurrent_halt() == 0 && + concurrent_root_problem_ready_.load(std::memory_order_acquire)) { + cuopt_assert(concurrent_root_problem_ != nullptr, "Concurrent root problem is not configured"); + gpu_root_concurrent_halt_.store(0, std::memory_order_release); + try { + cuopt_assert(concurrent_root_settings_ != nullptr, + "Concurrent root settings are not configured"); + const f_t remaining_time = + std::max(settings_.time_limit - toc(exploration_stats_.start_time), 0); + const f_t root_time_limit = + std::min(concurrent_root_max_time_, remaining_time * concurrent_root_time_ratio_); + auto result = solve_concurrent_root_relaxation(concurrent_root_problem_, + *concurrent_root_settings_, + root_time_limit, + &gpu_root_concurrent_halt_); + if (result.usable) { + set_root_relaxation_solution(result.primal, + result.dual, + result.reduced_cost, + result.solver_objective, + result.user_objective, + result.iterations, + result.method); + } + } catch (const std::exception& e) { + settings_.log.printf("Concurrent GPU root LP failed: %s\n", e.what()); + } + } + + // Wait until either the GPU root solve supplies a crossover point or CPU dual + // simplex finishes. If dual simplex wins, stop and join the GPU solve. while (!root_crossover_solution_set_.load(std::memory_order_acquire) && *get_root_concurrent_halt() == 0) { std::this_thread::sleep_for(std::chrono::milliseconds(1)); diff --git a/cpp/src/branch_and_bound/branch_and_bound.hpp b/cpp/src/branch_and_bound/branch_and_bound.hpp index 96b8a6d8fe..6c3233f380 100644 --- a/cpp/src/branch_and_bound/branch_and_bound.hpp +++ b/cpp/src/branch_and_bound/branch_and_bound.hpp @@ -81,6 +81,9 @@ struct clique_table_t; template struct mip_symmetry_t; +template +class problem_t; + template struct nondeterministic_policy_t; template @@ -145,6 +148,21 @@ class branch_and_bound_t { } void set_concurrent_lp_root_solve(bool enable) { enable_concurrent_lp_root_solve_ = enable; } + void configure_concurrent_lp_root_solve(problem_t* problem, + const pdlp_solver_settings_t& settings, + f_t max_time, + f_t time_ratio) + { + concurrent_root_problem_ = problem; + concurrent_root_settings_ = std::make_unique>(settings); + concurrent_root_max_time_ = max_time; + concurrent_root_time_ratio_ = time_ratio; + enable_concurrent_lp_root_solve_ = true; + } + void notify_concurrent_root_problem_ready() + { + concurrent_root_problem_ready_.store(true, std::memory_order_release); + } // Seed the global upper bound from an external source (e.g., early FJ during presolve). // `bound` must be in B&B's internal objective space. @@ -255,6 +273,12 @@ class branch_and_bound_t { omp_atomic_t root_lp_current_lower_bound_; omp_atomic_t solving_root_relaxation_{false}; bool enable_concurrent_lp_root_solve_{false}; + problem_t* concurrent_root_problem_{nullptr}; + std::unique_ptr> concurrent_root_settings_; + f_t concurrent_root_max_time_{0}; + f_t concurrent_root_time_ratio_{0}; + std::atomic concurrent_root_problem_ready_{false}; + std::atomic gpu_root_concurrent_halt_{0}; std::atomic root_concurrent_halt_{0}; std::atomic node_concurrent_halt_{0}; bool is_root_solution_set{false}; diff --git a/cpp/src/mip_heuristics/diversity/diversity_manager.cu b/cpp/src/mip_heuristics/diversity/diversity_manager.cu index ec82c4b423..047312a557 100644 --- a/cpp/src/mip_heuristics/diversity/diversity_manager.cu +++ b/cpp/src/mip_heuristics/diversity/diversity_manager.cu @@ -517,10 +517,7 @@ solution_t diversity_manager_t::run_solver() return population.best_feasible(); } - population.timer = timer; - const f_t time_limit = timer.remaining_time(); - const auto& hp = context.settings.heuristic_params; - const f_t lp_time_limit = std::min(hp.root_lp_max_time, time_limit * hp.root_lp_time_ratio); + population.timer = timer; // after every change to the problem, we should resize all the relevant vars // we need to encapsulate that to prevent repetitions recombine_stats.reset(); @@ -558,135 +555,158 @@ solution_t diversity_manager_t::run_solver() } else if (!fj_only_run) { convert_greater_to_less(*problem_ptr); - f_t absolute_tolerance = context.settings.tolerances.absolute_tolerance; - - pdlp_solver_settings_t pdlp_settings{}; - pdlp_settings.tolerances.absolute_dual_tolerance = absolute_tolerance; - pdlp_settings.tolerances.relative_dual_tolerance = - context.settings.tolerances.relative_tolerance; - pdlp_settings.tolerances.absolute_primal_tolerance = absolute_tolerance; - pdlp_settings.tolerances.relative_primal_tolerance = - context.settings.tolerances.relative_tolerance; - pdlp_settings.time_limit = lp_time_limit; - pdlp_settings.first_primal_feasible = false; - pdlp_settings.concurrent_halt = &global_concurrent_halt; - pdlp_settings.method = method_t::Concurrent; - pdlp_settings.inside_mip = true; - pdlp_settings.pdlp_solver_mode = pdlp_solver_mode_t::Stable2; - pdlp_settings.num_gpus = context.settings.num_gpus; - pdlp_settings.presolver = presolver_t::None; - pdlp_settings.per_constraint_residual = true; - set_pdlp_solver_mode(pdlp_settings); - timer_t lp_timer(lp_time_limit); - auto lp_result = solve_lp_with_method(*problem_ptr, pdlp_settings, lp_timer); - - // The concurrent root LP can fail to produce a usable solution -- e.g. the barrier - // hits a numerical error on an infeasible problem and PDLP returns NumericalError - // with empty primal/dual. In that case we must not copy or hand off the empty - // result (copying n elements from an empty buffer throws), and we must still - // release B&B's root-relaxation wait so it proceeds with its own dual-simplex root - // instead of spinning forever. - const bool root_lp_usable = - lp_result.get_termination_status() != pdlp_termination_status_t::NumericalError && - lp_result.get_primal_solution().size() == lp_optimal_solution.size() && - lp_result.get_dual_solution().size() == lp_dual_optimal_solution.size(); - - bool use_staged_simplex_solution = false; - { - std::lock_guard guard(relaxed_solution_mutex); - use_staged_simplex_solution = simplex_solution_exists.load(); - if (!use_staged_simplex_solution && root_lp_usable) { - raft::copy(lp_optimal_solution.data(), - lp_result.get_primal_solution().data(), - lp_optimal_solution.size(), - problem_ptr->handle_ptr->get_stream()); - raft::copy(lp_dual_optimal_solution.data(), - lp_result.get_dual_solution().data(), - lp_dual_optimal_solution.size(), - problem_ptr->handle_ptr->get_stream()); + if (context.branch_and_bound_ptr != nullptr && + context.branch_and_bound_ptr->enable_concurrent_lp_root_solve()) { + // B&B owns the CPU/GPU concurrent root solve. Signal that the GPU problem + // is ready, then wait for B&B to publish the winning root solution through + // set_simplex_solution_callback. + context.branch_and_bound_ptr->notify_concurrent_root_problem_ready(); + while (!simplex_solution_exists.load(std::memory_order_acquire) && !check_b_b_preemption()) { + std::this_thread::sleep_for(std::chrono::milliseconds(1)); } - } - if (use_staged_simplex_solution) { consume_staged_simplex_solution(lp_state); } - if (use_staged_simplex_solution || root_lp_usable) { - cuopt_assert(thrust::all_of(problem_ptr->handle_ptr->get_thrust_policy(), - lp_optimal_solution.begin(), - lp_optimal_solution.end(), - [] __host__ __device__(f_t val) { return std::isfinite(val); }), - "LP optimal solution contains non-finite values"); - } - ls.lp_optimal_exists = true; - if (!use_staged_simplex_solution) { - if (!root_lp_usable) { - // The concurrent root LP produced no usable solution. Do not hand an empty - // solution to B&B; instead release its root-relaxation wait loop so it falls - // back to its own dual-simplex root rather than deadlocking. - CUOPT_LOG_DEBUG("Root LP produced no usable solution (status %d); releasing B&B root solve", - (int)lp_result.get_termination_status()); + if (simplex_solution_exists.load(std::memory_order_acquire)) { + consume_staged_simplex_solution(lp_state); + ls.lp_optimal_exists = true; + } else { ls.lp_optimal_exists = false; - if (context.branch_and_bound_ptr != nullptr) { - context.branch_and_bound_ptr->set_root_concurrent_halt(1); + } + } else { + // Heuristics-only fallback: no B&B object exists to own the root solve. + const f_t time_limit = timer.remaining_time(); + const auto& hp = context.settings.heuristic_params; + const f_t lp_time_limit = std::min(hp.root_lp_max_time, time_limit * hp.root_lp_time_ratio); + f_t absolute_tolerance = context.settings.tolerances.absolute_tolerance; + + pdlp_solver_settings_t pdlp_settings{}; + pdlp_settings.tolerances.absolute_dual_tolerance = absolute_tolerance; + pdlp_settings.tolerances.relative_dual_tolerance = + context.settings.tolerances.relative_tolerance; + pdlp_settings.tolerances.absolute_primal_tolerance = absolute_tolerance; + pdlp_settings.tolerances.relative_primal_tolerance = + context.settings.tolerances.relative_tolerance; + pdlp_settings.time_limit = lp_time_limit; + pdlp_settings.first_primal_feasible = false; + pdlp_settings.concurrent_halt = &global_concurrent_halt; + pdlp_settings.method = method_t::Concurrent; + pdlp_settings.inside_mip = true; + pdlp_settings.pdlp_solver_mode = pdlp_solver_mode_t::Stable2; + pdlp_settings.num_gpus = context.settings.num_gpus; + pdlp_settings.presolver = presolver_t::None; + pdlp_settings.per_constraint_residual = true; + set_pdlp_solver_mode(pdlp_settings); + timer_t lp_timer(lp_time_limit); + auto lp_result = solve_lp_with_method(*problem_ptr, pdlp_settings, lp_timer); + + // The concurrent root LP can fail to produce a usable solution -- e.g. the barrier + // hits a numerical error on an infeasible problem and PDLP returns NumericalError + // with empty primal/dual. In that case we must not copy or hand off the empty + // result (copying n elements from an empty buffer throws), and we must still + // release B&B's root-relaxation wait so it proceeds with its own dual-simplex root + // instead of spinning forever. + const bool root_lp_usable = + lp_result.get_termination_status() != pdlp_termination_status_t::NumericalError && + lp_result.get_primal_solution().size() == lp_optimal_solution.size() && + lp_result.get_dual_solution().size() == lp_dual_optimal_solution.size(); + + bool use_staged_simplex_solution = false; + { + std::lock_guard guard(relaxed_solution_mutex); + use_staged_simplex_solution = simplex_solution_exists.load(); + if (!use_staged_simplex_solution && root_lp_usable) { + raft::copy(lp_optimal_solution.data(), + lp_result.get_primal_solution().data(), + lp_optimal_solution.size(), + problem_ptr->handle_ptr->get_stream()); + raft::copy(lp_dual_optimal_solution.data(), + lp_result.get_dual_solution().data(), + lp_dual_optimal_solution.size(), + problem_ptr->handle_ptr->get_stream()); + } + } + if (use_staged_simplex_solution) { consume_staged_simplex_solution(lp_state); } + if (use_staged_simplex_solution || root_lp_usable) { + cuopt_assert(thrust::all_of(problem_ptr->handle_ptr->get_thrust_policy(), + lp_optimal_solution.begin(), + lp_optimal_solution.end(), + [] __host__ __device__(f_t val) { return std::isfinite(val); }), + "LP optimal solution contains non-finite values"); + } + ls.lp_optimal_exists = true; + if (!use_staged_simplex_solution) { + if (!root_lp_usable) { + // The concurrent root LP produced no usable solution. Do not hand an empty + // solution to B&B; instead release its root-relaxation wait loop so it falls + // back to its own dual-simplex root rather than deadlocking. + CUOPT_LOG_DEBUG( + "Root LP produced no usable solution (status %d); releasing B&B root solve", + (int)lp_result.get_termination_status()); + ls.lp_optimal_exists = false; + if (context.branch_and_bound_ptr != nullptr) { + context.branch_and_bound_ptr->set_root_concurrent_halt(1); + } + } else if (lp_result.get_termination_status() == pdlp_termination_status_t::Optimal) { + solution_t lp_sol(*problem_ptr); + lp_sol.copy_new_assignment(lp_optimal_solution); + const bool consider_integrality = false; + lp_sol.compute_feasibility(consider_integrality); + if (lp_sol.get_feasible()) { set_new_user_bound(lp_result.get_objective_value()); } + } else if (lp_result.get_termination_status() == + pdlp_termination_status_t::PrimalInfeasible) { + CUOPT_LOG_ERROR("Problem is primal infeasible, continuing anyway!"); + ls.lp_optimal_exists = false; + } else if (lp_result.get_termination_status() == + pdlp_termination_status_t::DualInfeasible) { + CUOPT_LOG_ERROR("PDLP detected dual infeasibility, continuing anyway!"); + ls.lp_optimal_exists = false; + } else if (lp_result.get_termination_status() == pdlp_termination_status_t::TimeLimit) { + CUOPT_LOG_DEBUG( + "Initial LP run exceeded time limit, continuing solver with partial LP result!"); + // note to developer, in debug mode the LP run might be too slow and it might cause PDLP + // not to bring variables within the bounds } - } else if (lp_result.get_termination_status() == pdlp_termination_status_t::Optimal) { - solution_t lp_sol(*problem_ptr); - lp_sol.copy_new_assignment(lp_optimal_solution); - const bool consider_integrality = false; - lp_sol.compute_feasibility(consider_integrality); - if (lp_sol.get_feasible()) { set_new_user_bound(lp_result.get_objective_value()); } - } else if (lp_result.get_termination_status() == - pdlp_termination_status_t::PrimalInfeasible) { - CUOPT_LOG_ERROR("Problem is primal infeasible, continuing anyway!"); - ls.lp_optimal_exists = false; - } else if (lp_result.get_termination_status() == pdlp_termination_status_t::DualInfeasible) { - CUOPT_LOG_ERROR("PDLP detected dual infeasibility, continuing anyway!"); - ls.lp_optimal_exists = false; - } else if (lp_result.get_termination_status() == pdlp_termination_status_t::TimeLimit) { - CUOPT_LOG_DEBUG( - "Initial LP run exceeded time limit, continuing solver with partial LP result!"); - // note to developer, in debug mode the LP run might be too slow and it might cause PDLP - // not to bring variables within the bounds } - } - // Hand the root relaxation off to branch and bound when we have a usable solution - // (sets root_crossover_solution_set_, releasing B&B's wait). When the root LP failed - // the wait is instead released above via set_root_concurrent_halt, and a staged - // dual-simplex solution is owned by B&B already, so neither needs this hand-off. - if (!use_staged_simplex_solution && root_lp_usable && - problem_ptr->set_root_relaxation_solution_callback != nullptr) { - auto& d_primal_solution = lp_result.get_primal_solution(); - auto& d_dual_solution = lp_result.get_dual_solution(); - auto& d_reduced_costs = lp_result.get_reduced_cost(); - - std::vector host_primal(d_primal_solution.size()); - std::vector host_dual(d_dual_solution.size()); - std::vector host_reduced_costs(d_reduced_costs.size()); - raft::copy(host_primal.data(), - d_primal_solution.data(), - d_primal_solution.size(), - problem_ptr->handle_ptr->get_stream()); - raft::copy(host_dual.data(), - d_dual_solution.data(), - d_dual_solution.size(), - problem_ptr->handle_ptr->get_stream()); - raft::copy(host_reduced_costs.data(), - d_reduced_costs.data(), - d_reduced_costs.size(), - problem_ptr->handle_ptr->get_stream()); - problem_ptr->handle_ptr->sync_stream(); - - // PDLP returns user-space objective (it applies objective_scaling_factor internally) - auto user_obj = lp_result.get_objective_value(); - auto solver_obj = problem_ptr->get_solver_obj_from_user_obj(user_obj); - auto iterations = lp_result.get_additional_termination_information().number_of_steps_taken; - auto method = lp_result.get_additional_termination_information().solved_by; - // Set for the B&B (param4 expects solver space, param5 expects user space) - problem_ptr->set_root_relaxation_solution_callback( - host_primal, host_dual, host_reduced_costs, solver_obj, user_obj, iterations, method); - } + // Hand the root relaxation off to branch and bound when we have a usable solution + // (sets root_crossover_solution_set_, releasing B&B's wait). When the root LP failed + // the wait is instead released above via set_root_concurrent_halt, and a staged + // dual-simplex solution is owned by B&B already, so neither needs this hand-off. + if (!use_staged_simplex_solution && root_lp_usable && + problem_ptr->set_root_relaxation_solution_callback != nullptr) { + auto& d_primal_solution = lp_result.get_primal_solution(); + auto& d_dual_solution = lp_result.get_dual_solution(); + auto& d_reduced_costs = lp_result.get_reduced_cost(); + + std::vector host_primal(d_primal_solution.size()); + std::vector host_dual(d_dual_solution.size()); + std::vector host_reduced_costs(d_reduced_costs.size()); + raft::copy(host_primal.data(), + d_primal_solution.data(), + d_primal_solution.size(), + problem_ptr->handle_ptr->get_stream()); + raft::copy(host_dual.data(), + d_dual_solution.data(), + d_dual_solution.size(), + problem_ptr->handle_ptr->get_stream()); + raft::copy(host_reduced_costs.data(), + d_reduced_costs.data(), + d_reduced_costs.size(), + problem_ptr->handle_ptr->get_stream()); + problem_ptr->handle_ptr->sync_stream(); + + // PDLP returns user-space objective (it applies objective_scaling_factor internally) + auto user_obj = lp_result.get_objective_value(); + auto solver_obj = problem_ptr->get_solver_obj_from_user_obj(user_obj); + auto iterations = lp_result.get_additional_termination_information().number_of_steps_taken; + auto method = lp_result.get_additional_termination_information().solved_by; + // Set for the B&B (param4 expects solver space, param5 expects user space) + problem_ptr->set_root_relaxation_solution_callback( + host_primal, host_dual, host_reduced_costs, solver_obj, user_obj, iterations, method); + } - if (!use_staged_simplex_solution && root_lp_usable) { - // in case the pdlp returned var boudns that are out of bounds - clamp_within_var_bounds(lp_optimal_solution, problem_ptr, problem_ptr->handle_ptr); + if (!use_staged_simplex_solution && root_lp_usable) { + // in case the pdlp returned var boudns that are out of bounds + clamp_within_var_bounds(lp_optimal_solution, problem_ptr, problem_ptr->handle_ptr); + } } } diff --git a/cpp/src/mip_heuristics/solver.cu b/cpp/src/mip_heuristics/solver.cu index f8eac0c4d8..a4f0518465 100644 --- a/cpp/src/mip_heuristics/solver.cu +++ b/cpp/src/mip_heuristics/solver.cu @@ -442,7 +442,28 @@ solution_t mip_solver_t::run_solver() // Set the primal heuristics -> branch and bound callback if (context.settings.determinism_mode == CUOPT_MODE_OPPORTUNISTIC) { - branch_and_bound->set_concurrent_lp_root_solve(true); + pdlp_solver_settings_t concurrent_root_settings{}; + concurrent_root_settings.tolerances.absolute_dual_tolerance = + context.settings.tolerances.absolute_tolerance; + concurrent_root_settings.tolerances.relative_dual_tolerance = + context.settings.tolerances.relative_tolerance; + concurrent_root_settings.tolerances.absolute_primal_tolerance = + context.settings.tolerances.absolute_tolerance; + concurrent_root_settings.tolerances.relative_primal_tolerance = + context.settings.tolerances.relative_tolerance; + concurrent_root_settings.first_primal_feasible = false; + concurrent_root_settings.method = method_t::Concurrent; + concurrent_root_settings.inside_mip = true; + concurrent_root_settings.pdlp_solver_mode = pdlp_solver_mode_t::Stable2; + concurrent_root_settings.num_gpus = context.settings.num_gpus; + concurrent_root_settings.presolver = presolver_t::None; + concurrent_root_settings.per_constraint_residual = true; + set_pdlp_solver_mode(concurrent_root_settings); + branch_and_bound->configure_concurrent_lp_root_solve( + context.problem_ptr, + concurrent_root_settings, + context.settings.heuristic_params.root_lp_max_time, + context.settings.heuristic_params.root_lp_time_ratio); context.problem_ptr->branch_and_bound_callback = std::bind(&mip::branch_and_bound_t::set_solution_from_heuristics, @@ -461,17 +482,6 @@ solution_t mip_solver_t::run_solver() context.work_unit_scheduler_.register_context(branch_and_bound->get_work_unit_context()); // context.work_unit_scheduler_.verbose = true; - context.problem_ptr->set_root_relaxation_solution_callback = - std::bind(&mip::branch_and_bound_t::set_root_relaxation_solution, - branch_and_bound.get(), - std::placeholders::_1, - std::placeholders::_2, - std::placeholders::_3, - std::placeholders::_4, - std::placeholders::_5, - std::placeholders::_6, - std::placeholders::_7); - if (timer_.check_time_limit()) { CUOPT_LOG_INFO("Time limit reached during B&B setup"); context.stats.total_solve_time = timer_.elapsed_time(); From 45b717f0bf1fc70efb22549967dbd470bbf18cc6 Mon Sep 17 00:00:00 2001 From: Rajesh Gandham Date: Wed, 26 Aug 2026 15:53:28 -0700 Subject: [PATCH 2/4] Update the lower bound when PDLP/Barrier finish --- cpp/src/branch_and_bound/branch_and_bound.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/cpp/src/branch_and_bound/branch_and_bound.cpp b/cpp/src/branch_and_bound/branch_and_bound.cpp index 931df3cdfb..4b977a76c0 100644 --- a/cpp/src/branch_and_bound/branch_and_bound.cpp +++ b/cpp/src/branch_and_bound/branch_and_bound.cpp @@ -2845,6 +2845,9 @@ lp_status_t branch_and_bound_t::solve_root_relaxation( result.user_objective, result.iterations, result.method); + // Same as the old diversity-manager path: an Optimal GPU root LP is a + // valid MIP dual bound even if dual simplex / crossover has not finished. + if (result.optimal) { update_user_bound(result.solver_objective); } } } catch (const std::exception& e) { settings_.log.printf("Concurrent GPU root LP failed: %s\n", e.what()); From baad089f4e62e020f7374bf8af2932bce1eb554e Mon Sep 17 00:00:00 2001 From: Rajesh Gandham Date: Wed, 26 Aug 2026 15:58:18 -0700 Subject: [PATCH 3/4] Track concurrent GPU root solver sources. These files were used by CMake and the Optimal-bound hand-off but never committed, so they disappeared on branch checkout. --- .../concurrent_root_solver.cu | 60 +++++++++++++++++++ .../concurrent_root_solver.hpp | 37 ++++++++++++ 2 files changed, 97 insertions(+) create mode 100644 cpp/src/branch_and_bound/concurrent_root_solver.cu create mode 100644 cpp/src/branch_and_bound/concurrent_root_solver.hpp diff --git a/cpp/src/branch_and_bound/concurrent_root_solver.cu b/cpp/src/branch_and_bound/concurrent_root_solver.cu new file mode 100644 index 0000000000..98eefe7eaf --- /dev/null +++ b/cpp/src/branch_and_bound/concurrent_root_solver.cu @@ -0,0 +1,60 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include + +#include +#include + +#include + +namespace cuopt::mathematical_optimization::mip { + +template +concurrent_root_solution_t solve_concurrent_root_relaxation( + problem_t* problem, + const pdlp_solver_settings_t& settings, + f_t time_limit, + std::atomic* concurrent_halt) +{ + concurrent_root_solution_t result; + auto root_settings = settings; + root_settings.time_limit = time_limit; + root_settings.concurrent_halt = concurrent_halt; + + timer_t root_timer(time_limit); + auto lp_result = solve_lp_with_method(*problem, root_settings, root_timer); + const auto status = lp_result.get_termination_status(); + result.usable = status != pdlp_termination_status_t::NumericalError && + status != pdlp_termination_status_t::ConcurrentLimit && + lp_result.get_primal_solution().size() == static_cast(problem->n_variables) && + lp_result.get_dual_solution().size() == static_cast(problem->n_constraints); + result.optimal = status == pdlp_termination_status_t::Optimal; + if (!result.usable) { return result; } + + auto& d_primal = lp_result.get_primal_solution(); + auto& d_dual = lp_result.get_dual_solution(); + auto& d_reduced_cost = lp_result.get_reduced_cost(); + result.primal.resize(d_primal.size()); + result.dual.resize(d_dual.size()); + result.reduced_cost.resize(d_reduced_cost.size()); + auto stream = problem->handle_ptr->get_stream(); + raft::copy(result.primal.data(), d_primal.data(), d_primal.size(), stream); + raft::copy(result.dual.data(), d_dual.data(), d_dual.size(), stream); + raft::copy(result.reduced_cost.data(), d_reduced_cost.data(), d_reduced_cost.size(), stream); + problem->handle_ptr->sync_stream(); + + result.user_objective = lp_result.get_objective_value(); + result.solver_objective = problem->get_solver_obj_from_user_obj(result.user_objective); + result.iterations = lp_result.get_additional_termination_information().number_of_steps_taken; + result.method = lp_result.get_additional_termination_information().solved_by; + return result; +} + +template concurrent_root_solution_t solve_concurrent_root_relaxation( + problem_t*, const pdlp_solver_settings_t&, double, std::atomic*); + +} // namespace cuopt::mathematical_optimization::mip diff --git a/cpp/src/branch_and_bound/concurrent_root_solver.hpp b/cpp/src/branch_and_bound/concurrent_root_solver.hpp new file mode 100644 index 0000000000..9591f7ecf7 --- /dev/null +++ b/cpp/src/branch_and_bound/concurrent_root_solver.hpp @@ -0,0 +1,37 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include + +#include +#include + +namespace cuopt::mathematical_optimization::mip { + +template +class problem_t; + +template +struct concurrent_root_solution_t { + bool usable{false}; + bool optimal{false}; + std::vector primal; + std::vector dual; + std::vector reduced_cost; + f_t solver_objective{0}; + f_t user_objective{0}; + i_t iterations{0}; + method_t method{method_t::Unset}; +}; + +template +concurrent_root_solution_t solve_concurrent_root_relaxation( + problem_t* problem, + const pdlp_solver_settings_t& settings, + f_t time_limit, + std::atomic* concurrent_halt); + +} // namespace cuopt::mathematical_optimization::mip From fc317c2bd6d63617c16101c571bed434fb4629cf Mon Sep 17 00:00:00 2001 From: Rajesh Gandham Date: Wed, 26 Aug 2026 16:01:19 -0700 Subject: [PATCH 4/4] Cleanup --- cpp/src/branch_and_bound/concurrent_root_solver.cu | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/cpp/src/branch_and_bound/concurrent_root_solver.cu b/cpp/src/branch_and_bound/concurrent_root_solver.cu index 98eefe7eaf..88fa98c5a1 100644 --- a/cpp/src/branch_and_bound/concurrent_root_solver.cu +++ b/cpp/src/branch_and_bound/concurrent_root_solver.cu @@ -28,10 +28,11 @@ concurrent_root_solution_t solve_concurrent_root_relaxation( timer_t root_timer(time_limit); auto lp_result = solve_lp_with_method(*problem, root_settings, root_timer); const auto status = lp_result.get_termination_status(); - result.usable = status != pdlp_termination_status_t::NumericalError && - status != pdlp_termination_status_t::ConcurrentLimit && - lp_result.get_primal_solution().size() == static_cast(problem->n_variables) && - lp_result.get_dual_solution().size() == static_cast(problem->n_constraints); + result.usable = + status != pdlp_termination_status_t::NumericalError && + status != pdlp_termination_status_t::ConcurrentLimit && + lp_result.get_primal_solution().size() == static_cast(problem->n_variables) && + lp_result.get_dual_solution().size() == static_cast(problem->n_constraints); result.optimal = status == pdlp_termination_status_t::Optimal; if (!result.usable) { return result; }