From 4c6704f748493a3290a534e3a72ebc3ec93ca645 Mon Sep 17 00:00:00 2001 From: akifcorduk Date: Thu, 2 Oct 2025 08:48:56 -0700 Subject: [PATCH 01/10] remove left-over kernels and fix logs --- .../local_search/rounding/constraint_prop.cu | 2 +- cpp/src/mip/presolve/spmv_kernels.cuh | 229 ------------------ cpp/src/mip/presolve/third_party_presolve.cpp | 20 +- 3 files changed, 11 insertions(+), 240 deletions(-) delete mode 100644 cpp/src/mip/presolve/spmv_kernels.cuh diff --git a/cpp/src/mip/local_search/rounding/constraint_prop.cu b/cpp/src/mip/local_search/rounding/constraint_prop.cu index 4dfd1b216b..76d3916df6 100644 --- a/cpp/src/mip/local_search/rounding/constraint_prop.cu +++ b/cpp/src/mip/local_search/rounding/constraint_prop.cu @@ -907,7 +907,7 @@ bool constraint_prop_t::find_integer( CUOPT_LOG_DEBUG("Bounds propagation rounding: unset vars %lu", unset_integer_vars.size()); if (unset_integer_vars.size() == 0) { - CUOPT_LOG_ERROR("No integer variables provided in the bounds prop rounding"); + CUOPT_LOG_DEBUG("No integer variables provided in the bounds prop rounding"); expand_device_copy(orig_sol.assignment, sol.assignment, sol.handle_ptr->get_stream()); cuopt_func_call(orig_sol.test_variable_bounds()); return orig_sol.compute_feasibility(); diff --git a/cpp/src/mip/presolve/spmv_kernels.cuh b/cpp/src/mip/presolve/spmv_kernels.cuh deleted file mode 100644 index 3f11a8f365..0000000000 --- a/cpp/src/mip/presolve/spmv_kernels.cuh +++ /dev/null @@ -1,229 +0,0 @@ -/* - * SPDX-FileCopyrightText: Copyright (c) 2022-2025, NVIDIA CORPORATION & AFFILIATES. All rights - * reserved. SPDX-License-Identifier: Apache-2.0 - * - * Licensed under the Apache License, Version 2.0 (the "License"); - * you may not use this file except in compliance with the License. - * You may obtain a copy of the License at - * - * http://www.apache.org/licenses/LICENSE-2.0 - * - * Unless required by applicable law or agreed to in writing, software - * distributed under the License is distributed on an "AS IS" BASIS, - * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - * See the License for the specific language governing permissions and - * limitations under the License. - */ - -#pragma once - -#include -#include - -namespace cuopt::linear_programming::detail { - -template -__device__ f_t spmv(view_t view, raft::device_span input, i_t tid, i_t beg, i_t end) -{ - f_t out = 0.; - for (i_t i = tid + beg; i < end; i += MAX_EDGE_PER_CNST) { - auto coeff = view.coeff[i]; - auto var = view.elem[i]; - auto in = input[var]; - out += coeff * in; - } - return out; -} - -template -__global__ void lb_spmv_heavy_kernel(i_t id_range_beg, - raft::device_span ids, - raft::device_span pseudo_block_ids, - i_t work_per_block, - view_t view, - raft::device_span input, - raft::device_span tmp_out) -{ - auto idx = ids[blockIdx.x] + id_range_beg; - auto pseudo_block_id = pseudo_block_ids[blockIdx.x]; - i_t item_off_beg = view.offsets[idx] + work_per_block * pseudo_block_id; - i_t item_off_end = min(item_off_beg + work_per_block, view.offsets[idx + 1]); - - typedef cub::BlockReduce BlockReduce; - __shared__ typename BlockReduce::TempStorage temp_storage; - - auto out = spmv(view, input, threadIdx.x, item_off_beg, item_off_end); - - out = BlockReduce(temp_storage).Sum(out); - - if (threadIdx.x == 0) { tmp_out[blockIdx.x] = out; } -} - -template -__global__ void finalize_spmv_kernel(i_t heavy_beg_id, - raft::device_span item_offsets, - raft::device_span tmp_out, - view_t view, - raft::device_span output) -{ - using warp_reduce = cub::WarpReduce; - __shared__ typename warp_reduce::TempStorage temp_storage; - i_t idx = heavy_beg_id + blockIdx.x; - i_t item_idx = view.reorg_ids[idx]; - - i_t item_off_beg = item_offsets[blockIdx.x]; - i_t item_off_end = item_offsets[blockIdx.x + 1]; - f_t out = 0.; - for (i_t i = threadIdx.x + item_off_beg; i < item_off_end; i += blockDim.x) { - out += tmp_out[i]; - } - out = warp_reduce(temp_storage).Sum(out); - if (threadIdx.x == 0) { output[item_idx] = out; } -} - -template -__global__ void lb_spmv_block_kernel(i_t id_range_beg, - view_t view, - raft::device_span input, - raft::device_span output) - -{ - i_t idx = id_range_beg + blockIdx.x; - i_t item_idx = view.reorg_ids[idx]; - i_t item_off_beg = view.offsets[idx]; - i_t item_off_end = view.offsets[idx + 1]; - - typedef cub::BlockReduce BlockReduce; - __shared__ typename BlockReduce::TempStorage temp_storage; - - auto out = spmv(view, input, threadIdx.x, item_off_beg, item_off_end); - - out = BlockReduce(temp_storage).Sum(out); - - if (threadIdx.x == 0) { - // written to old index - output[item_idx] = out; - } -} - -template -__global__ void lb_spmv_sub_warp_kernel(i_t id_range_beg, - i_t id_range_end, - view_t view, - raft::device_span input, - raft::device_span output) -{ - constexpr i_t ids_per_block = BDIM / MAX_EDGE_PER_CNST; - i_t id_beg = blockIdx.x * ids_per_block + id_range_beg; - i_t idx = id_beg + (threadIdx.x / MAX_EDGE_PER_CNST); - i_t item_idx; - if (idx < id_range_end) { item_idx = view.reorg_ids[idx]; } - i_t p_tid = threadIdx.x % MAX_EDGE_PER_CNST; - - i_t head_flag = (p_tid == 0); - - using warp_reduce = cub::WarpReduce; - __shared__ typename warp_reduce::TempStorage temp_storage; - - f_t out = 0.; - - if (idx < id_range_end) { - i_t item_off_beg = view.offsets[idx]; - i_t item_off_end = view.offsets[idx + 1]; - out = spmv(view, input, p_tid, item_off_beg, item_off_end); - } - - out = warp_reduce(temp_storage).Sum(out); - - if (head_flag && (idx < id_range_end)) { output[item_idx] = out; } -} - -#if 1 - -#define BYTE_TO_BINARY(byte) \ - ((byte) & 0x80 ? '1' : '0'), ((byte) & 0x40 ? '1' : '0'), ((byte) & 0x20 ? '1' : '0'), \ - ((byte) & 0x10 ? '1' : '0'), ((byte) & 0x08 ? '1' : '0'), ((byte) & 0x04 ? '1' : '0'), \ - ((byte) & 0x02 ? '1' : '0'), ((byte) & 0x01 ? '1' : '0') - -template -__device__ __forceinline__ void get_sub_warp_bin(i_t* id_warp_beg, - i_t* id_range_end, - i_t* t_p_v, - raft::device_span warp_offsets, - raft::device_span bin_offsets) -{ - i_t warp_id = (blockDim.x * blockIdx.x + threadIdx.x) / 32; - i_t lane_id = threadIdx.x & 31; - bool pred = false; - if (lane_id < warp_offsets.size()) { pred = (warp_id >= warp_offsets[lane_id]); } - unsigned int m = __ballot_sync(0xffffffff, pred); - i_t seg = 31 - __clz(m); - i_t it_per_warp = (1 << (5 - seg)); // item per warp = 32/(2^seg) - if (5 - seg < 0) { - *t_p_v = 0; - return; - } - i_t beg = bin_offsets[seg] + (warp_id - warp_offsets[seg]) * it_per_warp; - i_t end = bin_offsets[seg + 1]; - *id_warp_beg = beg; - *id_range_end = end; - *t_p_v = (1 << seg); -} - -template -__device__ void spmv_sub_warp(i_t id_warp_beg, - i_t id_range_end, - view_t view, - raft::device_span input, - raft::device_span output) -{ - i_t lane_id = (threadIdx.x & 31); - i_t idx = id_warp_beg + (lane_id / MAX_EDGE_PER_CNST); - i_t item_idx; - if (idx < id_range_end) { item_idx = view.reorg_ids[idx]; } - i_t p_tid = lane_id & (MAX_EDGE_PER_CNST - 1); - - i_t head_flag = (p_tid == 0); - - using warp_reduce = cub::WarpReduce; - __shared__ typename warp_reduce::TempStorage temp_storage; - - f_t out = 0.; - - if (idx < id_range_end) { - i_t item_off_beg = view.offsets[idx]; - i_t item_off_end = view.offsets[idx + 1]; - out = spmv(view, input, p_tid, item_off_beg, item_off_end); - } - - out = warp_reduce(temp_storage).Sum(out); - - if (head_flag && (idx < id_range_end)) { output[item_idx] = out; } -} - -template -__global__ void lb_spmv_sub_warp_kernel(view_t view, - raft::device_span input, - raft::device_span output, - raft::device_span warp_item_offsets, - raft::device_span warp_item_id_offsets) -{ - i_t id_warp_beg, id_range_end, t_p_v; - get_sub_warp_bin( - &id_warp_beg, &id_range_end, &t_p_v, warp_item_offsets, warp_item_id_offsets); - - if (t_p_v == 1) { - spmv_sub_warp(id_warp_beg, id_range_end, view, input, output); - } else if (t_p_v == 2) { - spmv_sub_warp(id_warp_beg, id_range_end, view, input, output); - } else if (t_p_v == 4) { - spmv_sub_warp(id_warp_beg, id_range_end, view, input, output); - } else if (t_p_v == 8) { - spmv_sub_warp(id_warp_beg, id_range_end, view, input, output); - } else if (t_p_v == 16) { - spmv_sub_warp(id_warp_beg, id_range_end, view, input, output); - } -} -#endif - -} // namespace cuopt::linear_programming::detail diff --git a/cpp/src/mip/presolve/third_party_presolve.cpp b/cpp/src/mip/presolve/third_party_presolve.cpp index dc2d4b00e8..88de4912af 100644 --- a/cpp/src/mip/presolve/third_party_presolve.cpp +++ b/cpp/src/mip/presolve/third_party_presolve.cpp @@ -274,19 +274,19 @@ void check_presolve_status(const papilo::PresolveStatus& status) { switch (status) { case papilo::PresolveStatus::kUnchanged: - CUOPT_LOG_INFO("Presolve status:: did not result in any changes"); + CUOPT_LOG_INFO("Presolve status: did not result in any changes"); break; case papilo::PresolveStatus::kReduced: - CUOPT_LOG_INFO("Presolve status:: reduced the problem"); + CUOPT_LOG_INFO("Presolve status: reduced the problem"); break; case papilo::PresolveStatus::kUnbndOrInfeas: - CUOPT_LOG_INFO("Presolve status:: found an unbounded or infeasible problem"); + CUOPT_LOG_INFO("Presolve status: found an unbounded or infeasible problem"); break; case papilo::PresolveStatus::kInfeasible: - CUOPT_LOG_INFO("Presolve status:: found an infeasible problem"); + CUOPT_LOG_INFO("Presolve status: found an infeasible problem"); break; case papilo::PresolveStatus::kUnbounded: - CUOPT_LOG_INFO("Presolve status:: found an unbounded problem"); + CUOPT_LOG_INFO("Presolve status: found an unbounded problem"); break; } } @@ -294,10 +294,10 @@ void check_presolve_status(const papilo::PresolveStatus& status) void check_postsolve_status(const papilo::PostsolveStatus& status) { switch (status) { - case papilo::PostsolveStatus::kOk: CUOPT_LOG_INFO("Post-solve status:: succeeded"); break; + case papilo::PostsolveStatus::kOk: CUOPT_LOG_INFO("Post-solve status: succeeded"); break; case papilo::PostsolveStatus::kFailed: CUOPT_LOG_INFO( - "Post-solve status:: Post solved solution violates constraints. This is most likely due to " + "Post-solve status: Post solved solution violates constraints. This is most likely due to " "different tolerances."); break; } @@ -362,7 +362,7 @@ std::pair, bool> third_party_presolve_t papilo_problem = build_papilo_problem(op_problem); - CUOPT_LOG_INFO("Unpresolved problem:: %d constraints, %d variables, %d nonzeros", + CUOPT_LOG_INFO("Unpresolved problem: %d constraints, %d variables, %d nonzeros", papilo_problem.getNRows(), papilo_problem.getNCols(), papilo_problem.getConstraintMatrix().getNnz()); @@ -382,11 +382,11 @@ std::pair, bool> third_party_presolve_t(op_problem.get_handle_ptr()), false); } post_solve_storage_ = result.postsolve; - CUOPT_LOG_INFO("Presolve removed:: %d constraints, %d variables, %d nonzeros", + CUOPT_LOG_INFO("Presolve removed: %d constraints, %d variables, %d nonzeros", op_problem.get_n_constraints() - papilo_problem.getNRows(), op_problem.get_n_variables() - papilo_problem.getNCols(), op_problem.get_nnz() - papilo_problem.getConstraintMatrix().getNnz()); - CUOPT_LOG_INFO("Presolved problem:: %d constraints, %d variables, %d nonzeros", + CUOPT_LOG_INFO("Presolved problem: %d constraints, %d variables, %d nonzeros", papilo_problem.getNRows(), papilo_problem.getNCols(), papilo_problem.getConstraintMatrix().getNnz()); From d2baad3b9ee37d67e6ce92c47c352a2d16a23720 Mon Sep 17 00:00:00 2001 From: akifcorduk Date: Thu, 2 Oct 2025 01:23:33 -0700 Subject: [PATCH 02/10] fix gpu count --- benchmarks/linear_programming/cuopt/run_mip.cpp | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/benchmarks/linear_programming/cuopt/run_mip.cpp b/benchmarks/linear_programming/cuopt/run_mip.cpp index a4f52cb4ed..4a6121f8e2 100644 --- a/benchmarks/linear_programming/cuopt/run_mip.cpp +++ b/benchmarks/linear_programming/cuopt/run_mip.cpp @@ -80,10 +80,11 @@ void merge_result_files(const std::string& out_dir, void write_to_output_file(const std::string& out_dir, const std::string& base_filename, int gpu_id, + int n_gpus, int batch_id, const std::string& data) { - int output_id = batch_id * 8 + gpu_id; + int output_id = batch_id * n_gpus + gpu_id; std::string filename = out_dir + "/result_" + std::to_string(output_id) + ".txt"; std::ofstream outfile(filename, std::ios_base::app); if (outfile.is_open()) { @@ -149,6 +150,7 @@ std::vector> read_solution_from_dir(const std::string file_p int run_single_file(std::string file_path, int device, int batch_id, + int n_gpus, std::string out_dir, std::optional initial_solution_dir, bool heuristics_only, @@ -243,7 +245,7 @@ int run_single_file(std::string file_path, << obj_val << "," << benchmark_info.objective_of_initial_population << "," << benchmark_info.last_improvement_of_best_feasible << "," << benchmark_info.last_improvement_after_recombination << "\n"; - write_to_output_file(out_dir, base_filename, device, batch_id, ss.str()); + write_to_output_file(out_dir, base_filename, device, n_gpus, batch_id, ss.str()); CUOPT_LOG_INFO("Results written to the file %s", base_filename.c_str()); return sol_found; } @@ -251,6 +253,7 @@ int run_single_file(std::string file_path, void run_single_file_mp(std::string file_path, int device, int batch_id, + int n_gpus, std::string out_dir, std::optional input_file_dir, bool heuristics_only, @@ -265,6 +268,7 @@ void run_single_file_mp(std::string file_path, int sol_found = run_single_file(file_path, device, batch_id, + n_gpus, out_dir, input_file_dir, heuristics_only, @@ -462,6 +466,7 @@ int main(int argc, char* argv[]) run_single_file_mp(file_name, gpu_id, batch_num, + n_gpus, out_dir, initial_solution_file, heuristics_only, @@ -501,6 +506,7 @@ int main(int argc, char* argv[]) run_single_file(path, 0, 0, + n_gpus, out_dir, initial_solution_file, heuristics_only, From 1c27d01d7f3999bf0fca16bcea5ed6f20ad8a8c4 Mon Sep 17 00:00:00 2001 From: nicolas Date: Thu, 9 Oct 2025 16:13:34 +0200 Subject: [PATCH 03/10] fix starting variable bounds for diving. add backtracking parameter. --- cpp/src/dual_simplex/branch_and_bound.cpp | 34 ++++++++------- cpp/src/dual_simplex/branch_and_bound.hpp | 52 ++++++++++++++++++----- 2 files changed, 60 insertions(+), 26 deletions(-) diff --git a/cpp/src/dual_simplex/branch_and_bound.cpp b/cpp/src/dual_simplex/branch_and_bound.cpp index cf6fd69798..e306ad4969 100644 --- a/cpp/src/dual_simplex/branch_and_bound.cpp +++ b/cpp/src/dual_simplex/branch_and_bound.cpp @@ -576,10 +576,7 @@ node_status_t branch_and_bound_t::solve_node(search_tree_t& // two vectors at each node and potentially cause memory issues node_ptr->get_variable_bounds(leaf_problem.lower, leaf_problem.upper, bounds_changed); - i_t node_iter = 0; - f_t lp_start_time = tic(); - std::vector leaf_edge_norms = edge_norms_; // = node.steepest_edge_norms; - + std::vector leaf_edge_norms = edge_norms_; // = node.steepest_edge_norms; simplex_solver_settings_t lp_settings = settings_; lp_settings.set_log(false); lp_settings.cut_off = upper_bound + settings_.dual_tol; @@ -594,6 +591,9 @@ node_status_t branch_and_bound_t::solve_node(search_tree_t& dual::status_t lp_status = dual::status_t::DUAL_UNBOUNDED; if (feasible) { + i_t node_iter = 0; + f_t lp_start_time = tic(); + lp_status = dual_phase2(2, 0, lp_start_time, @@ -610,10 +610,10 @@ node_status_t branch_and_bound_t::solve_node(search_tree_t& leaf_problem, lp_start_time, lp_settings, leaf_solution, leaf_vstatus, leaf_edge_norms); lp_status = convert_lp_status_to_dual_status(second_status); } - } - stats_.total_lp_solve_time += toc(lp_start_time); - stats_.total_lp_iters += node_iter; + stats_.total_lp_solve_time += toc(lp_start_time); + stats_.total_lp_iters += node_iter; + } if (lp_status == dual::status_t::DUAL_UNBOUNDED) { // Node was infeasible. Do not branch @@ -866,7 +866,7 @@ void branch_and_bound_t::explore_subtree(i_t id, // would be better if we discard the node instead. if (get_heap_size() > settings_.num_bfs_threads) { mutex_dive_queue_.lock(); - dive_queue_.push(node->detach_copy()); + dive_queue_.emplace(node->detach_copy(), leaf_problem.lower, leaf_problem.upper); mutex_dive_queue_.unlock(); } @@ -943,23 +943,24 @@ void branch_and_bound_t::best_first_thread(i_t id, template void branch_and_bound_t::diving_thread(lp_problem_t& leaf_problem, - const csc_matrix_t& Arow) + const csc_matrix_t& Arow, + i_t backtracking) { logger_t log; log.log = false; while (status_ == mip_exploration_status_t::RUNNING && (active_subtrees_ > 0 || get_heap_size() > 0)) { - std::optional> start_node; + std::optional> start_node; mutex_dive_queue_.lock(); if (dive_queue_.size() > 0) { start_node = dive_queue_.pop(); } mutex_dive_queue_.unlock(); if (start_node.has_value()) { - if (get_upper_bound() < start_node->lower_bound) { continue; } + if (get_upper_bound() < start_node->node.lower_bound) { continue; } - search_tree_t subtree(std::move(start_node.value())); + search_tree_t subtree(std::move(start_node->node)); std::deque*> stack; stack.push_front(&subtree.root); @@ -985,16 +986,19 @@ void branch_and_bound_t::diving_thread(lp_problem_t& leaf_pr auto [first, second] = child_selection(node_ptr); stack.push_front(second); stack.push_front(first); + } + if (stack.size() > 1) { // If the diving thread is consuming the nodes faster than the // best first search, then we split the current subtree at the // lowest possible point and move to the queue, so it can // be picked by another thread. - if (dive_queue_.size() < min_diving_queue_size_) { + if (dive_queue_.size() < min_diving_queue_size_ || + (stack.front()->depth - stack.back()->depth) > backtracking) { mutex_dive_queue_.lock(); mip_node_t* new_node = stack.back(); stack.pop_back(); - dive_queue_.push(new_node->detach_copy()); + dive_queue_.emplace(new_node->detach_copy(), leaf_problem.lower, leaf_problem.upper); mutex_dive_queue_.unlock(); } } @@ -1192,7 +1196,7 @@ mip_status_t branch_and_bound_t::solve(mip_solution_t& solut for (i_t i = 0; i < settings_.num_diving_threads; i++) { #pragma omp task - diving_thread(leaf_problem, Arow); + diving_thread(leaf_problem, Arow, 10); } } } diff --git a/cpp/src/dual_simplex/branch_and_bound.hpp b/cpp/src/dual_simplex/branch_and_bound.hpp index 7b80f88fa9..ae70172e36 100644 --- a/cpp/src/dual_simplex/branch_and_bound.hpp +++ b/cpp/src/dual_simplex/branch_and_bound.hpp @@ -55,36 +55,64 @@ enum class mip_exploration_status_t { template void upper_bound_callback(f_t upper_bound); +template +struct diving_root_t { + mip_node_t node; + std::vector lp_lower; + std::vector lp_upper; + + diving_root_t(mip_node_t&& node, + const std::vector& lower, + const std::vector& upper) + : node(std::move(node)), lp_upper(upper), lp_lower(lower) + { + } + + friend bool operator>(const diving_root_t& a, const diving_root_t& b) + { + return a.node.lower_bound > b.node.lower_bound; + } +}; + // A min-heap for storing the starting nodes for the dives. -// This has a maximum size of 8192, such that the container +// This has a maximum size of 256, such that the container // will discard the least promising node if the queue is full. template class dive_queue_t { private: - std::vector> buffer; - static constexpr i_t max_size_ = 2048; + std::vector> buffer; + static constexpr i_t max_size_ = 256; public: dive_queue_t() { buffer.reserve(max_size_); } - void push(mip_node_t&& node) + void push(diving_root_t&& node) { buffer.push_back(std::move(node)); - std::push_heap(buffer.begin(), buffer.end(), node_compare_t()); + std::push_heap(buffer.begin(), buffer.end(), std::greater<>()); + if (buffer.size() > max_size()) { buffer.pop_back(); } + } + + void emplace(mip_node_t&& node, + const std::vector& lower, + const std::vector& upper) + { + buffer.emplace_back(std::move(node), lower, upper); + std::push_heap(buffer.begin(), buffer.end(), std::greater<>()); if (buffer.size() > max_size()) { buffer.pop_back(); } } - mip_node_t pop() + diving_root_t pop() { - std::pop_heap(buffer.begin(), buffer.end(), node_compare_t()); - mip_node_t node = std::move(buffer.back()); + std::pop_heap(buffer.begin(), buffer.end(), std::greater<>()); + diving_root_t node = std::move(buffer.back()); buffer.pop_back(); return node; } i_t size() const { return buffer.size(); } constexpr i_t max_size() const { return max_size_; } - const mip_node_t& top() const { return buffer.front(); } + const diving_root_t& top() const { return buffer.front(); } void clear() { buffer.clear(); } }; @@ -188,7 +216,7 @@ class branch_and_bound_t { // Set the final solution. mip_status_t set_final_solution(mip_solution_t& solution, f_t lower_bound); - // Update the incumbent solution with the new feasible solution. + // Update the incumbent solution with the new feasible solution // found during branch and bound. void add_feasible_solution(f_t leaf_objective, const std::vector& leaf_solution, @@ -222,7 +250,9 @@ class branch_and_bound_t { // Each diving thread pops the first node from the dive queue and then performs // a deep dive into the subtree determined by the node. - void diving_thread(lp_problem_t& leaf_problem, const csc_matrix_t& Arow); + void diving_thread(lp_problem_t& leaf_problem, + const csc_matrix_t& Arow, + i_t backtracking); // Solve the LP relaxation of a leaf node and update the tree. node_status_t solve_node(search_tree_t& search_tree, From 86a865516ef1a83a8350dab9a48de108de58310b Mon Sep 17 00:00:00 2001 From: nicolas Date: Thu, 9 Oct 2025 17:12:11 +0200 Subject: [PATCH 04/10] fix log during the ramp up phase --- cpp/src/dual_simplex/branch_and_bound.cpp | 30 +++++++++++------------ 1 file changed, 15 insertions(+), 15 deletions(-) diff --git a/cpp/src/dual_simplex/branch_and_bound.cpp b/cpp/src/dual_simplex/branch_and_bound.cpp index e306ad4969..c095c92715 100644 --- a/cpp/src/dual_simplex/branch_and_bound.cpp +++ b/cpp/src/dual_simplex/branch_and_bound.cpp @@ -695,17 +695,14 @@ void branch_and_bound_t::exploration_ramp_up(search_tree_t* i_t initial_heap_size) { if (status_ != mip_exploration_status_t::RUNNING) { return; } - if (omp_get_thread_num() == 0) { repair_heuristic_solutions(); } + repair_heuristic_solutions(); f_t lower_bound = node->lower_bound; f_t upper_bound = get_upper_bound(); f_t rel_gap = user_relative_gap(original_lp_, upper_bound, lower_bound); f_t abs_gap = upper_bound - lower_bound; - i_t nodes_explored = 0; - i_t nodes_unexplored = 0; - - nodes_explored = (stats_.nodes_explored++); - nodes_unexplored = (stats_.nodes_unexplored--); + i_t nodes_explored = (++stats_.nodes_explored); + i_t nodes_unexplored = (--stats_.nodes_unexplored); stats_.nodes_since_last_log++; if (lower_bound > upper_bound || rel_gap < settings_.relative_mip_gap_tol) { @@ -716,12 +713,17 @@ void branch_and_bound_t::exploration_ramp_up(search_tree_t* f_t now = toc(stats_.start_time); - if (omp_get_thread_num() == 0) { - f_t time_since_last_log = stats_.last_log == 0 ? 1.0 : toc(stats_.last_log); + f_t time_since_last_log = stats_.last_log == 0 ? 1.0 : toc(stats_.last_log); + + if (((stats_.nodes_since_last_log >= 10 || abs_gap < 10 * settings_.absolute_mip_gap_tol) && + (time_since_last_log >= 1)) || + (time_since_last_log > 30) || now > settings_.time_limit) { + // Check if no new node was explored until now. If this is the case, + // only the last thread should report the progress + if (stats_.nodes_explored.load() == nodes_explored) { + stats_.nodes_since_last_log = 0; + stats_.last_log = tic(); - if (((stats_.nodes_since_last_log >= 10 || abs_gap < 10 * settings_.absolute_mip_gap_tol) && - (time_since_last_log >= 1)) || - (time_since_last_log > 30) || now > settings_.time_limit) { f_t obj = compute_user_objective(original_lp_, upper_bound); f_t user_lower = compute_user_objective(original_lp_, root_objective_); std::string gap_user = user_mip_gap(obj, user_lower); @@ -735,8 +737,6 @@ void branch_and_bound_t::exploration_ramp_up(search_tree_t* nodes_explored > 0 ? stats_.total_lp_iters / nodes_explored : 0, gap_user.c_str(), now); - - stats_.nodes_since_last_log = 0; } } @@ -802,8 +802,8 @@ void branch_and_bound_t::explore_subtree(i_t id, // - The lower bound of the parent is lower or equal to its children assert(id < local_lower_bounds_.size()); local_lower_bounds_[id] = lower_bound; - i_t nodes_explored = stats_.nodes_explored++; - i_t nodes_unexplored = stats_.nodes_unexplored--; + i_t nodes_explored = (++stats_.nodes_explored); + i_t nodes_unexplored = (--stats_.nodes_unexplored); stats_.nodes_since_last_log++; if (lower_bound > upper_bound || rel_gap < settings_.relative_mip_gap_tol) { From 4044cc85ba589322532689a407a00a9e4714dabf Mon Sep 17 00:00:00 2001 From: nicolas Date: Thu, 9 Oct 2025 17:24:03 +0200 Subject: [PATCH 05/10] added comment --- cpp/src/dual_simplex/branch_and_bound.cpp | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/cpp/src/dual_simplex/branch_and_bound.cpp b/cpp/src/dual_simplex/branch_and_bound.cpp index c095c92715..efb4e5b4c7 100644 --- a/cpp/src/dual_simplex/branch_and_bound.cpp +++ b/cpp/src/dual_simplex/branch_and_bound.cpp @@ -695,6 +695,10 @@ void branch_and_bound_t::exploration_ramp_up(search_tree_t* i_t initial_heap_size) { if (status_ != mip_exploration_status_t::RUNNING) { return; } + + // Note that we do not know which thread will execute the + // `exploration_ramp_up` task, so we allow to any thread + // to repair the heuristic solution. repair_heuristic_solutions(); f_t lower_bound = node->lower_bound; @@ -711,8 +715,7 @@ void branch_and_bound_t::exploration_ramp_up(search_tree_t* return; } - f_t now = toc(stats_.start_time); - + f_t now = toc(stats_.start_time); f_t time_since_last_log = stats_.last_log == 0 ? 1.0 : toc(stats_.last_log); if (((stats_.nodes_since_last_log >= 10 || abs_gap < 10 * settings_.absolute_mip_gap_tol) && @@ -784,7 +787,7 @@ void branch_and_bound_t::explore_subtree(i_t id, stack.push_front(start_node); while (stack.size() > 0 && status_ == mip_exploration_status_t::RUNNING) { - if (omp_get_thread_num() == 0) { repair_heuristic_solutions(); } + if (id == 0) { repair_heuristic_solutions(); } mip_node_t* node_ptr = stack.front(); stack.pop_front(); From 1651db35e3f2f2618b93053faaaa0151bb532c94 Mon Sep 17 00:00:00 2001 From: nicolas Date: Thu, 9 Oct 2025 20:53:27 +0200 Subject: [PATCH 06/10] removed backtracking parameter due to a performance regression --- cpp/src/dual_simplex/branch_and_bound.cpp | 8 +++----- cpp/src/dual_simplex/branch_and_bound.hpp | 4 +--- 2 files changed, 4 insertions(+), 8 deletions(-) diff --git a/cpp/src/dual_simplex/branch_and_bound.cpp b/cpp/src/dual_simplex/branch_and_bound.cpp index efb4e5b4c7..0df4c082bc 100644 --- a/cpp/src/dual_simplex/branch_and_bound.cpp +++ b/cpp/src/dual_simplex/branch_and_bound.cpp @@ -946,8 +946,7 @@ void branch_and_bound_t::best_first_thread(i_t id, template void branch_and_bound_t::diving_thread(lp_problem_t& leaf_problem, - const csc_matrix_t& Arow, - i_t backtracking) + const csc_matrix_t& Arow) { logger_t log; log.log = false; @@ -996,8 +995,7 @@ void branch_and_bound_t::diving_thread(lp_problem_t& leaf_pr // best first search, then we split the current subtree at the // lowest possible point and move to the queue, so it can // be picked by another thread. - if (dive_queue_.size() < min_diving_queue_size_ || - (stack.front()->depth - stack.back()->depth) > backtracking) { + if (dive_queue_.size() < min_diving_queue_size_) { mutex_dive_queue_.lock(); mip_node_t* new_node = stack.back(); stack.pop_back(); @@ -1199,7 +1197,7 @@ mip_status_t branch_and_bound_t::solve(mip_solution_t& solut for (i_t i = 0; i < settings_.num_diving_threads; i++) { #pragma omp task - diving_thread(leaf_problem, Arow, 10); + diving_thread(leaf_problem, Arow); } } } diff --git a/cpp/src/dual_simplex/branch_and_bound.hpp b/cpp/src/dual_simplex/branch_and_bound.hpp index ae70172e36..5453e8b424 100644 --- a/cpp/src/dual_simplex/branch_and_bound.hpp +++ b/cpp/src/dual_simplex/branch_and_bound.hpp @@ -250,9 +250,7 @@ class branch_and_bound_t { // Each diving thread pops the first node from the dive queue and then performs // a deep dive into the subtree determined by the node. - void diving_thread(lp_problem_t& leaf_problem, - const csc_matrix_t& Arow, - i_t backtracking); + void diving_thread(lp_problem_t& leaf_problem, const csc_matrix_t& Arow); // Solve the LP relaxation of a leaf node and update the tree. node_status_t solve_node(search_tree_t& search_tree, From 717e9a474430f7c7c0086b361289283803e37415 Mon Sep 17 00:00:00 2001 From: nicolas Date: Thu, 9 Oct 2025 22:07:20 +0200 Subject: [PATCH 07/10] fix missing variable bounds --- cpp/src/dual_simplex/branch_and_bound.cpp | 17 +++++++++++++---- 1 file changed, 13 insertions(+), 4 deletions(-) diff --git a/cpp/src/dual_simplex/branch_and_bound.cpp b/cpp/src/dual_simplex/branch_and_bound.cpp index 0df4c082bc..78acd07426 100644 --- a/cpp/src/dual_simplex/branch_and_bound.cpp +++ b/cpp/src/dual_simplex/branch_and_bound.cpp @@ -566,10 +566,6 @@ node_status_t branch_and_bound_t::solve_node(search_tree_t& lp_solution_t leaf_solution(leaf_problem.num_rows, leaf_problem.num_cols); assert(leaf_vstatus.size() == leaf_problem.num_cols); - // Set the correct bounds for the leaf problem - leaf_problem.lower = original_lp_.lower; - leaf_problem.upper = original_lp_.upper; - std::vector bounds_changed(leaf_problem.num_cols, false); // Technically, we can get the already strengthened bounds from the node/parent instead of // getting it from the original problem and re-strengthening. But this requires storing @@ -747,6 +743,11 @@ void branch_and_bound_t::exploration_ramp_up(search_tree_t* status_ = mip_exploration_status_t::TIME_LIMIT; return; } + + // Set the correct bounds for the leaf problem + leaf_problem.lower = original_lp_.lower; + leaf_problem.upper = original_lp_.upper; + node_status_t node_status = solve_node(*search_tree, node, leaf_problem, Arow, upper_bound, settings_.log, 'B'); @@ -845,6 +846,10 @@ void branch_and_bound_t::explore_subtree(i_t id, return; } + // Set the correct bounds for the leaf problem + leaf_problem.lower = original_lp_.lower; + leaf_problem.upper = original_lp_.upper; + node_status_t node_status = solve_node(search_tree, node_ptr, leaf_problem, Arow, upper_bound, settings_.log, 'B'); @@ -978,6 +983,10 @@ void branch_and_bound_t::diving_thread(lp_problem_t& leaf_pr if (toc(stats_.start_time) > settings_.time_limit) { return; } + // Set the correct bounds for the leaf problem + leaf_problem.lower = start_node->lp_lower; + leaf_problem.upper = start_node->lp_upper; + node_status_t node_status = solve_node(subtree, node_ptr, leaf_problem, Arow, upper_bound, log, 'D'); From ff4cde2b54ae8ed2da4e8b4c58532e5c41076ffb Mon Sep 17 00:00:00 2001 From: akifcorduk Date: Fri, 10 Oct 2025 05:08:44 -0700 Subject: [PATCH 08/10] clique merge bug --- cpp/CMakeLists.txt | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index cb17f0c4a2..90bb1c57f7 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -190,9 +190,9 @@ FetchContent_Declare( # does not have some of the presolvers and settings that we need # Mainly, probing and clique merging. # This is the reason we are using the development branch - # commit from Oct 8, 2025. Once these changes are merged into the main branch, + # commit from cliquemergebug branch. Once these changes are merged into the main branch, #we can switch to the main branch. - GIT_TAG "24ccf5752656df0f15dd9aabe5b97feae829b9ec" + GIT_TAG "8f710e33d352bf319d30b9c57e70516222f3f5ca" GIT_PROGRESS TRUE SYSTEM ) @@ -201,8 +201,6 @@ find_package(TBB REQUIRED) set(BUILD_TESTING OFF CACHE BOOL "Disable test build for papilo") set(PAPILO_NO_BINARIES ON) option(LUSOL "Disable LUSOL" OFF) -# Disable TBB because of a bug in CliqueMerging parallel version -set(TBB OFF CACHE BOOL "Disable TBB for papilo") FetchContent_MakeAvailable(papilo) From c1ff4ea011705cceb64234efcc0b99cab7cc9f60 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Fri, 10 Oct 2025 10:32:38 -0500 Subject: [PATCH 09/10] push changes --- ci/test_python.sh | 3 +++ ci/test_wheel_cuopt_server.sh | 3 +++ python/cuopt/cuopt/tests/linear_programming/test_python_API.py | 1 + 3 files changed, 7 insertions(+) diff --git a/ci/test_python.sh b/ci/test_python.sh index 0d3d1e5963..7d504f4738 100755 --- a/ci/test_python.sh +++ b/ci/test_python.sh @@ -60,6 +60,9 @@ EXITCODE=0 trap "EXITCODE=1" ERR set +e +# Due to race condition in certain cases UCX might not be able to cleanup properly, so we set the number of threads to 1 +export OMP_NUM_THREADS=1 + rapids-logger "Test cuopt_cli" timeout 10m bash ./python/libcuopt/libcuopt/tests/test_cli.sh diff --git a/ci/test_wheel_cuopt_server.sh b/ci/test_wheel_cuopt_server.sh index 5f0b874ba3..de4a52f479 100755 --- a/ci/test_wheel_cuopt_server.sh +++ b/ci/test_wheel_cuopt_server.sh @@ -37,4 +37,7 @@ rapids-pip-retry install \ ./datasets/linear_programming/download_pdlp_test_dataset.sh ./datasets/mip/download_miplib_test_dataset.sh +# Due to race condition in certain cases UCX might not be able to cleanup properly, so we set the number of threads to 1 +export OMP_NUM_THREADS=1 + RAPIDS_DATASET_ROOT_DIR=./datasets timeout 30m python -m pytest --verbose --capture=no ./python/cuopt_server/cuopt_server/tests/ diff --git a/python/cuopt/cuopt/tests/linear_programming/test_python_API.py b/python/cuopt/cuopt/tests/linear_programming/test_python_API.py index c7ef8b99bf..42059bf3d3 100644 --- a/python/cuopt/cuopt/tests/linear_programming/test_python_API.py +++ b/python/cuopt/cuopt/tests/linear_programming/test_python_API.py @@ -406,6 +406,7 @@ def test_warm_start(): settings = SolverSettings() settings.set_parameter(CUOPT_PDLP_SOLVER_MODE, PDLPSolverMode.Stable2) + settings.set_parameter(CUOPT_METHOD, SolverMethod.PDLP) settings.set_optimality_tolerance(1e-3) settings.set_parameter(CUOPT_INFEASIBILITY_DETECTION, False) From ff5cfa129df295565b3dc7797362e7f9e0f854d1 Mon Sep 17 00:00:00 2001 From: Ramakrishna Prabhu Date: Fri, 10 Oct 2025 12:58:36 -0500 Subject: [PATCH 10/10] fix test in server as well --- python/cuopt_server/cuopt_server/tests/test_pdlp_warmstart.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/python/cuopt_server/cuopt_server/tests/test_pdlp_warmstart.py b/python/cuopt_server/cuopt_server/tests/test_pdlp_warmstart.py index be67894be1..cfc30fa1c8 100644 --- a/python/cuopt_server/cuopt_server/tests/test_pdlp_warmstart.py +++ b/python/cuopt_server/cuopt_server/tests/test_pdlp_warmstart.py @@ -22,6 +22,7 @@ from cuopt.linear_programming.solver.solver_parameters import ( CUOPT_INFEASIBILITY_DETECTION, CUOPT_PDLP_SOLVER_MODE, + CUOPT_METHOD, ) from cuopt.linear_programming.solver_settings import PDLPSolverMode @@ -45,6 +46,7 @@ def test_warmstart(cuoptproc): # noqa settings.set_optimality_tolerance(1e-4) settings.set_parameter(CUOPT_INFEASIBILITY_DETECTION, False) settings.set_parameter(CUOPT_PDLP_SOLVER_MODE, PDLPSolverMode.Stable2) + settings.set_parameter(CUOPT_METHOD, SolverMethod.PDLP) data["solver_config"] = settings.toDict() headers = {"CLIENT-VERSION": "custom"}