From 4c6704f748493a3290a534e3a72ebc3ec93ca645 Mon Sep 17 00:00:00 2001 From: akifcorduk Date: Thu, 2 Oct 2025 08:48:56 -0700 Subject: [PATCH 1/7] 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 4dfd1b216..76d3916df 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 3f11a8f36..000000000 --- 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 dc2d4b00e..88de4912a 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 64ef5cbaf880c461dd3d865f3341e23f72387a2b Mon Sep 17 00:00:00 2001 From: akifcorduk Date: Mon, 6 Oct 2025 09:52:18 -0700 Subject: [PATCH 2/7] fix_compile_and_log --- benchmarks/linear_programming/cuopt/run_mip.cpp | 4 ++-- cpp/src/dual_simplex/device_sparse_matrix.cuh | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/benchmarks/linear_programming/cuopt/run_mip.cpp b/benchmarks/linear_programming/cuopt/run_mip.cpp index a4f52cb4e..c7fb99370 100644 --- a/benchmarks/linear_programming/cuopt/run_mip.cpp +++ b/benchmarks/linear_programming/cuopt/run_mip.cpp @@ -207,11 +207,11 @@ int run_single_file(std::string file_path, settings.time_limit = time_limit; settings.heuristics_only = heuristics_only; - settings.num_cpu_threads = num_cpu_threads; + settings.num_cpu_threads = 1; settings.log_to_console = log_to_console; settings.tolerances.relative_tolerance = 1e-12; settings.tolerances.absolute_tolerance = 1e-6; - settings.presolve = true; + settings.presolve = false; cuopt::linear_programming::benchmark_info_t benchmark_info; settings.benchmark_info_ptr = &benchmark_info; auto start_run_solver = std::chrono::high_resolution_clock::now(); diff --git a/cpp/src/dual_simplex/device_sparse_matrix.cuh b/cpp/src/dual_simplex/device_sparse_matrix.cuh index f347f956b..00c198d3f 100644 --- a/cpp/src/dual_simplex/device_sparse_matrix.cuh +++ b/cpp/src/dual_simplex/device_sparse_matrix.cuh @@ -184,7 +184,7 @@ class device_csc_matrix_t { // Inclusive cumulative sum to have the corresponding column for each entry rmm::device_buffer d_temp_storage; - size_t temp_storage_bytes; + size_t temp_storage_bytes{0}; cub::DeviceScan::InclusiveSum( nullptr, temp_storage_bytes, col_index.data(), col_index.data(), col_index.size(), stream); d_temp_storage.resize(temp_storage_bytes, stream); From 2fae64443a46c845a34daf35ffdf70f2a026a0ec Mon Sep 17 00:00:00 2001 From: akifcorduk Date: Mon, 6 Oct 2025 09:54:37 -0700 Subject: [PATCH 3/7] add population log --- benchmarks/linear_programming/cuopt/run_mip.cpp | 4 ++-- cpp/src/mip/diversity/population.cu | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/benchmarks/linear_programming/cuopt/run_mip.cpp b/benchmarks/linear_programming/cuopt/run_mip.cpp index c7fb99370..a4f52cb4e 100644 --- a/benchmarks/linear_programming/cuopt/run_mip.cpp +++ b/benchmarks/linear_programming/cuopt/run_mip.cpp @@ -207,11 +207,11 @@ int run_single_file(std::string file_path, settings.time_limit = time_limit; settings.heuristics_only = heuristics_only; - settings.num_cpu_threads = 1; + settings.num_cpu_threads = num_cpu_threads; settings.log_to_console = log_to_console; settings.tolerances.relative_tolerance = 1e-12; settings.tolerances.absolute_tolerance = 1e-6; - settings.presolve = false; + settings.presolve = true; cuopt::linear_programming::benchmark_info_t benchmark_info; settings.benchmark_info_ptr = &benchmark_info; auto start_run_solver = std::chrono::high_resolution_clock::now(); diff --git a/cpp/src/mip/diversity/population.cu b/cpp/src/mip/diversity/population.cu index 94e8215ef..ae02ff738 100644 --- a/cpp/src/mip/diversity/population.cu +++ b/cpp/src/mip/diversity/population.cu @@ -211,7 +211,7 @@ std::vector> population_t::get_external_solutions sol.copy_new_assignment(h_entry.solution); sol.compute_feasibility(); if (!sol.get_feasible()) { - CUOPT_LOG_ERROR( + CUOPT_LOG_DEBUG( "External solution %d is infeasible, excess %g, obj %g, int viol %g, var viol %g, cstr " "viol %g, n_feasible %d/%d, integers %d/%d", counter, From 09f412b3a43023caf7deb35d10300673230752be Mon Sep 17 00:00:00 2001 From: akifcorduk Date: Tue, 7 Oct 2025 04:19:56 -0700 Subject: [PATCH 4/7] fix invalid ranges on form_col_index --- cpp/src/dual_simplex/device_sparse_matrix.cuh | 18 ++++++++++-------- 1 file changed, 10 insertions(+), 8 deletions(-) diff --git a/cpp/src/dual_simplex/device_sparse_matrix.cuh b/cpp/src/dual_simplex/device_sparse_matrix.cuh index 00c198d3f..8113db973 100644 --- a/cpp/src/dual_simplex/device_sparse_matrix.cuh +++ b/cpp/src/dual_simplex/device_sparse_matrix.cuh @@ -173,14 +173,16 @@ class device_csc_matrix_t { col_index.resize(x.size(), stream); RAFT_CUDA_TRY(cudaMemsetAsync(col_index.data(), 0, sizeof(i_t) * col_index.size(), stream)); // Scatter 1 when there is a col start in col_index - thrust::for_each( - rmm::exec_policy(stream), - thrust::make_counting_iterator(i_t(1)), // Skip the first 0 - thrust::make_counting_iterator(static_cast(col_start.size() - 1)), // Skip the end index - [span_col_start = cuopt::make_span(col_start), - span_col_index = cuopt::make_span(col_index)] __device__(i_t i) { - span_col_index[span_col_start[i]] = 1; - }); + if (col_start.size() > 2) { + thrust::for_each(rmm::exec_policy(stream), + thrust::make_counting_iterator(i_t(1)), // Skip the first 0 + thrust::make_counting_iterator( + static_cast(col_start.size() - 1)), // Skip the end index + [span_col_start = cuopt::make_span(col_start), + span_col_index = cuopt::make_span(col_index)] __device__(i_t i) { + span_col_index[span_col_start[i]] = 1; + }); + } // Inclusive cumulative sum to have the corresponding column for each entry rmm::device_buffer d_temp_storage; From 35e2009c1c1a89efde8c0c781b4a695a947138bf Mon Sep 17 00:00:00 2001 From: "Nicolas L. Guidotti" Date: Tue, 7 Oct 2025 02:31:49 +0200 Subject: [PATCH 5/7] Fixed delayed termination after timeout (#451) This PR fixes the delayed termination of the branch-and-bound algorithm after reaching the time limit. Authors: - Nicolas L. Guidotti (https://github.com/nguidotti) Approvers: - Chris Maes (https://github.com/chris-maes) URL: https://github.com/NVIDIA/cuopt/pull/451 --- cpp/src/dual_simplex/branch_and_bound.cpp | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/cpp/src/dual_simplex/branch_and_bound.cpp b/cpp/src/dual_simplex/branch_and_bound.cpp index 168ffa129..e5a32c33b 100644 --- a/cpp/src/dual_simplex/branch_and_bound.cpp +++ b/cpp/src/dual_simplex/branch_and_bound.cpp @@ -584,6 +584,7 @@ node_status_t branch_and_bound_t::solve_node(search_tree_t& lp_settings.set_log(false); lp_settings.cut_off = upper_bound + settings_.dual_tol; lp_settings.inside_mip = 2; + lp_settings.time_limit = settings_.time_limit - toc(stats_.start_time); // in B&B we only have equality constraints, leave it empty for default std::vector row_sense; @@ -739,7 +740,7 @@ void branch_and_bound_t::exploration_ramp_up(search_tree_t* } } - if (toc(stats_.start_time) > settings_.time_limit) { + if (now > settings_.time_limit) { status_ = mip_exploration_status_t::TIME_LIMIT; return; } @@ -836,9 +837,9 @@ void branch_and_bound_t::explore_subtree(i_t id, } } - if (toc(stats_.start_time) > settings_.time_limit) { + if (now > settings_.time_limit) { status_ = mip_exploration_status_t::TIME_LIMIT; - break; + return; } node_status_t node_status = @@ -846,7 +847,7 @@ void branch_and_bound_t::explore_subtree(i_t id, if (node_status == node_status_t::TIME_LIMIT) { status_ = mip_exploration_status_t::TIME_LIMIT; - break; + return; } else if (node_status == node_status_t::HAS_CHILDREN) { // The stack should only contain the children of the current parent. @@ -972,11 +973,13 @@ void branch_and_bound_t::diving_thread(lp_problem_t& leaf_pr continue; } + if (toc(stats_.start_time) > settings_.time_limit) { return; } + node_status_t node_status = solve_node(subtree, node_ptr, leaf_problem, Arow, upper_bound, log, 'D'); if (node_status == node_status_t::TIME_LIMIT) { - break; + return; } else if (node_status == node_status_t::HAS_CHILDREN) { auto [first, second] = child_selection(node_ptr); From 104877c7eeb3b3df2734ade8c9258c01b79d27c1 Mon Sep 17 00:00:00 2001 From: Alice Boucher Date: Tue, 7 Oct 2025 13:59:18 +0000 Subject: [PATCH 6/7] fix papilo + timeouts --- cpp/CMakeLists.txt | 1 + .../thirdparty/papilo-segfault-timeout.patch | 14 ++++++++ cpp/src/dual_simplex/branch_and_bound.cpp | 2 +- cpp/src/dual_simplex/pinned_host_allocator.cu | 5 +++ cpp/src/mip/local_search/local_search.cu | 32 +++++++++++-------- cpp/src/mip/local_search/local_search.cuh | 5 ++- 6 files changed, 44 insertions(+), 15 deletions(-) create mode 100644 cpp/cmake/thirdparty/papilo-segfault-timeout.patch diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index e03d2581f..c832dcb61 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -194,6 +194,7 @@ FetchContent_Declare( #we can switch to the main branch. GIT_TAG "34a40781fa14f8870cb6368cffb6c0eda2f47511" GIT_PROGRESS TRUE + PATCH_COMMAND git apply ${CMAKE_CURRENT_SOURCE_DIR}/cmake/thirdparty/papilo-segfault-timeout.patch SYSTEM ) diff --git a/cpp/cmake/thirdparty/papilo-segfault-timeout.patch b/cpp/cmake/thirdparty/papilo-segfault-timeout.patch new file mode 100644 index 000000000..8b716df7d --- /dev/null +++ b/cpp/cmake/thirdparty/papilo-segfault-timeout.patch @@ -0,0 +1,14 @@ +diff --git a/src/papilo/core/Presolve.hpp b/src/papilo/core/Presolve.hpp +index 7420744c..c0a53f48 100644 +--- a/src/papilo/core/Presolve.hpp ++++ b/src/papilo/core/Presolve.hpp +@@ -1275,7 +1275,8 @@ Presolve::applyPostponed( ProblemUpdate& probUpdate, const Timer& pr + { + probUpdate.setPostponeSubstitutions( false ); + +- for( int presolver = 0; presolver != (int) presolvers.size(); ++presolver ) ++ // postponedReductionToPresolver may be smaller than presolvers.size() if preempted by a time limit ++ for( int presolver = 0; presolver + 1 < (int)postponedReductionToPresolver.size(); ++presolver ) + { + int first = postponedReductionToPresolver[presolver]; + int last = postponedReductionToPresolver[presolver + 1]; diff --git a/cpp/src/dual_simplex/branch_and_bound.cpp b/cpp/src/dual_simplex/branch_and_bound.cpp index e5a32c33b..cf6fd6979 100644 --- a/cpp/src/dual_simplex/branch_and_bound.cpp +++ b/cpp/src/dual_simplex/branch_and_bound.cpp @@ -1038,7 +1038,6 @@ mip_status_t branch_and_bound_t::solve(mip_solution_t& solut original_lp_, stats_.start_time, lp_settings, root_relax_soln_, root_vstatus_, edge_norms_); stats_.total_lp_iters = root_relax_soln_.iterations; stats_.total_lp_solve_time = toc(stats_.start_time); - assert(root_vstatus_.size() == original_lp_.num_cols); if (root_status == lp_status_t::INFEASIBLE) { settings_.log.printf("MIP Infeasible\n"); // FIXME: rarely dual simplex detects infeasible whereas it is feasible. @@ -1062,6 +1061,7 @@ mip_status_t branch_and_bound_t::solve(mip_solution_t& solut return set_final_solution(solution, -inf); } + assert(root_vstatus_.size() == original_lp_.num_cols); set_uninitialized_steepest_edge_norms(edge_norms_); root_objective_ = compute_objective(original_lp_, root_relax_soln_.x); diff --git a/cpp/src/dual_simplex/pinned_host_allocator.cu b/cpp/src/dual_simplex/pinned_host_allocator.cu index 8821930d3..86eafc2b2 100644 --- a/cpp/src/dual_simplex/pinned_host_allocator.cu +++ b/cpp/src/dual_simplex/pinned_host_allocator.cu @@ -57,6 +57,11 @@ bool operator!=(const PinnedHostAllocator&, const PinnedHostAllocator&) no template class PinnedHostAllocator; template double vector_norm_inf>( const std::vector>& x); + +template bool operator==(const PinnedHostAllocator&, + const PinnedHostAllocator&) noexcept; +template bool operator!=(const PinnedHostAllocator&, + const PinnedHostAllocator&) noexcept; #endif template class PinnedHostAllocator; diff --git a/cpp/src/mip/local_search/local_search.cu b/cpp/src/mip/local_search/local_search.cu index 4b69330d8..ee055b974 100644 --- a/cpp/src/mip/local_search/local_search.cu +++ b/cpp/src/mip/local_search/local_search.cu @@ -225,8 +225,13 @@ void local_search_t::stop_cpufj_scratch_threads() template bool local_search_t::do_fj_solve(solution_t& solution, fj_t& in_fj, + f_t time_limit, const std::string& source) { + if (time_limit == 0.) return solution.get_feasible(); + + timer_t timer(time_limit); + auto h_weights = cuopt::host_copy(in_fj.cstr_weights, solution.handle_ptr->get_stream()); auto h_objective_weight = in_fj.objective_weight.value(solution.handle_ptr->get_stream()); for (auto& cpu_fj : ls_cpu_fj) { @@ -242,7 +247,8 @@ bool local_search_t::do_fj_solve(solution_t& solution, } // Run GPU solver and measure execution time - auto gpu_fj_start = std::chrono::high_resolution_clock::now(); + auto gpu_fj_start = std::chrono::high_resolution_clock::now(); + in_fj.settings.time_limit = timer.remaining_time(); in_fj.solve(solution); // Stop CPU solver @@ -320,9 +326,9 @@ void local_search_t::generate_fast_solution(solution_t& solu constraint_prop.apply_round(solution, 1., constr_prop_timer); if (solution.compute_feasibility()) { return; } if (timer.check_time_limit()) { return; }; - fj.settings.time_limit = std::min(3., timer.remaining_time()); + f_t time_limit = std::min(3., timer.remaining_time()); // run fj on the solution - do_fj_solve(solution, fj, "fast"); + do_fj_solve(solution, fj, time_limit, "fast"); // TODO check if FJ returns the same solution // check if the solution is feasible if (solution.compute_feasibility()) { return; } @@ -381,11 +387,11 @@ bool local_search_t::run_fj_until_timer(solution_t& solution bool is_feasible; fj.settings.n_of_minimums_for_exit = 1e6; fj.settings.mode = fj_mode_t::EXIT_NON_IMPROVING; - fj.settings.time_limit = timer.remaining_time() * 0.95; fj.settings.update_weights = false; fj.settings.feasibility_run = false; fj.copy_weights(weights, solution.handle_ptr); - do_fj_solve(solution, fj, "until_timer"); + f_t time_limit = timer.remaining_time() * 0.95; + do_fj_solve(solution, fj, time_limit, "until_timer"); CUOPT_LOG_DEBUG("Initial FJ feasibility done"); is_feasible = solution.compute_feasibility(); if (fj.settings.feasibility_run || timer.check_time_limit()) { return is_feasible; } @@ -410,11 +416,11 @@ bool local_search_t::run_fj_annealing(solution_t& solution, fj.settings.mode = fj_mode_t::EXIT_NON_IMPROVING; fj.settings.candidate_selection = fj_candidate_selection_t::FEASIBLE_FIRST; fj.settings.iteration_limit = ls_config.iteration_limit; - fj.settings.time_limit = std::min(10., timer.remaining_time()); fj.settings.parameters.allow_infeasibility_iterations = 100; fj.settings.update_weights = 1; fj.settings.baseline_objective_for_longer_run = ls_config.best_objective_of_parents; - do_fj_solve(solution, fj, "annealing"); + f_t time_limit = std::min(10., timer.remaining_time()); + do_fj_solve(solution, fj, time_limit, "annealing"); bool is_feasible = solution.compute_feasibility(); fj.settings = prev_settings; @@ -460,14 +466,14 @@ bool local_search_t::check_fj_on_lp_optimal(solution_t& solu solution.assign_random_within_bounds(perturbation_ratio); } cuopt_func_call(solution.test_variable_bounds(false)); - f_t lp_run_time_after_feasible = 1.; + f_t lp_run_time_after_feasible = std::min(1., timer.remaining_time()); timer_t bounds_prop_timer = timer_t(std::min(timer.remaining_time(), 10.)); bool is_feasible = constraint_prop.apply_round(solution, lp_run_time_after_feasible, bounds_prop_timer); if (!is_feasible) { const f_t lp_run_time = 2.; relaxed_lp_settings_t lp_settings; - lp_settings.time_limit = lp_run_time; + lp_settings.time_limit = std::min(lp_run_time, timer.remaining_time()); lp_settings.tolerance = solution.problem_ptr->tolerances.absolute_tolerance; run_lp_with_vars_fixed( *solution.problem_ptr, solution, solution.problem_ptr->integer_indices, lp_settings); @@ -479,8 +485,8 @@ bool local_search_t::check_fj_on_lp_optimal(solution_t& solu fj.settings.n_of_minimums_for_exit = 20000; fj.settings.update_weights = true; fj.settings.feasibility_run = false; - fj.settings.time_limit = std::min(30., timer.remaining_time()); - do_fj_solve(solution, fj, "on_lp_optimal"); + f_t time_limit = std::min(30., timer.remaining_time()); + do_fj_solve(solution, fj, time_limit, "on_lp_optimal"); return solution.get_feasible(); } @@ -496,8 +502,8 @@ bool local_search_t::run_fj_on_zero(solution_t& solution, ti fj.settings.n_of_minimums_for_exit = 20000; fj.settings.update_weights = true; fj.settings.feasibility_run = false; - fj.settings.time_limit = std::min(30., timer.remaining_time()); - bool is_feasible = do_fj_solve(solution, fj, "on_zero"); + f_t time_limit = std::min(30., timer.remaining_time()); + bool is_feasible = do_fj_solve(solution, fj, time_limit, "on_zero"); return is_feasible; } diff --git a/cpp/src/mip/local_search/local_search.cuh b/cpp/src/mip/local_search/local_search.cuh index e8dc016e2..1219a7d3f 100644 --- a/cpp/src/mip/local_search/local_search.cuh +++ b/cpp/src/mip/local_search/local_search.cuh @@ -113,7 +113,10 @@ class local_search_t { population_t* population_ptr = nullptr); void resize_vectors(problem_t& problem, const raft::handle_t* handle_ptr); - bool do_fj_solve(solution_t& solution, fj_t& fj, const std::string& source); + bool do_fj_solve(solution_t& solution, + fj_t& fj, + f_t time_limit, + const std::string& source); i_t ls_threads() const { return ls_cpu_fj.size() + scratch_cpu_fj.size(); } void save_solution_and_add_cutting_plane(solution_t& solution, From 36c808edc0338f2a789bf5f4a47ffeff36e89b7c Mon Sep 17 00:00:00 2001 From: Alice Boucher Date: Tue, 7 Oct 2025 14:25:06 +0000 Subject: [PATCH 7/7] remove papilo patch in favor of min 1s timeout --- cpp/CMakeLists.txt | 1 - cpp/cmake/thirdparty/papilo-segfault-timeout.patch | 14 -------------- cpp/src/linear_programming/solve.cu | 4 +++- 3 files changed, 3 insertions(+), 16 deletions(-) delete mode 100644 cpp/cmake/thirdparty/papilo-segfault-timeout.patch diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index c832dcb61..e03d2581f 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -194,7 +194,6 @@ FetchContent_Declare( #we can switch to the main branch. GIT_TAG "34a40781fa14f8870cb6368cffb6c0eda2f47511" GIT_PROGRESS TRUE - PATCH_COMMAND git apply ${CMAKE_CURRENT_SOURCE_DIR}/cmake/thirdparty/papilo-segfault-timeout.patch SYSTEM ) diff --git a/cpp/cmake/thirdparty/papilo-segfault-timeout.patch b/cpp/cmake/thirdparty/papilo-segfault-timeout.patch deleted file mode 100644 index 8b716df7d..000000000 --- a/cpp/cmake/thirdparty/papilo-segfault-timeout.patch +++ /dev/null @@ -1,14 +0,0 @@ -diff --git a/src/papilo/core/Presolve.hpp b/src/papilo/core/Presolve.hpp -index 7420744c..c0a53f48 100644 ---- a/src/papilo/core/Presolve.hpp -+++ b/src/papilo/core/Presolve.hpp -@@ -1275,7 +1275,8 @@ Presolve::applyPostponed( ProblemUpdate& probUpdate, const Timer& pr - { - probUpdate.setPostponeSubstitutions( false ); - -- for( int presolver = 0; presolver != (int) presolvers.size(); ++presolver ) -+ // postponedReductionToPresolver may be smaller than presolvers.size() if preempted by a time limit -+ for( int presolver = 0; presolver + 1 < (int)postponedReductionToPresolver.size(); ++presolver ) - { - int first = postponedReductionToPresolver[presolver]; - int last = postponedReductionToPresolver[presolver + 1]; diff --git a/cpp/src/linear_programming/solve.cu b/cpp/src/linear_programming/solve.cu index 089b5cb9a..22ec49986 100644 --- a/cpp/src/linear_programming/solve.cu +++ b/cpp/src/linear_programming/solve.cu @@ -821,7 +821,9 @@ optimization_problem_solution_t solve_lp(optimization_problem_t>(); auto [reduced_problem, feasible] = presolver->apply(op_problem,