From d7a7b0729ed5d98c76e6ed944e69cfa8549ccbc7 Mon Sep 17 00:00:00 2001 From: Trevor McKay Date: Wed, 5 Aug 2026 15:29:12 -0400 Subject: [PATCH] Add cooperative cancel for LP/MIP solves Poll a shared cancel flag like time limits so mid-solve cancel can unwind and avoid as much as possible the need for SIGKILL especially in the gRPC server, since SIGKILL can leave the GPU in an uninterruptible sleep in some cases. Remap limit statuses to Cancelled at solution finalization, preempt MIP heuristics on cancel. Signed-off-by: Trevor McKay --- .../mathematical_optimization/constants.h | 1 + .../mip/solver_settings.hpp | 5 + .../mip/solver_solution.hpp | 2 + .../pdlp/solver_settings.hpp | 3 + .../pdlp/solver_solution.hpp | 3 +- cpp/src/barrier/barrier.cu | 113 +- cpp/src/barrier/sparse_cholesky.cuh | 34 +- cpp/src/branch_and_bound/branch_and_bound.cpp | 32 +- cpp/src/branch_and_bound/branch_and_bound.hpp | 4 +- cpp/src/dual_simplex/crossover.cpp | 25 +- cpp/src/dual_simplex/phase2.cpp | 71 +- cpp/src/dual_simplex/primal.cpp | 16 +- cpp/src/dual_simplex/right_looking_lu.cpp | 21 +- .../dual_simplex/simplex_solver_settings.hpp | 27 + cpp/src/dual_simplex/solve.cpp | 44 +- cpp/src/dual_simplex/solve.hpp | 4 +- cpp/src/grpc/server/grpc_job_management.cpp | 64 +- cpp/src/grpc/server/grpc_worker.cpp | 46 +- .../diversity/diversity_manager.cu | 20 +- .../recombiners/bound_prop_recombiner.cuh | 6 +- .../diversity/recombiners/fp_recombiner.cuh | 3 +- .../feasibility_jump/early_cpufj.cu | 18 +- .../feasibility_jump/early_cpufj.cuh | 4 +- .../feasibility_jump/feasibility_jump.cu | 4 +- .../feasibility_pump/feasibility_pump.cu | 3 +- .../local_search/local_search.cu | 27 +- .../presolve/bounds_presolve.cu | 6 +- .../presolve/conflict_graph/clique_table.cu | 5 +- .../presolve/load_balanced_bounds_presolve.cu | 6 +- .../mip_heuristics/presolve/multi_probe.cu | 4 +- .../presolve/third_party_presolve.cpp | 15 +- .../mip_heuristics/relaxed_lp/relaxed_lp.cu | 3 +- .../mip_heuristics/relaxed_lp/relaxed_lp.cuh | 17 +- cpp/src/mip_heuristics/solution/solution.cu | 26 +- cpp/src/mip_heuristics/solution/solution.cuh | 4 + cpp/src/mip_heuristics/solve.cu | 34 +- cpp/src/mip_heuristics/solver.cu | 23 +- cpp/src/mip_heuristics/solver_solution.cu | 13 +- cpp/src/pdlp/pdlp.cu | 13 +- cpp/src/pdlp/solve.cu | 47 +- cpp/src/pdlp/solver_solution.cu | 1 + cpp/src/utilities/solve_limits.hpp | 193 ++ cpp/src/utilities/timer.hpp | 25 +- cpp/tests/mip/CMakeLists.txt | 3 + cpp/tests/mip/cooperative_cancel_test.cu | 188 ++ .../tests/fixtures/grpc_server_fixtures.py | 30 +- .../grpc_lifecycle_stress.py | 1790 +++++++++++++++++ .../linear_programming/test_grpc_client.py | 144 +- 48 files changed, 2939 insertions(+), 251 deletions(-) create mode 100644 cpp/src/utilities/solve_limits.hpp create mode 100644 cpp/tests/mip/cooperative_cancel_test.cu create mode 100644 python/cuopt/cuopt/tests/linear_programming/grpc_lifecycle_stress.py diff --git a/cpp/include/cuopt/mathematical_optimization/constants.h b/cpp/include/cuopt/mathematical_optimization/constants.h index 86ed6965c8..fca3a7dbd1 100644 --- a/cpp/include/cuopt/mathematical_optimization/constants.h +++ b/cpp/include/cuopt/mathematical_optimization/constants.h @@ -164,6 +164,7 @@ #define CUOPT_TERMINATION_STATUS_CONCURRENT_LIMIT 9 #define CUOPT_TERMINATION_STATUS_WORK_LIMIT 10 #define CUOPT_TERMINATION_STATUS_UNBOUNDED_OR_INFEASIBLE 11 +#define CUOPT_TERMINATION_STATUS_CANCELLED 12 /* @brief The objective sense constants */ #define CUOPT_MINIMIZE 1 diff --git a/cpp/include/cuopt/mathematical_optimization/mip/solver_settings.hpp b/cpp/include/cuopt/mathematical_optimization/mip/solver_settings.hpp index 8bb676cb94..d64aab655b 100644 --- a/cpp/include/cuopt/mathematical_optimization/mip/solver_settings.hpp +++ b/cpp/include/cuopt/mathematical_optimization/mip/solver_settings.hpp @@ -7,6 +7,7 @@ #pragma once +#include #include #include @@ -183,6 +184,10 @@ class mip_solver_settings_t { // benchmarks benchmark_info_t* benchmark_info_ptr = nullptr; + // Optional cooperative cancel (level-triggered). Non-owning. Cleared by + // solve_mip on return. Set true from another thread to request early exit. + std::atomic* cancel_requested{nullptr}; + // TODO check with Akif and Alice pdlp::pdlp_hyper_params_t hyper_params; diff --git a/cpp/include/cuopt/mathematical_optimization/mip/solver_solution.hpp b/cpp/include/cuopt/mathematical_optimization/mip/solver_solution.hpp index 1ad58b9e10..0b12134636 100644 --- a/cpp/include/cuopt/mathematical_optimization/mip/solver_solution.hpp +++ b/cpp/include/cuopt/mathematical_optimization/mip/solver_solution.hpp @@ -34,6 +34,7 @@ enum class mip_termination_status_t : int8_t { TimeLimit = CUOPT_TERMINATION_STATUS_TIME_LIMIT, WorkLimit = CUOPT_TERMINATION_STATUS_WORK_LIMIT, UnboundedOrInfeasible = CUOPT_TERMINATION_STATUS_UNBOUNDED_OR_INFEASIBLE, + Cancelled = CUOPT_TERMINATION_STATUS_CANCELLED, }; template @@ -65,6 +66,7 @@ class mip_solution_t : public base_solution_t { double get_total_solve_time() const; double get_presolve_time() const; mip_termination_status_t get_termination_status() const; + void set_termination_status(mip_termination_status_t termination_status); static std::string get_termination_status_string(mip_termination_status_t termination_status); std::string get_termination_status_string() const; const cuopt::logic_error& get_error_status() const; diff --git a/cpp/include/cuopt/mathematical_optimization/pdlp/solver_settings.hpp b/cpp/include/cuopt/mathematical_optimization/pdlp/solver_settings.hpp index ffcf3fad7a..abf904e462 100644 --- a/cpp/include/cuopt/mathematical_optimization/pdlp/solver_settings.hpp +++ b/cpp/include/cuopt/mathematical_optimization/pdlp/solver_settings.hpp @@ -340,6 +340,9 @@ class pdlp_solver_settings_t { bool inside_mip{false}; // For concurrent termination std::atomic* concurrent_halt{nullptr}; + // Optional cooperative cancel (level-triggered). Non-owning. Cleared by + // solve_lp on return. Set true from another thread to request early exit. + std::atomic* cancel_requested{nullptr}; // Shared strong branching solved flags for cooperative DS + PDLP cuda::std::span> shared_sb_solved; static constexpr f_t minimal_absolute_tolerance = 1.0e-12; diff --git a/cpp/include/cuopt/mathematical_optimization/pdlp/solver_solution.hpp b/cpp/include/cuopt/mathematical_optimization/pdlp/solver_solution.hpp index be48ea4baf..377f2d6159 100644 --- a/cpp/include/cuopt/mathematical_optimization/pdlp/solver_solution.hpp +++ b/cpp/include/cuopt/mathematical_optimization/pdlp/solver_solution.hpp @@ -37,7 +37,8 @@ enum class pdlp_termination_status_t : int8_t { TimeLimit = CUOPT_TERMINATION_STATUS_TIME_LIMIT, PrimalFeasible = CUOPT_TERMINATION_STATUS_PRIMAL_FEASIBLE, ConcurrentLimit = CUOPT_TERMINATION_STATUS_CONCURRENT_LIMIT, - UnboundedOrInfeasible = CUOPT_TERMINATION_STATUS_UNBOUNDED_OR_INFEASIBLE + UnboundedOrInfeasible = CUOPT_TERMINATION_STATUS_UNBOUNDED_OR_INFEASIBLE, + Cancelled = CUOPT_TERMINATION_STATUS_CANCELLED }; /** diff --git a/cpp/src/barrier/barrier.cu b/cpp/src/barrier/barrier.cu index 5db791f7e9..fd9b8b3919 100644 --- a/cpp/src/barrier/barrier.cu +++ b/cpp/src/barrier/barrier.cu @@ -500,7 +500,7 @@ class iteration_data_t { find_dense_columns( lp.A, settings, dense_columns_unordered, n_dense_rows, max_row_nz, estimated_nz_AAT); } - if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { return; } + if (settings.cancel_or_halt_requested()) { return; } #ifdef PRINT_INFO for (i_t j : dense_columns_unordered) { settings.log.printf("Dense column %6d\n", j); @@ -576,7 +576,7 @@ class iteration_data_t { if (n_upper_bounds > 0 || (has_Q && !use_augmented)) { inv_diag.sqrt(inv_sqrt_diag); } } - if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { return; } + if (settings.cancel_or_halt_requested()) { return; } { raft::common::nvtx::range scope("Barrier: LP Data: AD matrix setup"); @@ -642,7 +642,7 @@ class iteration_data_t { RAFT_CHECK_CUDA(handle_ptr->get_stream()); } - if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { return; } + if (settings.cancel_or_halt_requested()) { return; } { raft::common::nvtx::range scope("Barrier: LP Data: Cholesky init"); i_t factorization_size = @@ -651,7 +651,7 @@ class iteration_data_t { handle_ptr, settings, factorization_size); chol->set_positive_definite(false); } - if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { return; } + if (settings.cancel_or_halt_requested()) { return; } { raft::common::nvtx::range scope("Barrier: LP Data: symbolic analysis"); // Perform symbolic analysis @@ -662,14 +662,14 @@ class iteration_data_t { // Build the sparsity pattern of the augmented system form_augmented(true); } - if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { return; } + if (settings.cancel_or_halt_requested()) { return; } symbolic_status = chol->analyze(device_augmented); } else { { raft::common::nvtx::range form_scope("Barrier: LP Data: form ADAT"); form_adat(true); } - if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { return; } + if (settings.cancel_or_halt_requested()) { return; } symbolic_status = chol->analyze(device_ADAT); } } @@ -918,7 +918,7 @@ class iteration_data_t { }); RAFT_CHECK_CUDA(stream_view_); } - if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { return; } + if (settings_.cancel_or_halt_requested()) { return; } if (first_call) { raft::common::nvtx::range scope("Barrier: Form ADAT: cusparse init"); try { @@ -929,7 +929,7 @@ class iteration_data_t { return; } } - if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { return; } + if (settings_.cancel_or_halt_requested()) { return; } { raft::common::nvtx::range scope("Barrier: Form ADAT: ADAT multiply"); @@ -1023,9 +1023,7 @@ class iteration_data_t { dense_vector_t M_col(AD.m); solve_status = chol->solve(U_col, M_col); if (solve_status != 0) { return solve_status; } - if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { - return CONCURRENT_HALT_RETURN; - } + if (settings_.cancel_or_halt_requested()) { return CONCURRENT_HALT_RETURN; } M.set_column(k, M_col); if (debug) { @@ -1042,9 +1040,7 @@ class iteration_data_t { for (i_t k = 0; k < n_dense_columns; k++) { AD_dense.transpose_multiply( 1.0, M.values.data() + k * M.m, 0.0, H.values.data() + k * H.m); - if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { - return CONCURRENT_HALT_RETURN; - } + if (settings_.cancel_or_halt_requested()) { return CONCURRENT_HALT_RETURN; } } dense_vector_t e(n_dense_columns); @@ -1442,7 +1438,7 @@ class iteration_data_t { std::sort(column_nz_permutation.begin(), column_nz_permutation.end(), [&column_nz](i_t i, i_t j) { return column_nz[i] < column_nz[j]; }); - if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { return; } + if (settings.cancel_or_halt_requested()) { return; } // We then compute the exact sparsity pattern for columns of A whose where // the number of nonzeros is less than a threshold. This part can be done @@ -1473,7 +1469,7 @@ class iteration_data_t { // The best way to do that is to have A stored in CSR format. csr_matrix_t A_row(0, 0, 0); A.to_compressed_row(A_row); - if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { return; } + if (settings.cancel_or_halt_requested()) { return; } std::vector histogram(m + 1, 0); for (i_t j = 0; j < n; j++) { @@ -1545,7 +1541,7 @@ class iteration_data_t { delta_nz[j] += fill; // Capture contributions from A(:, j). j will be encountered multiple times } - if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { return; } + if (settings.cancel_or_halt_requested()) { return; } } int64_t sparse_nz_C = 0; @@ -1585,7 +1581,7 @@ class iteration_data_t { delta_nz[j] + static_cast( fill_estimate)); // Capture the estimated fill associated with column j } - if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { return; } + if (settings.cancel_or_halt_requested()) { return; } } int64_t estimated_nz_C = 0; @@ -1603,7 +1599,7 @@ class iteration_data_t { std::sort(permutation.begin(), permutation.end(), [&delta_nz](i_t i, i_t j) { return delta_nz[i] < delta_nz[j]; }); - if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { return; } + if (settings.cancel_or_halt_requested()) { return; } // Now we make a forward pass and compute the number of nonzeros in C // assuming we had included column j @@ -2786,9 +2782,7 @@ i_t barrier_solver_t::gpu_compute_search_direction(iteration_data_tfactorize(data.device_augmented); @@ -2806,9 +2800,7 @@ i_t barrier_solver_t::gpu_compute_search_direction(iteration_data_tfactorize(data.device_ADAT); @@ -4102,6 +4094,21 @@ lp_status_t barrier_solver_t::solve(f_t start_time, lp_solution_t std::optional { + const auto limit = settings.check_solve_limits(start_time); + if (limit == cuopt::solve_limit_reason_t::None) { return std::nullopt; } + if (limit == cuopt::solve_limit_reason_t::Cancelled) { + settings.log.printf("Barrier solve cancelled\n"); + return lp_status_t::CANCELLED; + } + if (limit == cuopt::solve_limit_reason_t::ConcurrentHalt) { + settings.log.printf("Barrier solver halted\n"); + return lp_status_t::CONCURRENT_LIMIT; + } + settings.log.printf("Barrier time limit exceeded\n"); + return lp_status_t::TIME_LIMIT; + }; + i_t n = lp.num_cols; i_t m = lp.num_rows; @@ -4136,10 +4143,7 @@ lp_status_t barrier_solver_t::solve(f_t start_time, lp_solution_t data( lp, num_upper_bounds, presolve_info.direct_free_variables, Q, settings); - if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { - settings.log.printf("Barrier solver halted\n"); - return lp_status_t::CONCURRENT_LIMIT; - } + if (auto st = barrier_limit_status()) { return *st; } if (data.indefinite_Q) { return lp_status_t::NUMERICAL_ISSUES; } if (data.symbolic_status != 0) { settings.log.printf("Error in symbolic analysis\n"); @@ -4155,20 +4159,10 @@ lp_status_t barrier_solver_t::solve(f_t start_time, lp_solution_t settings.time_limit) { - settings.log.printf("Barrier time limit exceeded\n"); - return lp_status_t::TIME_LIMIT; - } + if (auto st = barrier_limit_status()) { return *st; } i_t initial_status = initial_point(data); - if (toc(start_time) > settings.time_limit) { - settings.log.printf("Barrier time limit exceeded\n"); - return lp_status_t::TIME_LIMIT; - } - if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { - settings.log.printf("Barrier solver halted\n"); - return lp_status_t::CONCURRENT_LIMIT; - } + if (auto st = barrier_limit_status()) { return *st; } if (initial_status != 0) { settings.log.printf("Unable to compute initial point\n"); return lp_status_t::NUMERICAL_ISSUES; @@ -4274,14 +4268,7 @@ lp_status_t barrier_solver_t::solve(f_t start_time, lp_solution_t settings.time_limit) { - settings.log.printf("Barrier time limit exceeded\n"); - return lp_status_t::TIME_LIMIT; - } - if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { - settings.log.printf("Barrier solver halted\n"); - return lp_status_t::CONCURRENT_LIMIT; - } + if (auto st = barrier_limit_status()) { return *st; } // Compute the affine step. This is the call that (re)factorizes the // augmented system, so the IR residual here drives the adaptation of @@ -4295,10 +4282,7 @@ lp_status_t barrier_solver_t::solve(f_t start_time, lp_solution_t::solve(f_t start_time, lp_solution_t settings.time_limit) { - settings.log.printf("Barrier time limit exceeded\n"); - return lp_status_t::TIME_LIMIT; - } - if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { - settings.log.printf("Barrier solver halted\n"); - return lp_status_t::CONCURRENT_LIMIT; - } + if (auto st = barrier_limit_status()) { return *st; } f_t mu_aff, sigma, new_mu; compute_target_mu(data, mu, mu_aff, sigma, new_mu); @@ -4336,10 +4313,7 @@ lp_status_t barrier_solver_t::solve(f_t start_time, lp_solution_t::solve(f_t start_time, lp_solution_t settings.time_limit) { - settings.log.printf("Barrier time limit exceeded\n"); - return lp_status_t::TIME_LIMIT; - } - if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { - settings.log.printf("Barrier solver halted\n"); - return lp_status_t::CONCURRENT_LIMIT; - } + if (auto st = barrier_limit_status()) { return *st; } compute_final_direction(data); f_t step_primal, step_dual; diff --git a/cpp/src/barrier/sparse_cholesky.cuh b/cpp/src/barrier/sparse_cholesky.cuh index 3d88fef2ce..e5823894f9 100644 --- a/cpp/src/barrier/sparse_cholesky.cuh +++ b/cpp/src/barrier/sparse_cholesky.cuh @@ -482,9 +482,7 @@ class sparse_cholesky_cudss_t : public sparse_cholesky_base_t { raft::common::nvtx::range fun_scope("Barrier: cuDSS Analyze : CUDSS_PHASE_ANALYSIS"); status = cudssExecute(handle, CUDSS_PHASE_REORDERING, solverConfig, solverData, A, cudss_x, cudss_b); - if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { - return CONCURRENT_HALT_RETURN; - } + if (settings_.cancel_or_halt_requested()) { return CONCURRENT_HALT_RETURN; } if (status != CUDSS_STATUS_SUCCESS) { settings_.log.printf( "FAILED: CUDSS call ended unsuccessfully with status = %d, details: cuDSSExecute for " @@ -498,9 +496,7 @@ class sparse_cholesky_cudss_t : public sparse_cholesky_base_t { status = cudssExecute( handle, CUDSS_PHASE_SYMBOLIC_FACTORIZATION, solverConfig, solverData, A, cudss_x, cudss_b); - if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { - return CONCURRENT_HALT_RETURN; - } + if (settings_.cancel_or_halt_requested()) { return CONCURRENT_HALT_RETURN; } if (status != CUDSS_STATUS_SUCCESS) { settings_.log.printf( "FAILED: CUDSS call ended unsuccessfully with status = %d, details: cuDSSExecute for " @@ -556,9 +552,7 @@ class sparse_cholesky_cudss_t : public sparse_cholesky_base_t { f_t start_numeric = tic(); status = cudssExecute( handle, CUDSS_PHASE_FACTORIZATION, solverConfig, solverData, A, cudss_x, cudss_b); - if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { - return CONCURRENT_HALT_RETURN; - } + if (settings_.cancel_or_halt_requested()) { return CONCURRENT_HALT_RETURN; } if (status != CUDSS_STATUS_SUCCESS) { settings_.log.printf( "FAILED: CUDSS call ended unsuccessfully with status = %d, details: cuDSSExecute for " @@ -572,9 +566,7 @@ class sparse_cholesky_cudss_t : public sparse_cholesky_base_t { #endif f_t numeric_time = toc(start_numeric); - if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { - return CONCURRENT_HALT_RETURN; - } + if (settings_.cancel_or_halt_requested()) { return CONCURRENT_HALT_RETURN; } int info; size_t sizeWritten = 0; @@ -691,9 +683,7 @@ class sparse_cholesky_cudss_t : public sparse_cholesky_base_t { A_created = true; // Perform symbolic analysis - if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { - return CONCURRENT_HALT_RETURN; - } + if (settings_.cancel_or_halt_requested()) { return CONCURRENT_HALT_RETURN; } f_t start_analysis = tic(); CUDSS_CALL_AND_CHECK( cudssExecute(handle, CUDSS_PHASE_REORDERING, solverConfig, solverData, A, cudss_x, cudss_b), @@ -701,9 +691,7 @@ class sparse_cholesky_cudss_t : public sparse_cholesky_base_t { "cudssExecute for reordering"); f_t reorder_time = toc(start_analysis); - if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { - return CONCURRENT_HALT_RETURN; - } + if (settings_.cancel_or_halt_requested()) { return CONCURRENT_HALT_RETURN; } f_t start_symbolic = tic(); @@ -716,7 +704,7 @@ class sparse_cholesky_cudss_t : public sparse_cholesky_base_t { f_t symbolic_time = toc(start_symbolic); f_t analysis_time = toc(start_analysis); settings_.log.printf("Symbolic factorization time : %.2fs\n", symbolic_time); - if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { + if (settings_.cancel_or_halt_requested()) { RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); handle_ptr_->get_stream().synchronize(); return CONCURRENT_HALT_RETURN; @@ -767,9 +755,7 @@ class sparse_cholesky_cudss_t : public sparse_cholesky_base_t { "cudssExecute for factorization"); f_t numeric_time = toc(start_numeric); - if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { - return CONCURRENT_HALT_RETURN; - } + if (settings_.cancel_or_halt_requested()) { return CONCURRENT_HALT_RETURN; } int info; size_t sizeWritten = 0; @@ -832,9 +818,7 @@ class sparse_cholesky_cudss_t : public sparse_cholesky_base_t { cudssMatrixSetValues(cudss_x, x.data()), status, "cudssMatrixSetValues for x"); status = cudssExecute(handle, CUDSS_PHASE_SOLVE, solverConfig, solverData, A, cudss_x, cudss_b); - if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { - return CONCURRENT_HALT_RETURN; - } + if (settings_.cancel_or_halt_requested()) { return CONCURRENT_HALT_RETURN; } if (status != CUDSS_STATUS_SUCCESS) { settings_.log.printf( "FAILED: CUDSS call ended unsuccessfully with status = %d, details: cuDSSExecute for " diff --git a/cpp/src/branch_and_bound/branch_and_bound.cpp b/cpp/src/branch_and_bound/branch_and_bound.cpp index 4dc6bc67a8..4b8f12d350 100644 --- a/cpp/src/branch_and_bound/branch_and_bound.cpp +++ b/cpp/src/branch_and_bound/branch_and_bound.cpp @@ -33,6 +33,7 @@ #include #include #include +#include #include @@ -826,6 +827,10 @@ void branch_and_bound_t::set_final_solution(mip_solution_t& settings_.log.printf("Time limit reached. Stopping the solver...\n"); } + if (solver_status_ == mip_status_t::CANCELLED) { + settings_.log.printf("Solve cancelled. Stopping the solver...\n"); + } + if (solver_status_ == mip_status_t::WORK_LIMIT) { settings_.log.printf("Work limit reached. Stopping the solver...\n"); } @@ -1733,6 +1738,14 @@ void branch_and_bound_t::plunge_with(bfs_worker_t* worker, break; } + if (settings_.cancel_requested != nullptr && + settings_.cancel_requested->load(std::memory_order_acquire)) { + solver_status_ = mip_status_t::CANCELLED; + stack.push_front(node_ptr); + --exploration_stats_.nodes_being_solved; + break; + } + if (exploration_stats_.nodes_explored + exploration_stats_.nodes_being_solved > settings_.node_limit) { solver_status_ = mip_status_t::NODE_LIMIT; @@ -2800,6 +2813,15 @@ lp_status_t branch_and_bound_t::solve_root_relaxation( // Wait for the root relaxation solution to be sent by the diversity manager or dual simplex while (!root_crossover_solution_set_.load(std::memory_order_acquire) && *get_root_concurrent_halt() == 0) { + // Dual may early-return on time/cancel without crossover; also poll limits + // here so this sleep loop cannot outlive the solve budget. + if (cuopt::solve_limit_reached(toc(exploration_stats_.start_time), + settings_.time_limit, + settings_.cancel_requested, + nullptr)) { + set_root_concurrent_halt(1); + break; + } std::this_thread::sleep_for(std::chrono::milliseconds(1)); #pragma omp taskyield } @@ -3340,9 +3362,15 @@ mip_status_t branch_and_bound_t::solve(mip_solution_t& solut return mip_status_t::UNBOUNDED; } - if (root_status == lp_status_t::TIME_LIMIT) { + // Concurrent root: dual may return CONCURRENT_LIMIT when the waiter raises + // concurrent_halt on cancel/TL (no crossover). Treat like TIME_LIMIT/CANCELLED. + if (root_status == lp_status_t::TIME_LIMIT || root_status == lp_status_t::CANCELLED || + root_status == lp_status_t::CONCURRENT_LIMIT) { settings_.log.printf("\n"); - solver_status_ = mip_status_t::TIME_LIMIT; + solver_status_ = + (root_status == lp_status_t::CANCELLED || cuopt::cancel_flag_set(settings_.cancel_requested)) + ? mip_status_t::CANCELLED + : mip_status_t::TIME_LIMIT; set_final_solution(solution, -inf); signal_extend_cliques_.store(true, std::memory_order_release); #pragma omp taskwait depend(in : *clique_signal) diff --git a/cpp/src/branch_and_bound/branch_and_bound.hpp b/cpp/src/branch_and_bound/branch_and_bound.hpp index 96b8a6d8fe..ab247bf932 100644 --- a/cpp/src/branch_and_bound/branch_and_bound.hpp +++ b/cpp/src/branch_and_bound/branch_and_bound.hpp @@ -55,7 +55,8 @@ enum class mip_status_t { NUMERICAL = 6, // The solver encountered a numerical error UNSET = 7, // The status is not set WORK_LIMIT = 8, // The solver reached a deterministic work limit - SUBMIP_HALT = 9 // Halt the solver + SUBMIP_HALT = 9, // Halt the solver + CANCELLED = 10 // Cooperative cancel requested }; inline std::string mip_status_to_string(mip_status_t status) @@ -71,6 +72,7 @@ inline std::string mip_status_to_string(mip_status_t status) case mip_status_t::UNSET: return "UNSET"; case mip_status_t::WORK_LIMIT: return "WORK_LIMIT"; case mip_status_t::SUBMIP_HALT: return "SUBMIP_HALT"; + case mip_status_t::CANCELLED: return "CANCELLED"; } return "UNKNOWN"; } diff --git a/cpp/src/dual_simplex/crossover.cpp b/cpp/src/dual_simplex/crossover.cpp index e1ba272adf..e552f41303 100644 --- a/cpp/src/dual_simplex/crossover.cpp +++ b/cpp/src/dual_simplex/crossover.cpp @@ -607,7 +607,10 @@ i_t dual_push(const lp_problem_t& lp, settings.log.printf( "%d of %d dual pushes in %.2fs\n", num_pushes, total_superbasics, toc(start_time)); } - if (toc(start_time) > settings.time_limit) { + if (cuopt::solve_limit_reached(cuopt::mathematical_optimization::toc(start_time), + settings.time_limit, + settings.cancel_requested, + nullptr)) { settings.log.printf("Crossover time exceeded\n"); return TIME_LIMIT_RETURN; } @@ -984,7 +987,10 @@ i_t primal_push(const lp_problem_t& lp, last_print_time = tic(); } - if (toc(start_time) > settings.time_limit) { + if (cuopt::solve_limit_reached(cuopt::mathematical_optimization::toc(start_time), + settings.time_limit, + settings.cancel_requested, + nullptr)) { settings.log.printf("Crossover time limit exceeded\n"); return TIME_LIMIT_RETURN; } @@ -1363,7 +1369,10 @@ crossover_status_t crossover(const lp_problem_t& lp, } reorder_basic_list(q, basic_list); - if (toc(start_time) > settings.time_limit) { + if (cuopt::solve_limit_reached(cuopt::mathematical_optimization::toc(start_time), + settings.time_limit, + settings.cancel_requested, + nullptr)) { settings.log.printf("Time limit exceeded\n"); return crossover_status_t::TIME_LIMIT; } @@ -1425,7 +1434,10 @@ crossover_status_t crossover(const lp_problem_t& lp, std::vector edge_norms; dual_status_t status = dual_phase2(2, 0, start_time, lp, settings, vstatus, solution, dual_iter, edge_norms); - if (toc(start_time) > settings.time_limit) { + if (cuopt::solve_limit_reached(cuopt::mathematical_optimization::toc(start_time), + settings.time_limit, + settings.cancel_requested, + nullptr)) { settings.log.printf("Time limit exceeded\n"); return crossover_status_t::TIME_LIMIT; } @@ -1587,7 +1599,10 @@ crossover_status_t crossover(const lp_problem_t& lp, std::vector edge_norms; status = dual_phase2( 2, iter == 0 ? 1 : 0, start_time, lp, settings, vstatus, solution, iter, edge_norms); - if (toc(start_time) > settings.time_limit) { + if (cuopt::solve_limit_reached(cuopt::mathematical_optimization::toc(start_time), + settings.time_limit, + settings.cancel_requested, + nullptr)) { settings.log.printf("Time limit exceeded\n"); return crossover_status_t::TIME_LIMIT; } diff --git a/cpp/src/dual_simplex/phase2.cpp b/cpp/src/dual_simplex/phase2.cpp index a5f10c3229..0989a1b849 100644 --- a/cpp/src/dual_simplex/phase2.cpp +++ b/cpp/src/dual_simplex/phase2.cpp @@ -1420,7 +1420,12 @@ i_t initialize_steepest_edge_norms(const lp_problem_t& lp, last_log = tic(); settings.log.printf("Initialized %d of %d steepest edge norms in %.2fs\n", k, m, now); } - if (toc(start_time) > settings.time_limit) { return -1; } + if (cuopt::solve_limit_reached(cuopt::mathematical_optimization::toc(start_time), + settings.time_limit, + settings.cancel_requested, + nullptr)) { + return -1; + } if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { return CONCURRENT_HALT_RETURN; } @@ -2573,6 +2578,16 @@ dual_status_t dual_phase2_with_advanced_basis(i_t phase, dual_status_t status = dual_status_t::UNSET; + // Concurrent root MIP: dual and barrier/PDLP/crossover race. On any phase-2 + // exit (optimal, time limit, cancel→TIME_LIMIT early return, etc.), raise + // concurrent_halt so waiters in solve_root_relaxation are not stuck forever + // when dual returns before the success-path epilogue below. + auto root_concurrent_halt_guard = cuopt::scope_guard([&]() { + if (phase == 2 && settings.inside_mip == 1 && settings.concurrent_halt != nullptr) { + *settings.concurrent_halt = 1; + } + }); + nvtx_range_guard init_scope("DualSimplex::phase2_advanced_init"); settings.log.printf("Dual Simplex Phase %d\n", phase); @@ -2606,7 +2621,12 @@ dual_status_t dual_phase2_with_advanced_basis(i_t phase, if (refactor_status == TIME_LIMIT_RETURN) { return dual_status_t::TIME_LIMIT; } if (refactor_status > 0) { return dual_status_t::NUMERICAL; } - if (toc(start_time) > settings.time_limit) { return dual_status_t::TIME_LIMIT; } + if (cuopt::solve_limit_reached(cuopt::mathematical_optimization::toc(start_time), + settings.time_limit, + settings.cancel_requested, + nullptr)) { + return dual_status_t::TIME_LIMIT; + } } // Populate c_basic after basis is initialized @@ -2618,7 +2638,12 @@ dual_status_t dual_phase2_with_advanced_basis(i_t phase, // Solve B'*y = cB ft.b_transpose_solve(c_basic, y); - if (toc(start_time) > settings.time_limit) { return dual_status_t::TIME_LIMIT; } + if (cuopt::solve_limit_reached(cuopt::mathematical_optimization::toc(start_time), + settings.time_limit, + settings.cancel_requested, + nullptr)) { + return dual_status_t::TIME_LIMIT; + } constexpr bool print_norms = false; if constexpr (print_norms) { settings.log.printf( @@ -2674,7 +2699,12 @@ dual_status_t dual_phase2_with_advanced_basis(i_t phase, xB_workspace, phase2_work_estimate); - if (toc(start_time) > settings.time_limit) { return dual_status_t::TIME_LIMIT; } + if (cuopt::solve_limit_reached(cuopt::mathematical_optimization::toc(start_time), + settings.time_limit, + settings.cancel_requested, + nullptr)) { + return dual_status_t::TIME_LIMIT; + } if (print_norms) { settings.log.printf("|| x || %e\n", vector_norm2(x)); } #ifdef COMPUTE_PRIMAL_RESIDUAL @@ -2990,7 +3020,12 @@ dual_status_t dual_phase2_with_advanced_basis(i_t phase, break; } - if (toc(start_time) > settings.time_limit) { return dual_status_t::TIME_LIMIT; } + if (cuopt::solve_limit_reached(cuopt::mathematical_optimization::toc(start_time), + settings.time_limit, + settings.cancel_requested, + nullptr)) { + return dual_status_t::TIME_LIMIT; + } if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { return dual_status_t::CONCURRENT_LIMIT; @@ -3597,7 +3632,12 @@ dual_status_t dual_phase2_with_advanced_basis(i_t phase, if (refactor_status > 0) { should_recompute_x = true; settings.log.printf("Failed to factorize basis. Iteration %d\n", iter); - if (toc(start_time) > settings.time_limit) { return dual_status_t::TIME_LIMIT; } + if (cuopt::solve_limit_reached(cuopt::mathematical_optimization::toc(start_time), + settings.time_limit, + settings.cancel_requested, + nullptr)) { + return dual_status_t::TIME_LIMIT; + } i_t count = 0; i_t deficient_size = 0; while (true) { @@ -3612,7 +3652,12 @@ dual_status_t dual_phase2_with_advanced_basis(i_t phase, iter, static_cast(deficient_size)); - if (toc(start_time) > settings.time_limit) { return dual_status_t::TIME_LIMIT; } + if (cuopt::solve_limit_reached(cuopt::mathematical_optimization::toc(start_time), + settings.time_limit, + settings.cancel_requested, + nullptr)) { + return dual_status_t::TIME_LIMIT; + } settings.threshold_partial_pivoting_tol = 1.0; count++; @@ -3717,7 +3762,12 @@ dual_status_t dual_phase2_with_advanced_basis(i_t phase, return dual_status_t::WORK_LIMIT; } - if (now > settings.time_limit) { return dual_status_t::TIME_LIMIT; } + if (cuopt::solve_limit_reached(cuopt::mathematical_optimization::toc(start_time), + settings.time_limit, + settings.cancel_requested, + nullptr)) { + return dual_status_t::TIME_LIMIT; + } if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { return dual_status_t::CONCURRENT_LIMIT; @@ -3737,10 +3787,7 @@ dual_status_t dual_phase2_with_advanced_basis(i_t phase, 100.0 * dense_delta_z / (sparse_delta_z + dense_delta_z)); ft.print_stats(); } - if (settings.inside_mip == 1 && settings.concurrent_halt != nullptr) { - settings.log.debug("Setting concurrent halt in Dual Simplex Phase 2\n"); - *settings.concurrent_halt = 1; - } + // concurrent_halt for inside_mip root is set by root_concurrent_halt_guard } return status; } diff --git a/cpp/src/dual_simplex/primal.cpp b/cpp/src/dual_simplex/primal.cpp index 78c7107ca3..47554af7be 100644 --- a/cpp/src/dual_simplex/primal.cpp +++ b/cpp/src/dual_simplex/primal.cpp @@ -312,8 +312,12 @@ primal_status_t primal_phase2(i_t phase, } else if (rank == TIME_LIMIT_RETURN) { return primal_status_t::TIME_LIMIT; } else if (rank < 0) { - return toc(start_time) > settings.time_limit ? primal_status_t::TIME_LIMIT - : primal_status_t::NUMERICAL; + return cuopt::solve_limit_reached(cuopt::mathematical_optimization::toc(start_time), + settings.time_limit, + settings.cancel_requested, + nullptr) + ? primal_status_t::TIME_LIMIT + : primal_status_t::NUMERICAL; } else if (rank != m) { settings.log.debug("Failed to factorize basis. rank %d m %d\n", rank, m); basis_repair(lp.A, @@ -345,8 +349,12 @@ primal_status_t primal_phase2(i_t phase, return primal_status_t::TIME_LIMIT; } else if (rank < 0) { settings.log.printf("Failed to factorize basis after repair. rank %d m %d\n", rank, m); - return toc(start_time) > settings.time_limit ? primal_status_t::TIME_LIMIT - : primal_status_t::NUMERICAL; + return cuopt::solve_limit_reached(cuopt::mathematical_optimization::toc(start_time), + settings.time_limit, + settings.cancel_requested, + nullptr) + ? primal_status_t::TIME_LIMIT + : primal_status_t::NUMERICAL; } else { settings.log.debug("Basis repaired\n"); } diff --git a/cpp/src/dual_simplex/right_looking_lu.cpp b/cpp/src/dual_simplex/right_looking_lu.cpp index 6a717cd257..060f943096 100644 --- a/cpp/src/dual_simplex/right_looking_lu.cpp +++ b/cpp/src/dual_simplex/right_looking_lu.cpp @@ -857,7 +857,12 @@ i_t right_looking_lu(const csc_matrix_t& A, if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { return CONCURRENT_HALT_RETURN; } - if (toc(start_time) > settings.time_limit) { return TIME_LIMIT_RETURN; } + if (cuopt::solve_limit_reached(cuopt::mathematical_optimization::toc(start_time), + settings.time_limit, + settings.cancel_requested, + nullptr)) { + return TIME_LIMIT_RETURN; + } // Find pivot that satisfies // abs(pivot) >= abstol, // abs(pivot) >= threshold_tol * max abs[pivot column] @@ -1068,7 +1073,12 @@ i_t right_looking_lu_row_permutation_only(const csc_matrix_t& A, toc(factorization_start_time)); last_print = tic(); } - if (toc(start_time) > settings.time_limit) { return TIME_LIMIT_RETURN; } + if (cuopt::solve_limit_reached(cuopt::mathematical_optimization::toc(start_time), + settings.time_limit, + settings.cancel_requested, + nullptr)) { + return TIME_LIMIT_RETURN; + } if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { if (!settings.inside_mip) { settings.log.printf("Concurrent halt\n"); } return CONCURRENT_HALT_RETURN; @@ -1713,7 +1723,12 @@ i_t right_looking_ldlt(const csc_matrix_t& A, if (settings.concurrent_halt != nullptr && *settings.concurrent_halt == 1) { return CONCURRENT_HALT_RETURN; } - if (toc(start_time) > settings.time_limit) { return TIME_LIMIT_RETURN; } + if (cuopt::solve_limit_reached(cuopt::mathematical_optimization::toc(start_time), + settings.time_limit, + settings.cancel_requested, + nullptr)) { + return TIME_LIMIT_RETURN; + } // Find symmetric pivot i_t pivot_p = -1; diff --git a/cpp/src/dual_simplex/simplex_solver_settings.hpp b/cpp/src/dual_simplex/simplex_solver_settings.hpp index 9ec57505f9..67df1d28c2 100644 --- a/cpp/src/dual_simplex/simplex_solver_settings.hpp +++ b/cpp/src/dual_simplex/simplex_solver_settings.hpp @@ -11,7 +11,9 @@ #include #include +#include #include +#include #include #include @@ -19,6 +21,7 @@ #include #include #include +#include #include namespace cuopt::mathematical_optimization { @@ -120,6 +123,28 @@ struct simplex_solver_settings_t { void enable_log_to_file() { log.enable_log_to_file(); } void set_log_filename(const std::string& log_filename) { log.set_log_file(log_filename); } void close_log_file() { log.close_log_file(); } + + // Unified time / concurrent / cancel / optional iteration poll (tic/toc style). + cuopt::solve_limit_reason_t check_solve_limits( + double start_time, std::optional iterations = std::nullopt) const + { + return cuopt::check_solve_limits( + cuopt::mathematical_optimization::toc(start_time), + time_limit, + cancel_requested, + concurrent_halt, + iterations, + iterations.has_value() + ? std::optional{static_cast(iteration_limit)} + : std::nullopt); + } + + /** True if cooperative cancel or concurrent halt was requested (ignores time). */ + bool cancel_or_halt_requested() const noexcept + { + return cuopt::cancel_or_halt_requested(cancel_requested, concurrent_halt); + } + i_t iteration_limit; i_t node_limit; f_t time_limit; @@ -238,6 +263,8 @@ struct simplex_solver_settings_t { mutable logger_t log; std::atomic* concurrent_halt; // if nullptr ignored, if !nullptr, 0 if solver should // continue, 1 if solver should halt + // Optional cooperative cancel (level-triggered). Non-owning. + std::atomic* cancel_requested{nullptr}; // Optional non-owning pointer to run-level benchmark stats. benchmark_info_t* benchmark_info_ptr = nullptr; }; diff --git a/cpp/src/dual_simplex/solve.cpp b/cpp/src/dual_simplex/solve.cpp index 7907abd3b9..d855a14833 100644 --- a/cpp/src/dual_simplex/solve.cpp +++ b/cpp/src/dual_simplex/solve.cpp @@ -26,6 +26,8 @@ #include #include +#include + #include #include @@ -37,6 +39,18 @@ namespace cuopt::mathematical_optimization::simplex { namespace { +template +lp_status_t remap_lp_status_if_cancelled(const simplex_solver_settings_t& settings, + lp_status_t status) +{ + return cuopt::remap_limit_status_if_cancelled(settings.cancel_requested, + status, + lp_status_t::CANCELLED, + lp_status_t::TIME_LIMIT, + lp_status_t::CONCURRENT_LIMIT, + lp_status_t::ITERATION_LIMIT); +} + template void write_matlab(const std::string& filename, const simplex::lp_problem_t& lp) { @@ -154,8 +168,12 @@ lp_status_t solve_linear_program_with_advanced_basis( raft::common::nvtx::range scope_presolve("DualSimplex::presolve"); ok = presolve(original_lp, settings, presolved_lp, presolve_info); } - if (ok == CONCURRENT_HALT_RETURN) { return lp_status_t::CONCURRENT_LIMIT; } - if (ok == TIME_LIMIT_RETURN) { return lp_status_t::TIME_LIMIT; } + if (ok == CONCURRENT_HALT_RETURN) { + return remap_lp_status_if_cancelled(settings, lp_status_t::CONCURRENT_LIMIT); + } + if (ok == TIME_LIMIT_RETURN) { + return remap_lp_status_if_cancelled(settings, lp_status_t::TIME_LIMIT); + } if (ok == -1) { return lp_status_t::INFEASIBLE; } constexpr bool write_out_matlab = false; @@ -372,8 +390,12 @@ lp_status_t solve_linear_program_with_barrier(const user_problem_t& us presolve_info_t presolve_info; lp_problem_t presolved_lp(handle_ptr, 1, 1, 1); const i_t ok = presolve(original_lp, barrier_settings, presolved_lp, presolve_info); - if (ok == CONCURRENT_HALT_RETURN) { return lp_status_t::CONCURRENT_LIMIT; } - if (ok == TIME_LIMIT_RETURN) { return lp_status_t::TIME_LIMIT; } + if (ok == CONCURRENT_HALT_RETURN) { + return remap_lp_status_if_cancelled(settings, lp_status_t::CONCURRENT_LIMIT); + } + if (ok == TIME_LIMIT_RETURN) { + return remap_lp_status_if_cancelled(settings, lp_status_t::TIME_LIMIT); + } if (ok == -1) { return lp_status_t::INFEASIBLE; } // Apply columns scaling to the presolve LP @@ -602,10 +624,14 @@ lp_status_t solve_linear_program_with_barrier(const user_problem_t& us solution.iterations = barrier_solution.iterations; } - if (barrier_status == lp_status_t::CONCURRENT_LIMIT) { return lp_status_t::CONCURRENT_LIMIT; } + if (barrier_status == lp_status_t::CONCURRENT_LIMIT) { + return remap_lp_status_if_cancelled(settings, lp_status_t::CONCURRENT_LIMIT); + } // If we aren't doing crossover, we're done - if (!settings.crossover || barrier_lp.Q.n > 0) { return barrier_status; } + if (!settings.crossover || barrier_lp.Q.n > 0) { + return remap_lp_status_if_cancelled(settings, barrier_status); + } if (settings.crossover && barrier_status == lp_status_t::OPTIMAL) { { @@ -676,7 +702,7 @@ lp_status_t solve_linear_program_with_barrier(const user_problem_t& us settings.log.printf("Crossover status: %d\n", crossover_status); if (crossover_status == crossover_status_t::OPTIMAL) { barrier_status = lp_status_t::OPTIMAL; } } - return barrier_status; + return remap_lp_status_if_cancelled(settings, barrier_status); } template @@ -716,7 +742,7 @@ lp_status_t solve_linear_program(const user_problem_t& user_problem, original_lp, start_time, settings, lp_solution, vstatus, edge_norms); if (status == lp_status_t::CONCURRENT_LIMIT) { solution.iterations = lp_solution.iterations; - return lp_status_t::CONCURRENT_LIMIT; + return remap_lp_status_if_cancelled(settings, lp_status_t::CONCURRENT_LIMIT); } uncrush_primal_solution(user_problem, original_lp, lp_solution.x, solution.x); uncrush_dual_solution( @@ -726,7 +752,7 @@ lp_status_t solve_linear_program(const user_problem_t& user_problem, solution.iterations = lp_solution.iterations; solution.l2_primal_residual = lp_solution.l2_primal_residual; solution.l2_dual_residual = lp_solution.l2_dual_residual; - return status; + return remap_lp_status_if_cancelled(settings, status); } template diff --git a/cpp/src/dual_simplex/solve.hpp b/cpp/src/dual_simplex/solve.hpp index 90c2dbd690..bdd81354fe 100644 --- a/cpp/src/dual_simplex/solve.hpp +++ b/cpp/src/dual_simplex/solve.hpp @@ -33,7 +33,8 @@ enum class lp_status_t { CUTOFF = 7, CONCURRENT_LIMIT = 8, WORK_LIMIT = 9, - UNSET = 10 + UNSET = 10, + CANCELLED = 11 }; static std::string lp_status_to_string(lp_status_t status) @@ -50,6 +51,7 @@ static std::string lp_status_to_string(lp_status_t status) case lp_status_t::CONCURRENT_LIMIT: return "CONCURRENT_LIMIT"; case lp_status_t::WORK_LIMIT: return "WORK_LIMIT"; case lp_status_t::UNSET: return "UNSET"; + case lp_status_t::CANCELLED: return "CANCELLED"; } return "UNKNOWN"; } diff --git a/cpp/src/grpc/server/grpc_job_management.cpp b/cpp/src/grpc/server/grpc_job_management.cpp index d35840cc92..026ec31053 100644 --- a/cpp/src/grpc/server/grpc_job_management.cpp +++ b/cpp/src/grpc/server/grpc_job_management.cpp @@ -232,7 +232,7 @@ void delete_log_file(const std::string& job_id) int cancel_job(const std::string& job_id, JobStatus& job_status_out, std::string& message) { - std::lock_guard lock(tracker_mutex); + std::unique_lock lock(tracker_mutex); auto it = job_tracker.find(job_id); if (it == job_tracker.end()) { @@ -273,18 +273,20 @@ int cancel_job(const std::string& job_id, JobStatus& job_status_out, std::string continue; } - pid_t worker_pid = job_queue[i].worker_pid.load(std::memory_order_relaxed); - - if (worker_pid > 0 && job_queue[i].claimed.load(std::memory_order_relaxed)) { - if (config.verbose) { - SERVER_LOG_DEBUG( - "[Server] Cancelling running job %s (killing worker %d)", job_id.c_str(), worker_pid); + pid_t worker_pid = job_queue[i].worker_pid.load(std::memory_order_relaxed); + const bool running = worker_pid > 0 && job_queue[i].claimed.load(std::memory_order_relaxed); + + // Prefer cooperative cancel: set the SHM flag the solver polls. Only kill + // the worker if it does not unwind within a grace period. + job_queue[i].cancelled.store(true, std::memory_order_release); + if (config.verbose) { + if (running) { + SERVER_LOG_DEBUG("[Server] Cancelling running job %s (cooperative cancel; worker %d)", + job_id.c_str(), + worker_pid); + } else { + SERVER_LOG_DEBUG("[Server] Cancelling queued job %s", job_id.c_str()); } - job_queue[i].cancelled.store(true, std::memory_order_release); - kill(worker_pid, SIGKILL); - } else { - if (config.verbose) { SERVER_LOG_DEBUG("[Server] Cancelling queued job %s", job_id.c_str()); } - job_queue[i].cancelled.store(true, std::memory_order_release); } it->second.status = JobStatus::CANCELLED; @@ -310,6 +312,44 @@ int cancel_job(const std::string& job_id, JobStatus& job_status_out, std::string } } + if (running) { + lock.unlock(); + // Return to the client immediately after setting the cancel flag. + // Grace wait + SIGTERM/SIGKILL run in the background so Cancel/Delete + // RPCs are not blocked for up to kGrace (client deadlines are ~60s). + std::thread([slot = i, worker_pid, job_id]() { + // Fixed cooperative grace. Solvers poll `cancelled` and should unwind + // without help; SIGTERM/SIGKILL is only a last resort if the worker is + // still claimed after this window. + constexpr auto kGrace = std::chrono::seconds(120); + const auto deadline = std::chrono::steady_clock::now() + kGrace; + + while (std::chrono::steady_clock::now() < deadline) { + if (!job_queue[slot].ready.load(std::memory_order_acquire) || + !job_queue[slot].claimed.load(std::memory_order_acquire) || + strcmp(job_queue[slot].job_id, job_id.c_str()) != 0) { + return; + } + if (kill(worker_pid, 0) != 0 && errno == ESRCH) { return; } + std::this_thread::sleep_for(std::chrono::milliseconds(50)); + } + // Slot may have been recycled during the wait — re-check identity. + if (!job_queue[slot].ready.load(std::memory_order_acquire) || + strcmp(job_queue[slot].job_id, job_id.c_str()) != 0) { + return; + } + if (kill(worker_pid, 0) != 0 && errno == ESRCH) { return; } + SERVER_LOG_WARN( + "[Server] Job %s still running after cooperative cancel grace; " + "sending SIGTERM then SIGKILL to worker %d", + job_id.c_str(), + worker_pid); + kill(worker_pid, SIGTERM); + std::this_thread::sleep_for(std::chrono::milliseconds(500)); + if (kill(worker_pid, 0) == 0) { kill(worker_pid, SIGKILL); } + }).detach(); + } + return 0; } diff --git a/cpp/src/grpc/server/grpc_worker.cpp b/cpp/src/grpc/server/grpc_worker.cpp index 3d0d00f839..0becac63cc 100644 --- a/cpp/src/grpc/server/grpc_worker.cpp +++ b/cpp/src/grpc/server/grpc_worker.cpp @@ -99,7 +99,8 @@ struct SolveResult { cuopt::remote::ChunkedResultHeader header; std::map> arrays; std::string error_message; - bool success = false; + bool success = false; + bool cancelled = false; }; // --------------------------------------------------------------------------- @@ -365,12 +366,14 @@ static SolveResult run_mip_solve(DeserializedJob& dj, raft::handle_t& handle, const std::string& log_file, const std::string& job_id, - int worker_id) + int worker_id, + std::atomic* cancel_requested) { SolveResult sr; try { - dj.mip_settings.log_file = log_file; - dj.mip_settings.log_to_console = config.log_to_console; + dj.mip_settings.log_file = log_file; + dj.mip_settings.log_to_console = config.log_to_console; + dj.mip_settings.cancel_requested = cancel_requested; // Create a per-solve incumbent callback wired to this worker's // incumbent pipe. Destroyed automatically when sr is returned. @@ -394,6 +397,13 @@ static SolveResult run_mip_solve(DeserializedJob& dj, auto gpu_solution = cuopt::mathematical_optimization::solve_mip(*gpu_problem, dj.mip_settings); SERVER_LOG_INFO("[Worker] solve_mip done"); + if (gpu_solution.get_termination_status() == + cuopt::mathematical_optimization::mip_termination_status_t::Cancelled) { + sr.error_message = "Job was cancelled"; + sr.cancelled = true; + return sr; + } + // solve_mip_helper catches cuopt::logic_error internally and stashes it // in mip_solution_t::error_status_ rather than rethrow (matches the LP // path's solver-API contract). Forward the error back to the client @@ -441,12 +451,14 @@ static SolveResult run_mip_solve(DeserializedJob& dj, // Exceptions are caught and returned as error messages. static SolveResult run_lp_solve(DeserializedJob& dj, raft::handle_t& handle, - const std::string& log_file) + const std::string& log_file, + std::atomic* cancel_requested) { SolveResult sr; try { - dj.lp_settings.log_file = log_file; - dj.lp_settings.log_to_console = config.log_to_console; + dj.lp_settings.log_file = log_file; + dj.lp_settings.log_to_console = config.log_to_console; + dj.lp_settings.cancel_requested = cancel_requested; SERVER_LOG_INFO("[Worker] Converting CPU problem to GPU problem..."); auto gpu_problem = dj.problem.to_optimization_problem(&handle); @@ -455,6 +467,13 @@ static SolveResult run_lp_solve(DeserializedJob& dj, auto gpu_solution = cuopt::mathematical_optimization::solve_lp(*gpu_problem, dj.lp_settings); SERVER_LOG_INFO("[Worker] solve_lp done"); + if (gpu_solution.get_termination_status() == + cuopt::mathematical_optimization::pdlp_termination_status_t::Cancelled) { + sr.error_message = "Job was cancelled"; + sr.cancelled = true; + return sr; + } + // solve_lp / solve_qcqp catch cuopt::logic_error internally and stash it // in optimization_problem_solution_t::error_status_ rather than rethrow // (long-standing solver-API contract; see solve.cu). Forward the error @@ -654,11 +673,16 @@ void worker_process(int worker_id) std::string log_file = get_log_file_path(job_id); raft::handle_t handle; - SolveResult result = (problem_category == cuopt::remote::MIP) - ? run_mip_solve(deserialized, handle, log_file, job_id, worker_id) - : run_lp_solve(deserialized, handle, log_file); + SolveResult result = + (problem_category == cuopt::remote::MIP) + ? run_mip_solve(deserialized, handle, log_file, job_id, worker_id, &job.cancelled) + : run_lp_solve(deserialized, handle, log_file, &job.cancelled); - publish_result(result, job_id, worker_id); + if (result.cancelled || job.cancelled.load(std::memory_order_acquire)) { + store_simple_result(job_id, worker_id, RESULT_CANCELLED, "Job was cancelled"); + } else { + publish_result(result, job_id, worker_id); + } reset_job_slot(job); SERVER_LOG_INFO("[Worker %d] Completed job: %s (success: %d)", diff --git a/cpp/src/mip_heuristics/diversity/diversity_manager.cu b/cpp/src/mip_heuristics/diversity/diversity_manager.cu index 61c90944f4..384a0e1f42 100644 --- a/cpp/src/mip_heuristics/diversity/diversity_manager.cu +++ b/cpp/src/mip_heuristics/diversity/diversity_manager.cu @@ -19,7 +19,9 @@ #include #include +#include +#include #include #include @@ -248,6 +250,7 @@ void diversity_manager_t::add_user_given_solutions( lp_settings.tolerance = problem_ptr->tolerances.absolute_tolerance; lp_settings.save_state = false; lp_settings.return_first_feasible = true; + lp_settings.cancel_requested = context.settings.cancel_requested; run_lp_with_vars_fixed(*problem_ptr, sol, problem_ptr->integer_indices, @@ -284,7 +287,8 @@ bool diversity_manager_t::run_presolve(f_t time_limit, timer_t global_ { raft::common::nvtx::range fun_scope("run_presolve"); CUOPT_LOG_INFO("\nRunning cuOpt presolve"); - timer_t presolve_timer(time_limit); + + timer_t presolve_timer(time_limit, context.settings.cancel_requested); auto term_crit = ls.constraint_prop.bounds_update.solve(*problem_ptr); if (ls.constraint_prop.bounds_update.infeas_constraints_count > 0) { @@ -308,7 +312,7 @@ bool diversity_manager_t::run_presolve(f_t time_limit, timer_t global_ // Run probing cache before trivial presolve to discover variable implications const f_t max_time_on_probing = diversity_config.max_time_on_probing; f_t time_for_probing_cache = std::min(max_time_on_probing, time_limit); - timer_t probing_timer{time_for_probing_cache}; + timer_t probing_timer{time_for_probing_cache, context.settings.cancel_requested}; // this function computes probing cache, finds singletons, substitutions and changes the problem bool problem_is_infeasible = compute_probing_cache(ls.constraint_prop.bounds_update, *problem_ptr, probing_timer); @@ -377,7 +381,7 @@ void diversity_manager_t::generate_quick_feasible_solution() // min 1 second, max 10 seconds const f_t generate_fast_solution_time = std::min(diversity_config.max_fast_sol_time, std::max(1., timer.remaining_time() / 20.)); - timer_t sol_timer(generate_fast_solution_time); + timer_t sol_timer(generate_fast_solution_time, context.settings.cancel_requested); // do very short LP run to get somewhere close to the optimal point ls.generate_fast_solution(solution, sol_timer); if (solution.get_feasible()) { @@ -402,6 +406,12 @@ void diversity_manager_t::generate_quick_feasible_solution() template bool diversity_manager_t::check_b_b_preemption() { + if (context.settings.cancel_requested != nullptr && + context.settings.cancel_requested->load(std::memory_order_acquire)) { + // Stop heuristics in parallel with B&B when cancel is requested. + context.preempt_heuristic_solver_.store(true, std::memory_order_release); + population.preempt_heuristic_solver(); + } if (context.preempt_heuristic_solver_.load()) { if (population.current_size() == 0) { population.allocate_solutions(); } population.add_external_solutions_to_population(); @@ -547,6 +557,7 @@ solution_t diversity_manager_t::run_solver() pdlp_settings.time_limit = lp_time_limit; pdlp_settings.first_primal_feasible = false; pdlp_settings.concurrent_halt = &global_concurrent_halt; + pdlp_settings.cancel_requested = context.settings.cancel_requested; pdlp_settings.method = method_t::Concurrent; pdlp_settings.inside_mip = true; pdlp_settings.pdlp_solver_mode = pdlp_solver_mode_t::Stable2; @@ -554,7 +565,7 @@ solution_t diversity_manager_t::run_solver() 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); + timer_t lp_timer(lp_time_limit, context.settings.cancel_requested); 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 @@ -897,6 +908,7 @@ diversity_manager_t::recombine_and_local_search(solution_t& lp_settings.return_first_feasible = false; lp_settings.save_state = true; lp_settings.per_constraint_residual = true; + lp_settings.cancel_requested = context.settings.cancel_requested; run_lp_with_vars_fixed(*lp_offspring.problem_ptr, lp_offspring, lp_offspring.problem_ptr->integer_indices, diff --git a/cpp/src/mip_heuristics/diversity/recombiners/bound_prop_recombiner.cuh b/cpp/src/mip_heuristics/diversity/recombiners/bound_prop_recombiner.cuh index 42fd838105..24a1657ef3 100644 --- a/cpp/src/mip_heuristics/diversity/recombiners/bound_prop_recombiner.cuh +++ b/cpp/src/mip_heuristics/diversity/recombiners/bound_prop_recombiner.cuh @@ -177,7 +177,8 @@ class bound_prop_recombiner_t : public recombiner_t { if (guiding_solution.get_feasible() && !a.problem_ptr->expensive_to_fix_vars) { this->compute_vars_to_fix(offspring, vars_to_fix, n_vars_from_other, n_vars_from_guiding); auto [fixed_problem, fixed_assignment, variable_map] = offspring.fix_variables(vars_to_fix); - timer_t timer(bp_recombiner_config_t::bounds_prop_time_limit); + timer_t timer(bp_recombiner_config_t::bounds_prop_time_limit, + this->context.settings.cancel_requested); rmm::device_uvector old_assignment(offspring.assignment, offspring.handle_ptr->get_stream()); offspring.handle_ptr->sync_stream(); @@ -211,7 +212,8 @@ class bound_prop_recombiner_t : public recombiner_t { // "Feasible after unfix should be same as feasible after bounds prop!"); a.handle_ptr->sync_stream(); } else { - timer_t timer(bp_recombiner_config_t::bounds_prop_time_limit); + timer_t timer(bp_recombiner_config_t::bounds_prop_time_limit, + this->context.settings.cancel_requested); get_probing_values_for_infeasible( guiding_solution, other_solution, offspring, probing_values, n_vars_from_other); probing_config.probing_values = host_copy(probing_values, offspring.handle_ptr->get_stream()); diff --git a/cpp/src/mip_heuristics/diversity/recombiners/fp_recombiner.cuh b/cpp/src/mip_heuristics/diversity/recombiners/fp_recombiner.cuh index 85909c9d69..c5e2fc463d 100644 --- a/cpp/src/mip_heuristics/diversity/recombiners/fp_recombiner.cuh +++ b/cpp/src/mip_heuristics/diversity/recombiners/fp_recombiner.cuh @@ -76,6 +76,7 @@ class fp_recombiner_t : public recombiner_t { lp_settings.return_first_feasible = true; lp_settings.save_state = true; lp_settings.check_infeasibility = true; + lp_settings.cancel_requested = this->context.settings.cancel_requested; // run lp with infeasibility detection on auto lp_response = get_relaxed_lp_solution(fixed_problem, fixed_assignment, offspring.lp_state, lp_settings); @@ -96,7 +97,7 @@ class fp_recombiner_t : public recombiner_t { offspring.handle_ptr->sync_stream(); offspring.assignment = std::move(fixed_assignment); cuopt_func_call(offspring.test_variable_bounds(false)); - timer_t timer(fp_recombiner_config_t::fp_time_limit); + timer_t timer(fp_recombiner_config_t::fp_time_limit, this->context.settings.cancel_requested); fp.timer = timer; fp.cycle_queue.reset(offspring); fp.reset(); diff --git a/cpp/src/mip_heuristics/feasibility_jump/early_cpufj.cu b/cpp/src/mip_heuristics/feasibility_jump/early_cpufj.cu index ba14e657d5..5270d539ab 100644 --- a/cpp/src/mip_heuristics/feasibility_jump/early_cpufj.cu +++ b/cpp/src/mip_heuristics/feasibility_jump/early_cpufj.cu @@ -15,9 +15,11 @@ template early_cpufj_t::early_cpufj_t( const optimization_problem_t& op_problem, const typename mip_solver_settings_t::tolerances_t& tolerances, - early_incumbent_callback_t incumbent_callback) + early_incumbent_callback_t incumbent_callback, + std::atomic* cancel_requested) : early_heuristic_t>( - op_problem, tolerances, std::move(incumbent_callback)) + op_problem, tolerances, std::move(incumbent_callback)), + cancel_requested_(cancel_requested) { } @@ -36,7 +38,12 @@ void early_cpufj_t::start() this->preemption_flag_.store(false); this->start_time_ = std::chrono::steady_clock::now(); - fj_cpu_ = init_fj_cpu_standalone(*this->problem_ptr_, *this->solution_ptr_, preemption_flag_); + // Prefer the gRPC cancel flag as the climber's preemption signal so cancel + // aborts early CPUFJ during Papilo without waiting for stop()/taskwait. + std::atomic& preempt_ref = + (cancel_requested_ != nullptr) ? *cancel_requested_ : preemption_flag_; + + fj_cpu_ = init_fj_cpu_standalone(*this->problem_ptr_, *this->solution_ptr_, preempt_ref); fj_cpu_->log_prefix = "[Early CPUFJ] "; @@ -56,6 +63,11 @@ void early_cpufj_t::stop() if (!fj_cpu_) { return; } preemption_flag_.store(true); + // The climber exits when either its preempt atomic or halted is true. If + // cancel_requested_ was wired as that preempt atomic, writing + // preemption_flag_ here does not affect the climber; fj_cpu_->halted below + // still stops it on the normal (non-cancel) path. On cancel, + // *cancel_requested_ is already true so the climber is already exiting. fj_cpu_->halted = true; #pragma omp taskwait depend(in : *fj_cpu_) // Wait for the early CPUFJ task to finish diff --git a/cpp/src/mip_heuristics/feasibility_jump/early_cpufj.cuh b/cpp/src/mip_heuristics/feasibility_jump/early_cpufj.cuh index e2bb2c07b2..4ccfecfa72 100644 --- a/cpp/src/mip_heuristics/feasibility_jump/early_cpufj.cuh +++ b/cpp/src/mip_heuristics/feasibility_jump/early_cpufj.cuh @@ -20,7 +20,8 @@ class early_cpufj_t : public early_heuristic_t public: early_cpufj_t(const optimization_problem_t& op_problem, const typename mip_solver_settings_t::tolerances_t& tolerances, - early_incumbent_callback_t incumbent_callback); + early_incumbent_callback_t incumbent_callback, + std::atomic* cancel_requested = nullptr); ~early_cpufj_t(); @@ -32,6 +33,7 @@ class early_cpufj_t : public early_heuristic_t private: std::unique_ptr> fj_cpu_; std::atomic preemption_flag_{false}; + std::atomic* cancel_requested_{nullptr}; }; } // namespace cuopt::mathematical_optimization::mip diff --git a/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cu b/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cu index a6665e57e1..f3c6d02f27 100644 --- a/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cu +++ b/cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cu @@ -840,7 +840,7 @@ i_t fj_t::host_loop(solution_t& solution, i_t climber_idx) data.incumbent_quality.set_value_async(obj, handle_ptr->get_stream()); - timer_t timer(settings.time_limit); + timer_t timer(settings.time_limit, context.settings.cancel_requested); i_t steps; bool limit_reached = false; for (steps = 0; steps < std::numeric_limits::max(); steps += iterations_per_graph) { @@ -1053,7 +1053,7 @@ template i_t fj_t::solve(solution_t& solution) { raft::common::nvtx::range scope("fj_solve"); - timer_t timer(settings.time_limit); + timer_t timer(settings.time_limit, context.settings.cancel_requested); handle_ptr = const_cast(solution.handle_ptr); pb_ptr = solution.problem_ptr; last_reported_objective_ = std::numeric_limits::infinity(); diff --git a/cpp/src/mip_heuristics/local_search/feasibility_pump/feasibility_pump.cu b/cpp/src/mip_heuristics/local_search/feasibility_pump/feasibility_pump.cu index bc47027ed5..e64ecdcb55 100644 --- a/cpp/src/mip_heuristics/local_search/feasibility_pump/feasibility_pump.cu +++ b/cpp/src/mip_heuristics/local_search/feasibility_pump/feasibility_pump.cu @@ -248,7 +248,8 @@ bool feasibility_pump_t::round(solution_t& solution) { bool result; CUOPT_LOG_DEBUG("Rounding the point"); - timer_t bounds_prop_timer(std::max(0.05, std::min(0.5, timer.remaining_time() / 10.))); + timer_t bounds_prop_timer(std::max(0.05, std::min(0.5, timer.remaining_time() / 10.)), + context.settings.cancel_requested); const f_t lp_run_time_after_feasible = 0.; bool old_var = constraint_prop.round_all_vars; f_t old_time = constraint_prop.max_time_for_bounds_prop; diff --git a/cpp/src/mip_heuristics/local_search/local_search.cu b/cpp/src/mip_heuristics/local_search/local_search.cu index 75c4185949..72ea346a59 100644 --- a/cpp/src/mip_heuristics/local_search/local_search.cu +++ b/cpp/src/mip_heuristics/local_search/local_search.cu @@ -145,7 +145,11 @@ void local_search_t::stop_cpufj_scratch_threads() { if (omp_get_num_threads() < CUOPT_MIP_FJ_REQUIRED_THREAD_COUNT) return; + // Climbers are null until the matching start_* fills them (and start_lptopt + // is independent of start_scratch). Skip unset slots so stop is safe if a + // start never ran. for (size_t i = 0; i < scratch_cpu_fj.size(); ++i) { + if (!scratch_cpu_fj[i]) continue; scratch_cpu_fj[i]->halted = true; #pragma omp taskwait depend(in : *scratch_cpu_fj[i]) // Wait for each scratch CPU FJ task to finish } @@ -231,7 +235,7 @@ bool local_search_t::do_fj_solve(solution_t& solution, { if (time_limit == 0.) return solution.get_feasible(); - timer_t timer(time_limit); + timer_t timer(time_limit, context.settings.cancel_requested); const auto old_n_cstr_weights = in_fj.cstr_weights.size(); const auto expected_n_cstr_weights = static_cast(solution.problem_ptr->n_constraints); // in case this is the first time run, resize @@ -349,7 +353,8 @@ void local_search_t::generate_fast_solution(solution_t& solu fj.settings.feasibility_run = true; fj.settings.time_limit = std::min(30., timer.remaining_time()); while (!context.diversity_manager_ptr->check_b_b_preemption() && !timer.check_time_limit()) { - timer_t constr_prop_timer = timer_t(std::min(timer.remaining_time(), 2.)); + timer_t constr_prop_timer(std::min(timer.remaining_time(), 2.), + context.settings.cancel_requested); // do constraint prop on lp optimal solution constraint_prop.apply_round(solution, 1., constr_prop_timer); if (solution.compute_feasibility()) { return; } @@ -376,10 +381,10 @@ bool local_search_t::run_local_search(solution_t& solution, if (!solution.get_feasible()) { if (ls_config.at_least_one_parent_feasible) { fj_settings.time_limit = 0.5; - timer = timer_t(fj_settings.time_limit); + timer = timer_t(fj_settings.time_limit, context.settings.cancel_requested); } else { fj_settings.time_limit = 0.25; - timer = timer_t(fj_settings.time_limit); + timer = timer_t(fj_settings.time_limit, context.settings.cancel_requested); } } else { fj_settings.time_limit = std::min(1., timer.remaining_time()); @@ -498,14 +503,16 @@ bool local_search_t::check_fj_on_lp_optimal(solution_t& solu } cuopt_func_call(solution.test_variable_bounds(false)); 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.)); + timer_t bounds_prop_timer(std::min(timer.remaining_time(), 10.), + context.settings.cancel_requested); 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 = std::min(lp_run_time, timer.remaining_time()); - lp_settings.tolerance = solution.problem_ptr->tolerances.absolute_tolerance; + lp_settings.time_limit = std::min(lp_run_time, timer.remaining_time()); + lp_settings.tolerance = solution.problem_ptr->tolerances.absolute_tolerance; + lp_settings.cancel_requested = context.settings.cancel_requested; run_lp_with_vars_fixed( *solution.problem_ptr, solution, solution.problem_ptr->integer_indices, lp_settings); } else { @@ -569,7 +576,7 @@ bool local_search_t::run_staged_fp(solution_t& solution, } CUOPT_LOG_DEBUG("Running staged FP from beginning it %d", i); fp.relax_general_integers(solution); - timer_t binary_timer(timer.remaining_time() / 3); + timer_t binary_timer(timer.remaining_time() / 3, context.settings.cancel_requested); i_t binary_it_counter = 0; for (; binary_it_counter < 100; ++binary_it_counter) { population_ptr->add_external_solutions_to_population(); @@ -747,7 +754,7 @@ bool local_search_t::run_fp(solution_t& solution, is_feasible ? solution.get_objective() : std::numeric_limits::max(); rmm::device_uvector best_solution(solution.assignment, solution.handle_ptr->get_stream()); problem_t* old_problem_ptr = solution.problem_ptr; - fp.timer = timer_t(timer.remaining_time()); + fp.timer = timer_t(timer.remaining_time(), context.settings.cancel_requested); // if it has not been initialized yet, create a new problem and move it to the cut problem if (!problem_with_objective_cut.cutting_plane_added) { problem_with_objective_cut = std::move(problem_t(*old_problem_ptr)); @@ -848,7 +855,7 @@ bool local_search_t::generate_solution(solution_t& solution, { raft::common::nvtx::range fun_scope("generate_solution"); cuopt_assert(population_ptr != nullptr, "Population pointer must not be null"); - timer_t timer(time_limit); + timer_t timer(time_limit, context.settings.cancel_requested); auto n_vars = solution.problem_ptr->n_variables; auto n_binary_vars = solution.problem_ptr->get_n_binary_variables(); auto n_integer_vars = solution.problem_ptr->n_integer_vars; diff --git a/cpp/src/mip_heuristics/presolve/bounds_presolve.cu b/cpp/src/mip_heuristics/presolve/bounds_presolve.cu index e5a7f249f1..1c3bc51226 100644 --- a/cpp/src/mip_heuristics/presolve/bounds_presolve.cu +++ b/cpp/src/mip_heuristics/presolve/bounds_presolve.cu @@ -232,7 +232,7 @@ termination_criterion_t bound_presolve_t::solve(problem_t& p i_t var_idx) { auto& handle_ptr = pb.handle_ptr; - timer_t timer(settings.time_limit); + timer_t timer(settings.time_limit, context.settings.cancel_requested); copy_input_bounds(pb); upd.lb.set_element_async(var_idx, var_lb, handle_ptr->get_stream()); upd.ub.set_element_async(var_idx, var_ub, handle_ptr->get_stream()); @@ -245,7 +245,7 @@ termination_criterion_t bound_presolve_t::solve( const std::vector>& var_probe_val_pairs, bool use_host_bounds) { - timer_t timer(settings.time_limit); + timer_t timer(settings.time_limit, context.settings.cancel_requested); auto& handle_ptr = pb.handle_ptr; if (use_host_bounds) { update_device_bounds(handle_ptr); @@ -260,7 +260,7 @@ termination_criterion_t bound_presolve_t::solve( template termination_criterion_t bound_presolve_t::solve(problem_t& pb) { - timer_t timer(settings.time_limit); + timer_t timer(settings.time_limit, context.settings.cancel_requested); auto& handle_ptr = pb.handle_ptr; copy_input_bounds(pb); return bound_update_loop(pb, timer); diff --git a/cpp/src/mip_heuristics/presolve/conflict_graph/clique_table.cu b/cpp/src/mip_heuristics/presolve/conflict_graph/clique_table.cu index a8e6997572..4179c1984c 100644 --- a/cpp/src/mip_heuristics/presolve/conflict_graph/clique_table.cu +++ b/cpp/src/mip_heuristics/presolve/conflict_graph/clique_table.cu @@ -674,8 +674,8 @@ void find_initial_cliques(user_problem_t& problem, cuopt::timer_t& timer, omp_atomic_t* signal_extend) { - cuopt::timer_t stage_timer(std::numeric_limits::infinity()); #ifdef DEBUG_CLIQUE_TABLE + cuopt::timer_t stage_timer(std::numeric_limits::infinity(), timer.get_cancel_requested()); double t_fill = 0.; double t_coeff = 0.; double t_sort = 0.; @@ -707,7 +707,8 @@ void find_initial_cliques(user_problem_t& problem, clique_config.max_clique_size_for_extension); clique_table->tolerances = tolerances; double time_limit_for_additional_cliques = timer.remaining_time() / 2; - cuopt::timer_t additional_cliques_timer(time_limit_for_additional_cliques); + cuopt::timer_t additional_cliques_timer(time_limit_for_additional_cliques, + timer.get_cancel_requested()); double find_work_estimate = 0.0; // Always build base cliques in full; signal_extend only gates the extension phase. for (const auto& knapsack_constraint : knapsack_constraints) { diff --git a/cpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve.cu b/cpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve.cu index 017bf32e91..7a2884c9ad 100644 --- a/cpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve.cu +++ b/cpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve.cu @@ -627,7 +627,7 @@ termination_criterion_t load_balanced_bounds_presolve_t::solve(f_t var f_t var_ub, i_t var_idx) { - timer_t timer(settings.time_limit); + timer_t timer(settings.time_limit, context.settings.cancel_requested); auto& handle_ptr = pb->handle_ptr; copy_input_bounds(*pb); vars_bnd.set_element_async(2 * var_idx, var_lb, handle_ptr->get_stream()); @@ -639,7 +639,7 @@ template termination_criterion_t load_balanced_bounds_presolve_t::solve( raft::device_span input_bounds) { - timer_t timer(settings.time_limit); + timer_t timer(settings.time_limit, context.settings.cancel_requested); auto& handle_ptr = pb->handle_ptr; if (input_bounds.size() != 0) { raft::copy(vars_bnd.data(), input_bounds.data(), input_bounds.size(), handle_ptr->get_stream()); @@ -668,7 +668,7 @@ template termination_criterion_t load_balanced_bounds_presolve_t::solve( const std::vector>& var_probe_val_pairs, bool use_host_bounds) { - timer_t timer(settings.time_limit); + timer_t timer(settings.time_limit, context.settings.cancel_requested); auto& handle_ptr = pb->handle_ptr; if (use_host_bounds) { update_device_bounds(handle_ptr); diff --git a/cpp/src/mip_heuristics/presolve/multi_probe.cu b/cpp/src/mip_heuristics/presolve/multi_probe.cu index f1adf28650..a35a10cbd3 100644 --- a/cpp/src/mip_heuristics/presolve/multi_probe.cu +++ b/cpp/src/mip_heuristics/presolve/multi_probe.cu @@ -380,7 +380,7 @@ termination_criterion_t multi_probe_t::solve_for_interval( const std::tuple, std::pair>& var_interval_vals, const raft::handle_t* handle_ptr) { - timer_t timer(settings.time_limit); + timer_t timer(settings.time_limit, context.settings.cancel_requested); copy_problem_into_probing_buffers(pb, handle_ptr); set_interval_bounds(var_interval_vals, pb, handle_ptr); @@ -394,7 +394,7 @@ termination_criterion_t multi_probe_t::solve( const std::tuple, std::vector, std::vector>& var_probe_vals, bool use_host_bounds) { - timer_t timer(settings.time_limit); + timer_t timer(settings.time_limit, context.settings.cancel_requested); auto& handle_ptr = pb.handle_ptr; if (use_host_bounds) { update_device_bounds(handle_ptr); diff --git a/cpp/src/mip_heuristics/presolve/third_party_presolve.cpp b/cpp/src/mip_heuristics/presolve/third_party_presolve.cpp index 6a5bd341fb..f407adc25e 100644 --- a/cpp/src/mip_heuristics/presolve/third_party_presolve.cpp +++ b/cpp/src/mip_heuristics/presolve/third_party_presolve.cpp @@ -745,11 +745,20 @@ void set_presolve_parameters( }; // Papilo has work unit measurements for probing. Because of this when the first batch fails to // produce any reductions, the algorithm stops. To avoid stopping the algorithm, we set a - // minimum badge size to a huge value. The time limit makes sure that we exit if it takes too - // long + // minimum badge size to a large value. The time limit makes sure that we exit if it takes too + // long. + // + // Also cap maxbadgesize so presolve.tlim / cancel can be observed between badges. Without a + // cap, one badge can be ~ncols/2 and run far past tlim (observed on seymour1.mps), blocking + // cooperative cancel until the badge finishes. The alternative would be a much longer + // watchdog grace before SIGKILL while Papilo ignores cancel mid-badge; capping badges is + // preferable. Keep minbadgesize <= maxbadgesize. Cap may slow probing on huge MIPs vs + // unlimited badges. if (reduction_allowed("probing")) { - int min_badgesize = std::max(ncols / 2, 32); + constexpr int max_badgesize = 512; + int min_badgesize = std::min(std::max(ncols / 2, 32), max_badgesize); params.setParameter("probing.minbadgesize", min_badgesize); + params.setParameter("probing.maxbadgesize", max_badgesize); } if (reduction_allowed("cliquemerging")) { params.setParameter("cliquemerging.enabled", true); diff --git a/cpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cu b/cpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cu index 14a56cb903..617372c0ec 100644 --- a/cpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cu +++ b/cpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cu @@ -49,6 +49,7 @@ optimization_problem_solution_t get_relaxed_lp_solution( pdlp_settings.tolerances.relative_dual_tolerance = settings.tolerance / tolerance_divisor; pdlp_settings.time_limit = settings.time_limit; pdlp_settings.concurrent_halt = settings.concurrent_halt; + pdlp_settings.cancel_requested = settings.cancel_requested; pdlp_settings.per_constraint_residual = settings.per_constraint_residual; pdlp_settings.first_primal_feasible = settings.return_first_feasible; pdlp_settings.pdlp_solver_mode = pdlp_solver_mode_t::Stable2; @@ -84,7 +85,7 @@ optimization_problem_solution_t get_relaxed_lp_solution( // before LP flush the logs as it takes quite some time cuopt::default_logger().flush(); // temporarily add timer - auto start_time = timer_t(pdlp_settings.time_limit); + auto start_time = timer_t(pdlp_settings.time_limit, pdlp_settings.cancel_requested); lp_solver.set_inside_mip(true); auto solver_response = lp_solver.run_solver(start_time); diff --git a/cpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cuh b/cpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cuh index e13cc04aa8..c32f68a7e5 100644 --- a/cpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cuh +++ b/cpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cuh @@ -17,14 +17,15 @@ namespace cuopt::mathematical_optimization::mip { struct relaxed_lp_settings_t { - double tolerance = 1e-4; - double time_limit = 1.0; - bool check_infeasibility = true; - bool return_first_feasible = false; - bool save_state = true; - bool per_constraint_residual = true; - bool has_initial_primal = true; - std::atomic* concurrent_halt = nullptr; + double tolerance = 1e-4; + double time_limit = 1.0; + bool check_infeasibility = true; + bool return_first_feasible = false; + bool save_state = true; + bool per_constraint_residual = true; + bool has_initial_primal = true; + std::atomic* concurrent_halt = nullptr; + std::atomic* cancel_requested = nullptr; }; template diff --git a/cpp/src/mip_heuristics/solution/solution.cu b/cpp/src/mip_heuristics/solution/solution.cu index 3b00fca7a8..a0df11b6f2 100644 --- a/cpp/src/mip_heuristics/solution/solution.cu +++ b/cpp/src/mip_heuristics/solution/solution.cu @@ -75,6 +75,7 @@ solution_t::solution_t(const solution_t& other) h_infeasibility_cost(other.h_infeasibility_cost), is_feasible(other.is_feasible), is_problem_fully_reduced(other.is_problem_fully_reduced), + cancelled_(other.cancelled_), is_scaled_(other.is_scaled_), post_process_completed(other.post_process_completed), lp_state(other.lp_state) @@ -104,6 +105,7 @@ void solution_t::copy_from(const solution_t& other_sol) handle_ptr->get_stream()); is_feasible = other_sol.is_feasible; is_problem_fully_reduced = other_sol.is_problem_fully_reduced; + cancelled_ = other_sol.cancelled_; is_scaled_ = other_sol.is_scaled_; post_process_completed = other_sol.post_process_completed; expand_device_copy( @@ -204,6 +206,18 @@ void solution_t::set_problem_fully_reduced() is_problem_fully_reduced = true; } +template +void solution_t::set_cancelled() +{ + cancelled_ = true; +} + +template +bool solution_t::get_cancelled() const +{ + return cancelled_; +} + template std::vector solution_t::get_host_assignment() { @@ -623,6 +637,7 @@ cuopt::mathematical_optimization::mip_solution_t solution_t: auto term_reason = not_optimal ? mip_termination_status_t::FeasibleFound : mip_termination_status_t::Optimal; if (is_problem_fully_reduced) { term_reason = mip_termination_status_t::Optimal; } + if (cancelled_) { term_reason = mip_termination_status_t::Cancelled; } auto sol = cuopt::mathematical_optimization::mip_solution_t(std::move(assignment), problem_ptr->var_names, @@ -636,11 +651,14 @@ cuopt::mathematical_optimization::mip_solution_t solution_t: if (log_stats) { sol.log_detailed_summary(); } return sol; } else { + mip_termination_status_t term_reason = mip_termination_status_t::TimeLimit; + if (is_problem_fully_reduced) { + term_reason = mip_termination_status_t::Infeasible; + } else if (cancelled_) { + term_reason = mip_termination_status_t::Cancelled; + } return cuopt::mathematical_optimization::mip_solution_t{ - is_problem_fully_reduced ? mip_termination_status_t::Infeasible - : mip_termination_status_t::TimeLimit, - stats, - handle_ptr->get_stream()}; + term_reason, stats, handle_ptr->get_stream()}; } } diff --git a/cpp/src/mip_heuristics/solution/solution.cuh b/cpp/src/mip_heuristics/solution/solution.cuh index f243937d3e..3469f52d3b 100644 --- a/cpp/src/mip_heuristics/solution/solution.cuh +++ b/cpp/src/mip_heuristics/solution/solution.cuh @@ -69,6 +69,9 @@ class solution_t { bool get_problem_fully_reduced(); // sets the is_problem_fully_reduced flag to 1 void set_problem_fully_reduced(); + // Mark that cooperative cancel requested early exit (surfaces as Cancelled). + void set_cancelled(); + bool get_cancelled() const; // computes the number of integral variables that have integral value i_t compute_number_of_integers(); // computes the l2 residual value from the excess values @@ -143,6 +146,7 @@ class solution_t { f_t h_infeasibility_cost = 0.; bool is_feasible = false; bool is_problem_fully_reduced{false}; + bool cancelled_{false}; bool is_scaled_{false}; bool post_process_completed{false}; lp_state_t lp_state; diff --git a/cpp/src/mip_heuristics/solve.cu b/cpp/src/mip_heuristics/solve.cu index b970eb1bf8..3aca418419 100644 --- a/cpp/src/mip_heuristics/solve.cu +++ b/cpp/src/mip_heuristics/solve.cu @@ -28,6 +28,7 @@ #include #include #include +#include #include #include @@ -63,6 +64,7 @@ #include #include +#include #include #include @@ -239,6 +241,7 @@ mip_solution_t run_mip_solver( auto stats = solver.get_solver_stats(); stats.total_solve_time = timer.elapsed_time(); sol.post_process_completed = true; + if (cuopt::cancel_flag_set(settings.cancel_requested)) { sol.set_cancelled(); } return sol.get_solution(false, stats, false); } @@ -291,8 +294,10 @@ mip_solution_t run_mip_solver( user_assignment, no_bound); }; - early_cpufj = std::make_unique>( - *problem.original_problem_ptr, settings.get_tolerances(), incumbent_callback); + early_cpufj = std::make_unique>(*problem.original_problem_ptr, + settings.get_tolerances(), + incumbent_callback, + settings.cancel_requested); // Convert initial_upper_bound from user-space to the CPUFJ's solver-space (papilo-presolved). // problem.get_solver_obj_from_user_obj uses the papilo offset/scale (matching the CPUFJ). if (std::isfinite(initial_upper_bound)) { @@ -373,7 +378,7 @@ mip_solution_t solve_mip_helper(optimization_problem_t& op_p if (settings.seed >= 0) { cuopt::seed_generator::set_seed(settings.seed); } raft::common::nvtx::range fun_scope("Running solver"); - auto timer = timer_t(time_limit); + auto timer = timer_t(time_limit, settings.cancel_requested); problem_checking_t::check_problem_representation(op_problem); problem_checking_t::check_initial_solution_representation(op_problem, settings); @@ -546,7 +551,7 @@ mip_solution_t solve_mip_helper(optimization_problem_t& op_p // Start early CPUFJ on original problem (will restart on presolved problem after Papilo) early_cpufj = std::make_unique>( - op_problem, settings.get_tolerances(), early_fj_callback); + op_problem, settings.get_tolerances(), early_fj_callback, settings.cancel_requested); early_cpufj->start(); CUOPT_LOG_DEBUG("Started early CPUFJ on original problem"); @@ -653,6 +658,16 @@ mip_solution_t solve_mip_helper(optimization_problem_t& op_p CUOPT_LOG_INFO("Writing presolved problem to file: %s", settings.presolve_file.c_str()); presolve_result_opt->reduced_problem.write_to_mps(settings.presolve_file); } + + if (cuopt::cancel_flag_set(settings.cancel_requested)) { + CUOPT_LOG_INFO("Solve cancelled during or after presolve"); + solver_stats_t stats{}; + stats.total_solve_time = timer.elapsed_time(); + stats.presolve_time = presolve_time; + return mip_solution_t( + mip_termination_status_t::Cancelled, stats, op_problem.get_handle_ptr()->get_stream()); + } + // early_best_user_obj is in user-space. // run_mip_solver stores it in context.initial_upper_bound and converts to target spaces as // needed. @@ -668,7 +683,8 @@ mip_solution_t solve_mip_helper(optimization_problem_t& op_p if (run_presolve) { auto status_to_skip = sol.get_termination_status() == mip_termination_status_t::TimeLimit || sol.get_termination_status() == mip_termination_status_t::WorkLimit || - sol.get_termination_status() == mip_termination_status_t::Infeasible; + sol.get_termination_status() == mip_termination_status_t::Infeasible || + sol.get_termination_status() == mip_termination_status_t::Cancelled; auto primal_solution = cuopt::device_copy(sol.get_solution(), op_problem.get_handle_ptr()->get_stream()); rmm::device_uvector dual_solution(0, op_problem.get_handle_ptr()->get_stream()); @@ -784,6 +800,14 @@ mip_solution_t solve_mip_helper(optimization_problem_t& op_p sol.write_to_sol_file(settings.sol_file, op_problem.get_handle_ptr()->get_stream()); } + // Cooperative cancel wins over whatever status the solver settled on + // (TimeLimit, FeasibleFound, etc.): callers that set cancel_requested expect + // Cancelled so they can distinguish cancel from a natural early exit. + // Keep the existing solution/stats; only remap the termination status. + if (cuopt::cancel_flag_set(settings.cancel_requested)) { + sol.set_termination_status(mip_termination_status_t::Cancelled); + } + return sol; } catch (const cuopt::logic_error& e) { CUOPT_LOG_ERROR("Error in solve_mip: %s", e.what()); diff --git a/cpp/src/mip_heuristics/solver.cu b/cpp/src/mip_heuristics/solver.cu index b045190e9f..4493711a26 100644 --- a/cpp/src/mip_heuristics/solver.cu +++ b/cpp/src/mip_heuristics/solver.cu @@ -25,6 +25,7 @@ #include #include +#include #include #include @@ -252,6 +253,7 @@ solution_t mip_solver_t::run_solver() if (timer_.check_time_limit()) { CUOPT_LOG_INFO("Time limit reached after presolve"); + if (cuopt::cancel_flag_set(context.settings.cancel_requested)) { sol.set_cancelled(); } context.stats.total_solve_time = timer_.elapsed_time(); context.problem_ptr->post_process_solution(sol); return sol; @@ -261,10 +263,11 @@ solution_t mip_solver_t::run_solver() if (run_presolve && context.problem_ptr->n_integer_vars == 0) { CUOPT_LOG_INFO("Problem reduced to a LP, running concurrent LP"); pdlp_solver_settings_t settings{}; - settings.time_limit = timer_.remaining_time(); - auto lp_timer = timer_t(settings.time_limit); - settings.method = method_t::Concurrent; - settings.presolver = presolver_t::None; + settings.time_limit = timer_.remaining_time(); + settings.cancel_requested = context.settings.cancel_requested; + auto lp_timer = timer_t(settings.time_limit, context.settings.cancel_requested); + settings.method = method_t::Concurrent; + settings.presolver = presolver_t::None; auto opt_sol = solve_lp_with_method(*context.problem_ptr, settings, lp_timer); @@ -346,6 +349,7 @@ solution_t mip_solver_t::run_solver() // Fill in the settings for branch and bound branch_and_bound_settings.time_limit = timer_.get_time_limit(); + branch_and_bound_settings.cancel_requested = context.settings.cancel_requested; branch_and_bound_settings.node_limit = context.settings.node_limit; branch_and_bound_settings.num_threads = std::max(num_threads - 1, 1); branch_and_bound_settings.print_presolve_stats = false; @@ -477,6 +481,7 @@ solution_t mip_solver_t::run_solver() if (timer_.check_time_limit()) { CUOPT_LOG_INFO("Time limit reached during B&B setup"); + if (cuopt::cancel_flag_set(context.settings.cancel_requested)) { sol.set_cancelled(); } context.stats.total_solve_time = timer_.elapsed_time(); context.problem_ptr->post_process_solution(sol); return sol; @@ -497,6 +502,14 @@ solution_t mip_solver_t::run_solver() sol = dm.run_solver(); } // implicit barrier for all tasks created in B&B and heuristics + const bool cancel_requested = cuopt::cancel_flag_set(context.settings.cancel_requested) || + branch_and_bound_status == mip::mip_status_t::CANCELLED; + if (cancel_requested) { + context.preempt_heuristic_solver_.store(true, std::memory_order_release); + dm.population.preempt_heuristic_solver(); + sol.set_cancelled(); + } + if (!context.settings.heuristics_only && branch_and_bound->has_solver_space_incumbent()) { solution_t branch_and_bound_sol(*context.problem_ptr); branch_and_bound_sol.copy_new_assignment(branch_and_bound_solution.x); @@ -505,6 +518,7 @@ solution_t mip_solver_t::run_solver() if (branch_and_bound_sol.get_feasible() && (!sol.get_feasible() || branch_and_bound_sol.get_objective() < sol.get_objective())) { sol = std::move(branch_and_bound_sol); + if (cancel_requested) { sol.set_cancelled(); } } } @@ -527,6 +541,7 @@ solution_t mip_solver_t::run_solver() if (!is_feasible.value(sol.handle_ptr->get_stream())) { CUOPT_LOG_ERROR( "Solution is not feasible due to variable bounds, returning infeasible solution!"); + if (cancel_requested) { sol.set_cancelled(); } context.stats.total_solve_time = timer_.elapsed_time(); context.problem_ptr->post_process_solution(sol); return sol; diff --git a/cpp/src/mip_heuristics/solver_solution.cu b/cpp/src/mip_heuristics/solver_solution.cu index 1997d684dc..10c1c56640 100644 --- a/cpp/src/mip_heuristics/solver_solution.cu +++ b/cpp/src/mip_heuristics/solver_solution.cu @@ -127,6 +127,12 @@ mip_termination_status_t mip_solution_t::get_termination_status() cons return termination_status_; } +template +void mip_solution_t::set_termination_status(mip_termination_status_t termination_status) +{ + termination_status_ = termination_status; +} + template std::string mip_solution_t::get_termination_status_string( mip_termination_status_t termination_status) @@ -139,8 +145,9 @@ std::string mip_solution_t::get_termination_status_string( case mip_termination_status_t::TimeLimit: return "TimeLimit"; case mip_termination_status_t::WorkLimit: return "WorkLimit"; case mip_termination_status_t::Unbounded: return "Unbounded"; - case mip_termination_status_t::UnboundedOrInfeasible: - return "UnboundedOrInfeasible"; + case mip_termination_status_t::UnboundedOrInfeasible: return "UnboundedOrInfeasible"; + case mip_termination_status_t::Cancelled: + return "Cancelled"; // Do not implement default case to trigger compile time error if new enum is added } return std::string(); @@ -258,6 +265,8 @@ void mip_solution_t::log_detailed_summary() const CUOPT_LOG_INFO("No feasible solution was found within the time limit.\n"); break; + case mip_termination_status_t::Cancelled: CUOPT_LOG_INFO("Solve was cancelled.\n"); break; + case mip_termination_status_t::WorkLimit: CUOPT_LOG_INFO("No feasible solution was found within the work limit.\n"); break; diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index 62b50825e7..1ac0bf626a 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -24,6 +24,7 @@ #include #include +#include #include #include @@ -614,12 +615,17 @@ template std::optional> pdlp_solver_t::check_limits( const timer_t& timer) { + // Cancel exits the same way as the time limit; solve_lp / run_concurrent remap the + // final status to Cancelled. + const bool cancelled = cuopt::cancel_flag_set(settings_.cancel_requested); + // Check for time limit - if (time_limit_reached(timer)) { + if (cancelled || time_limit_reached(timer)) { if (settings_.save_best_primal_so_far) { #ifdef PDLP_VERBOSE_MODE RAFT_CUDA_TRY(cudaDeviceSynchronize()); - std::cout << "Time Limit reached, returning best primal so far" << std::endl; + std::cout << (cancelled ? "Cancelled" : "Time Limit reached") + << ", returning best primal so far" << std::endl; #endif return std::move(best_primal_solution_so_far); } @@ -630,7 +636,8 @@ std::optional> pdlp_solver_t #ifdef PDLP_VERBOSE_MODE RAFT_CUDA_TRY(cudaDeviceSynchronize()); - std::cout << "Time Limit reached, returning current solution" << std::endl; + std::cout << (cancelled ? "Cancelled" : "Time Limit reached") << ", returning current solution" + << std::endl; #endif return current_termination_strategy_.fill_return_problem_solution( internal_solver_iterations_, diff --git a/cpp/src/pdlp/solve.cu b/cpp/src/pdlp/solve.cu index 9e54bb1a11..a1e2bd3d87 100644 --- a/cpp/src/pdlp/solve.cu +++ b/cpp/src/pdlp/solve.cu @@ -39,6 +39,7 @@ #include #include #include +#include #include #include @@ -377,6 +378,7 @@ optimization_problem_solution_t convert_dual_simplex_sol( case simplex::lp_status_t::ITERATION_LIMIT: return pdlp_termination_status_t::IterationLimit; case simplex::lp_status_t::CONCURRENT_LIMIT: return pdlp_termination_status_t::ConcurrentLimit; + case simplex::lp_status_t::CANCELLED: return pdlp_termination_status_t::Cancelled; case simplex::lp_status_t::UNBOUNDED_OR_INFEASIBLE: return pdlp_termination_status_t::UnboundedOrInfeasible; default: return pdlp_termination_status_t::NumericalError; @@ -502,6 +504,7 @@ std::tuple, simplex::lp_status_t, f_t, f_t, f_t barrier_settings.time_limit = settings.time_limit; barrier_settings.iteration_limit = settings.iteration_limit; barrier_settings.concurrent_halt = settings.concurrent_halt; + barrier_settings.cancel_requested = settings.cancel_requested; barrier_settings.folding = settings.folding; barrier_settings.augmented = settings.augmented; barrier_settings.dualize = settings.dualize; @@ -544,6 +547,13 @@ std::tuple, simplex::lp_status_t, f_t, f_t, f_t *settings.concurrent_halt = 1; } + status = cuopt::remap_limit_status_if_cancelled(settings.cancel_requested, + status, + simplex::lp_status_t::CANCELLED, + simplex::lp_status_t::TIME_LIMIT, + simplex::lp_status_t::CONCURRENT_LIMIT, + simplex::lp_status_t::ITERATION_LIMIT); + return {std::move(solution), status, timer.elapsed_time(), norm_user_objective, norm_rhs}; } @@ -594,9 +604,10 @@ std::tuple, simplex::lp_status_t, f_t, f_t, f_t f_t norm_rhs = vector_norm2(user_problem.rhs); simplex::simplex_solver_settings_t dual_simplex_settings; - dual_simplex_settings.time_limit = settings.time_limit; - dual_simplex_settings.iteration_limit = settings.iteration_limit; - dual_simplex_settings.concurrent_halt = settings.concurrent_halt; + dual_simplex_settings.time_limit = settings.time_limit; + dual_simplex_settings.iteration_limit = settings.iteration_limit; + dual_simplex_settings.concurrent_halt = settings.concurrent_halt; + dual_simplex_settings.cancel_requested = settings.cancel_requested; if (dual_simplex_settings.concurrent_halt != nullptr) { // Don't show the dual simplex log in concurrent mode. Show the PDLP log instead dual_simplex_settings.log.log = false; @@ -617,6 +628,13 @@ std::tuple, simplex::lp_status_t, f_t, f_t, f_t *settings.concurrent_halt = 1; } + status = cuopt::remap_limit_status_if_cancelled(settings.cancel_requested, + status, + simplex::lp_status_t::CANCELLED, + simplex::lp_status_t::TIME_LIMIT, + simplex::lp_status_t::CONCURRENT_LIMIT, + simplex::lp_status_t::ITERATION_LIMIT); + return {std::move(solution), status, timer.elapsed_time(), norm_user_objective, norm_rhs}; } @@ -845,9 +863,10 @@ optimization_problem_solution_t run_pdlp(mip::problem_t& pro simplex::lp_solution_t initial_solution(1, 1); translate_to_crossover_problem(problem, sol, lp, initial_solution); simplex::simplex_solver_settings_t dual_simplex_settings; - dual_simplex_settings.time_limit = settings.time_limit; - dual_simplex_settings.iteration_limit = settings.iteration_limit; - dual_simplex_settings.concurrent_halt = settings.concurrent_halt; + dual_simplex_settings.time_limit = settings.time_limit; + dual_simplex_settings.iteration_limit = settings.iteration_limit; + dual_simplex_settings.concurrent_halt = settings.concurrent_halt; + dual_simplex_settings.cancel_requested = settings.cancel_requested; simplex::lp_solution_t vertex_solution(lp.num_rows, lp.num_cols); std::vector vstatus(lp.num_cols); simplex::crossover_status_t crossover_status = simplex::crossover(lp, @@ -1694,6 +1713,11 @@ optimization_problem_solution_t run_concurrent( f_t end_time = timer.elapsed_time(); CUOPT_LOG_CONDITIONAL_INFO(!settings.inside_mip, "Concurrent time: %.3fs", end_time); + auto remap_pdlp_if_cancelled = [&](optimization_problem_solution_t& sol) { + if (!cuopt::cancel_flag_set(settings.cancel_requested)) { return; } + sol.set_termination_status(pdlp_termination_status_t::Cancelled); + }; + const auto dual_simplex_status = !settings.inside_mip ? std::get<1>(*sol_dual_simplex_ptr) : simplex::lp_status_t::CONCURRENT_LIMIT; const auto barrier_status = @@ -1717,6 +1741,7 @@ optimization_problem_solution_t run_concurrent( CUOPT_LOG_CONDITIONAL_INFO(!settings.inside_mip, "Solved with dual simplex"); sol_pdlp.copy_from(problem.handle_ptr, sol_dual_simplex); sol_pdlp.set_solve_time(end_time); + remap_pdlp_if_cancelled(sol_pdlp); CUOPT_LOG_CONDITIONAL_INFO( !settings.inside_mip, "Status: %s Objective: %.8e Iterations: %d Time: %.3fs", @@ -1748,6 +1773,7 @@ optimization_problem_solution_t run_concurrent( CUOPT_LOG_CONDITIONAL_INFO(!settings.inside_mip, "Solved with barrier"); sol_pdlp.copy_from(problem.handle_ptr, sol_barrier); sol_pdlp.set_solve_time(end_time); + remap_pdlp_if_cancelled(sol_pdlp); CUOPT_LOG_CONDITIONAL_INFO( !settings.inside_mip, "Status: %s Objective: %.8e Iterations: %d Time: %.3fs", @@ -1770,6 +1796,7 @@ optimization_problem_solution_t run_concurrent( sol_dual_simplex_ptr.reset(); sol_barrier_ptr.reset(); CUOPT_LOG_CONDITIONAL_INFO(!settings.inside_mip, "Solved with PDLP"); + remap_pdlp_if_cancelled(sol_pdlp); return sol_pdlp; } else if (!settings.inside_mip && sol_pdlp.get_termination_status() == pdlp_termination_status_t::ConcurrentLimit) { @@ -1784,11 +1811,13 @@ optimization_problem_solution_t run_concurrent( method_t::DualSimplex); sol_dual_simplex_ptr.reset(); CUOPT_LOG_CONDITIONAL_INFO(!settings.inside_mip, "Using dual simplex solve info"); + remap_pdlp_if_cancelled(sol_dual_simplex); return sol_dual_simplex; } else { sol_dual_simplex_ptr.reset(); sol_barrier_ptr.reset(); CUOPT_LOG_CONDITIONAL_INFO(!settings.inside_mip, "Using PDLP solve info"); + remap_pdlp_if_cancelled(sol_pdlp); return sol_pdlp; } } @@ -2022,7 +2051,7 @@ optimization_problem_solution_t solve_lp( } validate_new_bounds(op_problem, settings); - auto lp_timer = cuopt::timer_t(settings.time_limit); + auto lp_timer = cuopt::timer_t(settings.time_limit, settings.cancel_requested); std::optional> problem; // handle default presolve if (settings.presolver == presolver_t::Default) { @@ -2175,6 +2204,10 @@ optimization_problem_solution_t solve_lp( solution.write_to_sol_file(settings.sol_file, op_problem.get_handle_ptr()->get_stream()); } + if (cuopt::cancel_flag_set(settings.cancel_requested)) { + solution.set_termination_status(pdlp_termination_status_t::Cancelled); + } + return solution; } catch (const cuopt::logic_error& e) { CUOPT_LOG_ERROR("Error in solve_lp: %s", e.what()); diff --git a/cpp/src/pdlp/solver_solution.cu b/cpp/src/pdlp/solver_solution.cu index 08e5ee00a8..517547a275 100644 --- a/cpp/src/pdlp/solver_solution.cu +++ b/cpp/src/pdlp/solver_solution.cu @@ -318,6 +318,7 @@ std::string optimization_problem_solution_t::get_termination_status_st case pdlp_termination_status_t::PrimalFeasible: return "Primal Feasible"; case pdlp_termination_status_t::ConcurrentLimit: return "Concurrent Limit"; case pdlp_termination_status_t::UnboundedOrInfeasible: return "UnboundedOrInfeasible"; + case pdlp_termination_status_t::Cancelled: return "Cancelled"; case pdlp_termination_status_t::NoTermination: return "NoTermination"; // Do not implement default case to trigger compile time error if new enum is added diff --git a/cpp/src/utilities/solve_limits.hpp b/cpp/src/utilities/solve_limits.hpp new file mode 100644 index 0000000000..305ddae7f4 --- /dev/null +++ b/cpp/src/utilities/solve_limits.hpp @@ -0,0 +1,193 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ + +#pragma once + +#include +#include "timer.hpp" + +#include +#include +#include +#include +#include + +namespace cuopt { + +/** + * @brief Unified early-exit reasons polled by solvers (alongside optimality checks). + * + * Priority when multiple apply: Cancelled > ConcurrentHalt > TimeLimit > IterationLimit. + */ +enum class solve_limit_reason_t : int8_t { + None = 0, + Cancelled = 1, + ConcurrentHalt = 2, + TimeLimit = 3, + IterationLimit = 4, +}; + +namespace detail { + +inline bool atomic_flag_set(const std::atomic* flag) noexcept +{ + return flag != nullptr && flag->load(std::memory_order_acquire); +} + +inline bool concurrent_halt_set(const std::atomic* halt) noexcept +{ + return halt != nullptr && halt->load(std::memory_order_acquire) == 1; +} + +} // namespace detail + +/** + * @brief Check cancel / concurrent halt / time / iteration budgets. + * + * Timer overload used by PDLP and MIP heuristics (`timer_t`). + */ +inline solve_limit_reason_t check_solve_limits( + const timer_t& timer, + const std::atomic* cancel_requested = nullptr, + const std::atomic* concurrent_halt = nullptr, + std::optional iterations = std::nullopt, + std::optional iteration_limit = std::nullopt) noexcept +{ + const std::atomic* cancel = + cancel_requested != nullptr ? cancel_requested : timer.get_cancel_requested(); + if (detail::atomic_flag_set(cancel)) { return solve_limit_reason_t::Cancelled; } + if (detail::concurrent_halt_set(concurrent_halt)) { return solve_limit_reason_t::ConcurrentHalt; } + if (timer.time_exhausted()) { return solve_limit_reason_t::TimeLimit; } + if (iterations.has_value() && iteration_limit.has_value() && + iterations.value() >= iteration_limit.value()) { + return solve_limit_reason_t::IterationLimit; + } + return solve_limit_reason_t::None; +} + +/** + * @brief Same checks using tic/toc wall-clock style (barrier, dual simplex, B&B). + * + * @param elapsed_seconds Elapsed time, e.g. `toc(start_time)`. + * @param time_limit Wall-clock budget in seconds. + */ +inline solve_limit_reason_t check_solve_limits( + double elapsed_seconds, + double time_limit, + const std::atomic* cancel_requested = nullptr, + const std::atomic* concurrent_halt = nullptr, + std::optional iterations = std::nullopt, + std::optional iteration_limit = std::nullopt) noexcept +{ + if (detail::atomic_flag_set(cancel_requested)) { return solve_limit_reason_t::Cancelled; } + if (detail::concurrent_halt_set(concurrent_halt)) { return solve_limit_reason_t::ConcurrentHalt; } + if (elapsed_seconds > time_limit) { return solve_limit_reason_t::TimeLimit; } + if (iterations.has_value() && iteration_limit.has_value() && + iterations.value() >= iteration_limit.value()) { + return solve_limit_reason_t::IterationLimit; + } + return solve_limit_reason_t::None; +} + +inline bool cancel_or_halt_requested(const std::atomic* cancel = nullptr, + const std::atomic* halt = nullptr) noexcept +{ + return detail::atomic_flag_set(cancel) || detail::concurrent_halt_set(halt); +} + +/** + * @brief Convenience bool for loops that previously used only `timer.check_time_limit()`. + * + * If `cancel` is null, uses the cancel pointer embedded on `timer` (if any). + */ +inline bool solve_limit_reached(const timer_t& timer, + const std::atomic* cancel = nullptr, + const std::atomic* halt = nullptr) noexcept +{ + return check_solve_limits(timer, cancel, halt) != solve_limit_reason_t::None; +} + +inline bool solve_limit_reached(double elapsed_seconds, + double time_limit, + const std::atomic* cancel = nullptr, + const std::atomic* halt = nullptr) noexcept +{ + return check_solve_limits(elapsed_seconds, time_limit, cancel, halt) != + solve_limit_reason_t::None; +} + +/** Map to public LP/MIP termination constant values where defined. */ +inline int solve_limit_to_termination_status(solve_limit_reason_t reason) noexcept +{ + switch (reason) { + case solve_limit_reason_t::Cancelled: return CUOPT_TERMINATION_STATUS_CANCELLED; + case solve_limit_reason_t::ConcurrentHalt: return CUOPT_TERMINATION_STATUS_CONCURRENT_LIMIT; + case solve_limit_reason_t::TimeLimit: return CUOPT_TERMINATION_STATUS_TIME_LIMIT; + case solve_limit_reason_t::IterationLimit: return CUOPT_TERMINATION_STATUS_ITERATION_LIMIT; + case solve_limit_reason_t::None: return CUOPT_TERMINATION_STATUS_NO_TERMINATION; + } + return CUOPT_TERMINATION_STATUS_NO_TERMINATION; +} + +inline bool cancel_flag_set(const std::atomic* cancel) noexcept +{ + return detail::atomic_flag_set(cancel); +} + +/** + * @brief Remap limit-like statuses to Cancelled when the cancel flag is set. + * + * Call at solution finalization so inner loops that collapse cancel into + * TIME_LIMIT / CONCURRENT_LIMIT still surface Cancelled to callers. + */ +template +inline Status remap_limit_status_if_cancelled(const std::atomic* cancel, + Status status, + Status cancelled_status, + Status time_limit_status, + Status concurrent_limit_status) noexcept +{ + if (!detail::atomic_flag_set(cancel)) { return status; } + if (status == time_limit_status || status == concurrent_limit_status || + status == cancelled_status) { + return cancelled_status; + } + return status; +} + +template +inline Status remap_limit_status_if_cancelled(const std::atomic* cancel, + Status status, + Status cancelled_status, + Status time_limit_status, + Status concurrent_limit_status, + Status iteration_limit_status) noexcept +{ + if (!detail::atomic_flag_set(cancel)) { return status; } + if (status == time_limit_status || status == concurrent_limit_status || + status == cancelled_status || status == iteration_limit_status) { + return cancelled_status; + } + return status; +} + +/** + * @brief Optional helper to clear a cancel flag after a top-level solve returns. + * + * gRPC clears via job-slot reset instead so the worker can still observe the + * flag after solve_* returns. Standalone callers may use this RAII guard. + */ +struct cancel_flag_clear_on_exit_t { + std::atomic* flag{nullptr}; + explicit cancel_flag_clear_on_exit_t(std::atomic* f) noexcept : flag(f) {} + cancel_flag_clear_on_exit_t(const cancel_flag_clear_on_exit_t&) = delete; + cancel_flag_clear_on_exit_t& operator=(const cancel_flag_clear_on_exit_t&) = delete; + ~cancel_flag_clear_on_exit_t() + { + if (flag != nullptr) { flag->store(false, std::memory_order_release); } + } +}; + +} // namespace cuopt diff --git a/cpp/src/utilities/timer.hpp b/cpp/src/utilities/timer.hpp index b7ab6a63bd..311c5d9e20 100644 --- a/cpp/src/utilities/timer.hpp +++ b/cpp/src/utilities/timer.hpp @@ -6,6 +6,7 @@ /* clang-format on */ #pragma once +#include #include #include @@ -19,10 +20,10 @@ class timer_t { public: timer_t() = delete; timer_t(const timer_t&) = default; - timer_t(double time_limit_) + timer_t(double time_limit_) : timer_t(time_limit_, nullptr) {} + timer_t(double time_limit_, std::atomic* cancel_requested) + : time_limit(time_limit_), begin(steady_clock::now()), cancel_requested_(cancel_requested) { - time_limit = time_limit_; - begin = steady_clock::now(); } void print_debug(std::string msg) const @@ -34,7 +35,20 @@ class timer_t { elapsed_time()); } - bool check_time_limit() const noexcept { return elapsed_time() >= time_limit; } + /** True if cancel was requested or the wall-clock budget is exhausted. */ + bool check_time_limit() const noexcept + { + if (cancel_requested()) { return true; } + return time_exhausted(); + } + + /** Wall-clock budget only (ignores cancel). */ + bool time_exhausted() const noexcept { return elapsed_time() >= time_limit; } + + bool cancel_requested() const noexcept + { + return cancel_requested_ != nullptr && cancel_requested_->load(std::memory_order_acquire); + } bool check_half_time() const noexcept { return elapsed_time() >= time_limit / 2; } @@ -55,6 +69,8 @@ class timer_t { double get_time_limit() const noexcept { return time_limit; } + std::atomic* get_cancel_requested() const noexcept { return cancel_requested_; } + double get_tic_start() const noexcept { /** @@ -87,6 +103,7 @@ class timer_t { private: double time_limit; steady_clock::time_point begin; + std::atomic* cancel_requested_{nullptr}; }; } // namespace cuopt diff --git a/cpp/tests/mip/CMakeLists.txt b/cpp/tests/mip/CMakeLists.txt index 91fb87a953..20eacd7052 100644 --- a/cpp/tests/mip/CMakeLists.txt +++ b/cpp/tests/mip/CMakeLists.txt @@ -18,3 +18,6 @@ ConfigureTest(DOC_EXAMPLE_TEST ConfigureTest(HEURISTICS_HYPER_PARAMS_TEST ${CMAKE_CURRENT_SOURCE_DIR}/heuristics_hyper_params_test.cu LABELS numopt) +ConfigureTest(COOPERATIVE_CANCEL_TEST + ${CMAKE_CURRENT_SOURCE_DIR}/cooperative_cancel_test.cu + LABELS numopt) diff --git a/cpp/tests/mip/cooperative_cancel_test.cu b/cpp/tests/mip/cooperative_cancel_test.cu new file mode 100644 index 0000000000..f3e170c93b --- /dev/null +++ b/cpp/tests/mip/cooperative_cancel_test.cu @@ -0,0 +1,188 @@ +/* clang-format off */ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: Apache-2.0 + */ +/* clang-format on */ + +#include "../linear_programming/utilities/pdlp_test_utilities.cuh" + +#include +#include +#include +#include +#include +#include + +#include + +#include + +#include +#include +#include +#include +#include +#include +#include + +namespace cuopt::mathematical_optimization::test { + +namespace { + +using clock = std::chrono::steady_clock; + +double elapsed_seconds(clock::time_point t0) +{ + return std::chrono::duration(clock::now() - t0).count(); +} + +bool file_exists(const std::string& rel_path) +{ + return std::filesystem::exists(make_path_absolute(rel_path)); +} + +} // namespace + +// Mid-solve cancel on local MIP solves across several MIPLIB-style instances. +// Each instance gets a long time limit so cancel (not TL) is the expected exit. +TEST(cooperative_cancel, mip_mid_solve_cancel_loop) +{ + constexpr double long_time_limit_s = 300.0; + constexpr double cancel_after_s = 2.0; + // Allow Papilo / early FJ / root LP some room after the flag flips. + constexpr double max_total_s = 90.0; + + const std::vector instances = { + "mip/neos5.mps", + "mip/gen-ip054.mps", + "mip/swath1.mps", + "mip/seymour1.mps", + "mip/ns1208400.mps", + "mip/rmatr200-p5.mps", + }; + + int ran = 0; + for (const auto& rel : instances) { + if (!file_exists(rel)) { + GTEST_LOG_(WARNING) << "Skipping missing dataset " << rel; + continue; + } + ++ran; + + const raft::handle_t handle{}; + auto path = make_path_absolute(rel); + auto problem = io::read_mps(path, false); + handle.sync_stream(); + + mip_solver_settings_t settings; + settings.time_limit = long_time_limit_s; + settings.log_to_console = false; + std::atomic cancel{false}; + settings.cancel_requested = &cancel; + + std::optional> solution; + const auto t0 = clock::now(); + std::thread worker([&] { + solution = solve_mip(&handle, problem, settings); + handle.sync_stream(); + }); + + std::this_thread::sleep_for(std::chrono::duration(cancel_after_s)); + cancel.store(true, std::memory_order_release); + worker.join(); + const double secs = elapsed_seconds(t0); + + ASSERT_TRUE(solution.has_value()) << rel << " produced no solution object"; + EXPECT_GE(secs, cancel_after_s * 0.9) << rel << " finished before cancel was set"; + EXPECT_LT(secs, max_total_s) << rel << " did not unwind promptly after cancel (" << secs + << "s)"; + EXPECT_EQ(solution->get_termination_status(), mip_termination_status_t::Cancelled) + << rel << " status=" << static_cast(solution->get_termination_status()) << " after " + << secs << "s"; + } + ASSERT_GE(ran, 3) << "Need at least a few MIP datasets under RAPIDS_DATASET_ROOT_DIR"; +} + +// Pre-armed cancel should not burn the wall-clock budget. +TEST(cooperative_cancel, mip_preset_cancel_exits_early) +{ + constexpr double long_time_limit_s = 300.0; + + const std::string rel = "mip/seymour1.mps"; + if (!file_exists(rel)) { GTEST_SKIP() << "Missing " << rel; } + + const raft::handle_t handle{}; + auto problem = io::read_mps(make_path_absolute(rel), false); + handle.sync_stream(); + + mip_solver_settings_t settings; + settings.time_limit = long_time_limit_s; + settings.log_to_console = false; + std::atomic cancel{true}; + settings.cancel_requested = &cancel; + + const auto t0 = clock::now(); + auto solution = solve_mip(&handle, problem, settings); + handle.sync_stream(); + const double secs = elapsed_seconds(t0); + + EXPECT_LT(secs, 60.0) << "pre-set cancel should exit early, took " << secs << "s"; + EXPECT_EQ(solution.get_termination_status(), mip_termination_status_t::Cancelled); +} + +// Local LP cancel on RAPIDS datasets that stay busy under a long time limit. +// Easy Optinals (savsched1/ex10/...) finish before cancel and are omitted. +TEST(cooperative_cancel, lp_mid_solve_cancel_loop) +{ + constexpr double long_time_limit_s = 300.0; + constexpr double cancel_after_s = 2.0; + constexpr double max_total_s = 90.0; + + const std::vector instances = { + "linear_programming/scpm1/scpm1.mps", + }; + + int ran = 0; + for (const auto& rel : instances) { + if (!file_exists(rel)) { + GTEST_LOG_(WARNING) << "Skipping missing LP " << rel; + continue; + } + ++ran; + + const raft::handle_t handle{}; + auto problem = io::read_mps(make_path_absolute(rel), false); + handle.sync_stream(); + + pdlp_solver_settings_t settings; + settings.time_limit = long_time_limit_s; + settings.log_to_console = false; + std::atomic cancel{false}; + settings.cancel_requested = &cancel; + + std::optional> solution; + const auto t0 = clock::now(); + std::thread worker([&] { + solution = solve_lp(&handle, problem, settings); + handle.sync_stream(); + }); + + std::this_thread::sleep_for(std::chrono::duration(cancel_after_s)); + cancel.store(true, std::memory_order_release); + worker.join(); + const double secs = elapsed_seconds(t0); + + ASSERT_TRUE(solution.has_value()) << rel << " produced no solution object"; + EXPECT_GE(secs, cancel_after_s * 0.9) << rel << " finished before cancel was set"; + EXPECT_LT(secs, max_total_s) << rel << " did not unwind promptly after cancel (" << secs + << "s)"; + EXPECT_EQ(solution->get_termination_status(), pdlp_termination_status_t::Cancelled) + << rel << " status=" << static_cast(solution->get_termination_status()) << " after " + << secs << "s"; + } + if (ran == 0) { GTEST_SKIP() << "No hard LP datasets found under RAPIDS_DATASET_ROOT_DIR"; } + ASSERT_GE(ran, 1) << "Need at least one hard LP for cancel coverage"; +} + +} // namespace cuopt::mathematical_optimization::test diff --git a/python/cuopt/cuopt/tests/fixtures/grpc_server_fixtures.py b/python/cuopt/cuopt/tests/fixtures/grpc_server_fixtures.py index a37a6a43d4..15c5b08bbe 100644 --- a/python/cuopt/cuopt/tests/fixtures/grpc_server_fixtures.py +++ b/python/cuopt/cuopt/tests/fixtures/grpc_server_fixtures.py @@ -1,4 +1,4 @@ -# SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 """ @@ -35,6 +35,7 @@ def test_foo(self, grpc_server): GRPC_PORT_OFFSET_CLIENT = 800 GRPC_PORT_OFFSET_TLS = 850 GRPC_PORT_OFFSET_MTLS = 900 +GRPC_PORT_OFFSET_COOP_CANCEL = 950 def find_grpc_server(): @@ -265,23 +266,30 @@ def start_tls_grpc_server(port_offset, cert_dir, require_client_cert=False): return proc, port -def start_grpc_server(port_offset): - """Locate the server, start it on BASE + port_offset, return (proc, client_env).""" +def start_grpc_server(port_offset, server_log_path=None, workers=1): + """Locate the server, start it on BASE + port_offset, return (proc, client_env). + + When ``server_log_path`` is set, pass ``--server-log`` so tests can assert + on cooperative-cancel / worker-restart lines. + """ server_bin = find_grpc_server() if server_bin is None: pytest.skip("cuopt_grpc_server not found") port = int(os.environ.get("CUOPT_TEST_PORT_BASE", "18000")) + port_offset client_env = client_remote_env(port) + cmd = [ + server_bin, + "--port", + str(port), + "--workers", + str(workers), + "--log-to-console", + ] + if server_log_path is not None: + cmd.extend(["--server-log", str(server_log_path)]) proc = spawn_server( - [ - server_bin, - "--port", - str(port), - "--workers", - "1", - "--log-to-console", - ], + cmd, env=server_env(), ) time.sleep(0.5) diff --git a/python/cuopt/cuopt/tests/linear_programming/grpc_lifecycle_stress.py b/python/cuopt/cuopt/tests/linear_programming/grpc_lifecycle_stress.py new file mode 100644 index 0000000000..4fc1a5c0d9 --- /dev/null +++ b/python/cuopt/cuopt/tests/linear_programming/grpc_lifecycle_stress.py @@ -0,0 +1,1790 @@ +#!/usr/bin/env python3 +# SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +r""" +Stress harness for cuopt_grpc_server cancel / delete / controlled-shutdown. + +Starts a multi-worker gRPC server, submits LP/MIP jobs of mixed sizes with +jittered timing, then repeatedly: + + 1. Cancels a PROCESSING job — expects CANCELLED status for that job_id, + result() failure, server log ``Marked job CANCELLED``, and a healthy + follow-up solve (cooperative cancel preferred; kill+restart is fallback). + 2. Deletes a PROCESSING job — expects NOT_FOUND for that job_id (delete + includes cancel), result() failure, cancel-start log for that job_id, + and a healthy follow-up solve. + 3. Fetches results via Client.result() (unified GetResult, no client-side + is_mip) for both a completed LP and a completed MIP, checking objectives. + 4. Sends SIGINT with jobs mid-solve — expects the server process to exit + promptly (no intermediate-log requirement). + +Example (from repo root, with the cuOpt env active):: + + python python/cuopt/cuopt/tests/linear_programming/grpc_lifecycle_stress.py \ + --workers 2 --loops 12 --port 15051 + +Environment: + CUOPT_GRPC_SERVER_PATH Override path to cuopt_grpc_server + RAPIDS_DATASET_ROOT_DIR Dataset root (default: /datasets or ./datasets) +""" + +from __future__ import annotations + +import argparse +import os +import random +import re +import shutil +import signal +import socket +import subprocess +import sys +import threading +import time +from dataclasses import dataclass, field +from pathlib import Path +from typing import Optional + +from cuopt.grpc.linear_programming import Client, GrpcError, JobStatus +from cuopt.linear_programming import Read, SolverSettings +from cuopt.linear_programming.problem import INTEGER, MAXIMIZE, Problem +from cuopt.linear_programming.solver.solver_parameters import CUOPT_TIME_LIMIT + +# --------------------------------------------------------------------------- +# Paths / constants +# --------------------------------------------------------------------------- + +_REPO_ROOT = Path(__file__).resolve().parents[5] +_DEFAULT_DATASETS = _REPO_ROOT / "datasets" +if not _DEFAULT_DATASETS.is_dir(): + _DEFAULT_DATASETS = Path.cwd() / "datasets" + +RESTARTED_WORKER_RE = re.compile( + r"Restarted worker\s+(\d+)\s+with PID\s+(\d+)" +) +# Legacy immediate-kill cancel (pre-cooperative). +CANCEL_KILL_RE = re.compile( + r"Cancelling running job\s+(\S+)\s+\(killing worker\s+(\d+)\)" +) +# Cooperative cancel: SHM flag set; worker may unwind without SIGKILL. +CANCEL_COOP_RE = re.compile( + r"Cancelling running job\s+(\S+)\s+\(cooperative cancel; worker\s+(\d+)\)" +) +# Fallback after cooperative grace period. +CANCEL_FALLBACK_KILL_RE = re.compile( + r"Job\s+(\S+)\s+still running after cooperative cancel grace;.*" + r"SIGKILL to worker\s+(\d+)" +) +# Result-path confirmation that the job was recorded as cancelled. +JOB_TRACKER_CANCELLED_RE = re.compile( + r"Marked job CANCELLED in job_tracker:\s+(\S+)\s+msg=" +) +SHUTDOWN_SIGNAL_RE = re.compile(r"Shutdown signal\s+(\d+)\s+received") +WORKER_KILLED_RE = re.compile(r"Worker\s+(\d+)\s+killed by signal\s+(\d+)") +WORKER_PROCESSING_RE = re.compile( + r"\[Worker\s+(\d+)\]\s+Processing job:\s+(\S+)" +) +WORKER_COMPLETED_RE = re.compile( + r"\[Worker\s+(\d+)\]\s+Completed job:\s+(\S+)" +) + + +def note_unexpected_worker_segfault(server_log_text: str) -> str: + """Diagnostic note when a worker dies with SIGSEGV during cooperative cancel.""" + notes = [] + for m in WORKER_KILLED_RE.finditer(server_log_text): + sig = int(m.group(2)) + if sig == 11: + notes.append(f"unexpected worker SIGSEGV pid={m.group(1)}") + return "; ".join(notes) + + +# Solver progress lines that mean the MIP/LP is actively producing intermediates +# (not just the startup / scaling banner). MIP emits +# "New solution from ... Objective +N ..."; LP/PDLP emits iteration lines. +INTERMEDIATE_LOG_RE = re.compile( + r"(New solution from|Incumbent|Best bound|Best objective|PDLP iteration|Iteration\s+\d+)", + re.IGNORECASE, +) + + +@dataclass +class ProblemSpec: + name: str + path: Optional[Path] = None + kind: str = "mip" # "mip" | "lp" | "synth" + time_limit: float = 120.0 + synth_factory: Optional[str] = None + + +@dataclass +class TrialResult: + action: str + ok: bool + detail: str + elapsed_s: float = 0.0 + worker_restarted: Optional[bool] = None + client_status_ok: Optional[bool] = None + server_exited: Optional[bool] = None + mid_job_at_signal: Optional[bool] = None + + +@dataclass +class Summary: + trials: list[TrialResult] = field(default_factory=list) + + def add(self, trial: TrialResult) -> None: + self.trials.append(trial) + flag = "PASS" if trial.ok else "FAIL" + extras = [] + if trial.worker_restarted is not None: + extras.append(f"worker_restarted={trial.worker_restarted}") + if trial.client_status_ok is not None: + extras.append(f"client_status_ok={trial.client_status_ok}") + if trial.server_exited is not None: + extras.append(f"server_exited={trial.server_exited}") + if trial.mid_job_at_signal is not None: + extras.append(f"mid_job={trial.mid_job_at_signal}") + extra = (", " + ", ".join(extras)) if extras else "" + log( + f"[{flag}] {trial.action}: {trial.detail} " + f"({trial.elapsed_s:.2f}s{extra})" + ) + + def print_report(self) -> int: + total = len(self.trials) + passed = sum(1 for t in self.trials if t.ok) + failed = total - passed + by_action: dict[str, list[TrialResult]] = {} + for t in self.trials: + by_action.setdefault(t.action, []).append(t) + + log("=" * 72) + log(f"SUMMARY: {passed}/{total} passed, {failed} failed") + for action, items in sorted(by_action.items()): + a_pass = sum(1 for t in items if t.ok) + restarts = [ + t.worker_restarted + for t in items + if t.worker_restarted is not None + ] + exits = [ + t.server_exited for t in items if t.server_exited is not None + ] + mid = [ + t.mid_job_at_signal + for t in items + if t.mid_job_at_signal is not None + ] + bits = [f"{a_pass}/{len(items)} ok"] + if restarts: + bits.append(f"restarts={sum(restarts)}/{len(restarts)}") + if exits: + bits.append(f"exits={sum(exits)}/{len(exits)}") + if mid: + bits.append(f"signaled_mid_job={sum(mid)}/{len(mid)}") + log(f" {action}: " + ", ".join(bits)) + log("=" * 72) + return 0 if failed == 0 else 1 + + +_log_lock = threading.Lock() + + +def log(msg: str) -> None: + with _log_lock: + print(f"{time.strftime('%H:%M:%S')} {msg}", flush=True) + + +# --------------------------------------------------------------------------- +# Server process helpers +# --------------------------------------------------------------------------- + + +def find_grpc_server() -> Optional[str]: + env_path = os.environ.get("CUOPT_GRPC_SERVER_PATH") + if env_path and os.path.isfile(env_path) and os.access(env_path, os.X_OK): + return env_path + found = shutil.which("cuopt_grpc_server") + if found: + return found + candidates = [ + _REPO_ROOT / "cpp" / "build" / "cuopt_grpc_server", + _REPO_ROOT / ".cuopt_env" / "bin" / "cuopt_grpc_server", + Path(os.environ.get("CONDA_PREFIX", "")) / "bin" / "cuopt_grpc_server", + ] + for c in candidates: + if c.is_file() and os.access(c, os.X_OK): + return str(c) + return None + + +def wait_for_port(port: int, timeout: float = 30.0) -> bool: + deadline = time.monotonic() + timeout + while time.monotonic() < deadline: + try: + with socket.create_connection(("127.0.0.1", port), timeout=1): + return True + except OSError: + time.sleep(0.2) + return False + + +def wait_for_client(port: int, timeout: float = 45.0) -> bool: + deadline = time.monotonic() + timeout + while time.monotonic() < deadline: + if not wait_for_port(port, timeout=1): + time.sleep(0.2) + continue + try: + client = Client("localhost", port) + del client + return True + except GrpcError: + time.sleep(0.2) + return False + + +def _set_pdeathsig() -> None: + try: + import ctypes + + ctypes.CDLL("libc.so.6", use_errno=True).prctl(1, signal.SIGKILL) + except Exception: + pass + + +class ServerHandle: + def __init__( + self, + workers: int, + port: int, + server_log: Path, + max_message_mb: int = 256, + ): + self.workers = workers + self.port = port + self.server_log = server_log + self.max_message_mb = max_message_mb + self.proc: Optional[subprocess.Popen] = None + self._log_pos = 0 + + @property + def pid(self) -> Optional[int]: + return None if self.proc is None else self.proc.pid + + def is_running(self) -> bool: + return self.proc is not None and self.proc.poll() is None + + def start(self) -> None: + if self.is_running(): + return + binary = find_grpc_server() + if binary is None: + raise RuntimeError( + "cuopt_grpc_server not found; set CUOPT_GRPC_SERVER_PATH or " + "activate the cuOpt env" + ) + self.server_log.parent.mkdir(parents=True, exist_ok=True) + # Truncate so each start has a clean scan window. + self.server_log.write_text("") + self._log_pos = 0 + env = os.environ.copy() + for key in list(env): + if key.startswith("CUOPT_TLS_") or key.startswith("CUOPT_REMOTE_"): + env.pop(key) + cmd = [ + binary, + "--port", + str(self.port), + "--workers", + str(self.workers), + "--max-message-mb", + str(self.max_message_mb), + "--server-log", + str(self.server_log), + "--log-to-console", + ] + log(f"Starting server: {' '.join(cmd)}") + self.proc = subprocess.Popen( + cmd, + stdout=subprocess.PIPE, + stderr=subprocess.STDOUT, + env=env, + start_new_session=True, + preexec_fn=_set_pdeathsig, + text=True, + bufsize=1, + ) + # Drain stdout in background so the pipe cannot fill; operational + # detail also lives in --server-log. + threading.Thread( + target=self._drain_stdout, name="server-stdout", daemon=True + ).start() + time.sleep(0.5) + if self.proc.poll() is not None: + raise RuntimeError( + f"cuopt_grpc_server exited immediately (rc={self.proc.returncode})" + ) + if not wait_for_client(self.port, timeout=45): + self.force_kill() + raise RuntimeError("gRPC server did not become ready in time") + log( + f"Server ready pid={self.proc.pid} port={self.port} workers={self.workers}" + ) + + def _drain_stdout(self) -> None: + assert self.proc is not None and self.proc.stdout is not None + try: + for line in self.proc.stdout: + # Mirror a few high-signal lines to the harness log. + if any( + s in line + for s in ( + "Restarted worker", + "Cancelling running", + "Shutdown signal", + "killed by signal", + "Using CUDA device", + ) + ): + log(f"[server] {line.rstrip()}") + except Exception: + pass + + def mark_log(self) -> int: + """Remember current end of server log for later delta scans.""" + try: + self._log_pos = self.server_log.stat().st_size + except FileNotFoundError: + self._log_pos = 0 + return self._log_pos + + def read_log_delta(self) -> str: + try: + with open( + self.server_log, "r", encoding="utf-8", errors="replace" + ) as f: + f.seek(self._log_pos) + return f.read() + except FileNotFoundError: + return "" + + def wait_log_match( + self, pattern: re.Pattern, timeout: float = 15.0 + ) -> Optional[re.Match]: + deadline = time.monotonic() + timeout + while time.monotonic() < deadline: + text = self.read_log_delta() + m = pattern.search(text) + if m: + return m + if not self.is_running() and pattern is not SHUTDOWN_SIGNAL_RE: + # Process died unexpectedly during a non-shutdown wait. + text = self.read_log_delta() + return pattern.search(text) + time.sleep(0.1) + return pattern.search(self.read_log_delta()) + + def send_sigint(self) -> None: + if self.proc is None: + return + log(f"Sending SIGINT to server pid={self.proc.pid}") + os.kill(self.proc.pid, signal.SIGINT) + + def wait_exited(self, timeout: float = 15.0) -> bool: + if self.proc is None: + return True + deadline = time.monotonic() + timeout + while time.monotonic() < deadline: + if self.proc.poll() is not None: + log(f"Server exited rc={self.proc.returncode}") + return True + time.sleep(0.05) + return False + + def force_kill(self) -> None: + if self.proc is None: + return + try: + pgid = os.getpgid(self.proc.pid) + except (ProcessLookupError, OSError): + return + for sig in (signal.SIGTERM, signal.SIGKILL): + try: + os.killpg(pgid, sig) + except (ProcessLookupError, OSError): + return + try: + self.proc.wait(timeout=5) + return + except subprocess.TimeoutExpired: + continue + + +# --------------------------------------------------------------------------- +# Problems +# --------------------------------------------------------------------------- + + +def _synth_tiny_lp() -> Problem: + problem = Problem("stress_tiny_lp") + x = problem.addVariable(lb=0.0, ub=2.0, name="x") + y = problem.addVariable(lb=0.0, name="y") + problem.addConstraint(3 * x + 4 * y <= 5.4, name="c1") + problem.addConstraint(2.7 * x + 10.1 * y <= 4.9, name="c2") + problem.setObjective(0.2 * x + 0.1 * y, sense=MAXIMIZE) + return problem + + +def _synth_small_mip() -> Problem: + problem = Problem("stress_small_mip") + x = problem.addVariable(lb=0, ub=50, vtype=INTEGER, name="x") + y = problem.addVariable(lb=0, ub=50, vtype=INTEGER, name="y") + problem.addConstraint(x + y <= 80, name="c1") + problem.addConstraint(2 * x - y >= 0, name="c2") + problem.setObjective(x + 3 * y, sense=MAXIMIZE) + return problem + + +def _synth_known_mip() -> Problem: + """Same MIP as test_grpc_client.test_mip_submit_and_result (opt = 15).""" + problem = Problem("stress_known_mip") + x = problem.addVariable(lb=0, ub=10, vtype=INTEGER, name="x") + y = problem.addVariable(lb=0, ub=10, vtype=INTEGER, name="y") + problem.addConstraint(x + y <= 10, name="c1") + problem.addConstraint(x - y >= 0, name="c2") + problem.setObjective(x + 2 * y, sense=MAXIMIZE) + return problem + + +def build_problem_pool(datasets: Path) -> list[ProblemSpec]: + mip = datasets / "mip" + lp = datasets / "linear_programming" + specs = [ + ProblemSpec( + "synth_tiny_lp", + kind="synth", + time_limit=5.0, + synth_factory="tiny_lp", + ), + ProblemSpec( + "synth_small_mip", + kind="synth", + time_limit=30.0, + synth_factory="small_mip", + ), + ProblemSpec( + "afiro", path=lp / "afiro_original.mps", kind="lp", time_limit=15.0 + ), + ProblemSpec( + "neos5", + path=mip / "neos5-free-bound.mps", + kind="mip", + time_limit=120.0, + ), + ProblemSpec( + "gen-ip054", + path=mip / "gen-ip054.mps", + kind="mip", + time_limit=60.0, + ), + ProblemSpec( + "swath1", path=mip / "swath1.mps", kind="mip", time_limit=120.0 + ), + ProblemSpec( + "seymour1", path=mip / "seymour1.mps", kind="mip", time_limit=120.0 + ), + ProblemSpec( + "ns1208400", + path=mip / "ns1208400.mps", + kind="mip", + time_limit=120.0, + ), + ProblemSpec( + "rmatr200-p5", + path=mip / "rmatr200-p5.mps", + kind="mip", + time_limit=180.0, + ), + ] + available = [] + for s in specs: + if s.kind == "synth": + available.append(s) + elif s.path is not None and s.path.is_file(): + available.append(s) + else: + log(f"Skipping missing dataset: {s.name} ({s.path})") + if not available: + raise RuntimeError(f"No problems available under {datasets}") + return available + + +def load_problem(spec: ProblemSpec): + if spec.synth_factory == "tiny_lp": + return _synth_tiny_lp() + if spec.synth_factory == "small_mip": + return _synth_small_mip() + if spec.synth_factory == "known_mip": + return _synth_known_mip() + assert spec.path is not None + return Read(str(spec.path)) + + +def make_settings(time_limit: float) -> SolverSettings: + settings = SolverSettings() + settings.set_parameter(CUOPT_TIME_LIMIT, float(time_limit)) + return settings + + +def _approx( + a: float, b: float, rel: float = 1e-3, abs_tol: float = 1e-5 +) -> bool: + scale = max(abs(a), abs(b), 1.0) + return abs(a - b) <= max(abs_tol, rel * scale) + + +def fetch_and_check_result( + client: Client, + problem: Problem, + settings: SolverSettings, + var_names: list[str], + expected_obj: float, + label: str, + timeout: float = 120.0, + server: Optional["ServerHandle"] = None, +) -> tuple[bool, str]: + """ + Submit → wait COMPLETED → Client.result() (unified GetResult, no is_mip) + and verify the primal objective. Exercises LP and MIP through the same API. + """ + job_id = client.submit(problem, settings) + log(f"Result-check submit {label} job_id={job_id}") + try: + terminal = client.wait(job_id, timeout=int(timeout)) + if terminal != JobStatus.COMPLETED: + extra = "" + try: + client.result(job_id) + except GrpcError as e: + extra = f"; result_error={e}" + if server is not None: + for line in server.read_log_delta().splitlines(): + if job_id in line and ( + "FAILED" in line or "Memory" in line + ): + extra += f"; server_log={line.strip()}" + break + return ( + False, + f"{label}: wait returned {terminal.name}, expected COMPLETED{extra}", + ) + + # Unified result path: server response selects LP vs MIP solution. + solution = client.result(job_id, var_names) + if solution is None: + return False, f"{label}: result() returned None after COMPLETED" + + try: + obj = float(solution.get_primal_objective()) + except Exception as e: + return False, f"{label}: get_primal_objective failed: {e}" + + if not _approx(obj, expected_obj): + return ( + False, + f"{label}: objective {obj} != expected {expected_obj}", + ) + + if var_names: + try: + vars_ = solution.get_vars() + except Exception as e: + return False, f"{label}: get_vars failed: {e}" + missing = [n for n in var_names if n not in vars_] + if missing: + return ( + False, + f"{label}: missing vars {missing} in {list(vars_.keys())}", + ) + return True, f"{label}: obj={obj} vars={vars_} job_id={job_id}" + + return True, f"{label}: obj={obj} job_id={job_id}" + except GrpcError as e: + return False, f"{label}: GrpcError: {e}" + finally: + try: + client.delete(job_id) + except GrpcError: + pass + + +def ensure_workers_healthy( + server: ServerHandle, pool: list[ProblemSpec], attempts: int = 2 +) -> bool: + """ + Submit a tiny LP and require COMPLETED. If the GPU is poisoned from a + prior SIGKILL, restart the server and retry. + """ + for attempt in range(1, attempts + 1): + if not server.is_running(): + server.start() + client = make_client(server.port) + ok, detail = fetch_and_check_result( + client, + _synth_tiny_lp(), + make_settings(10.0), + ["x", "y"], + expected_obj=0.36, + label=f"health(attempt={attempt})", + server=server, + ) + if ok: + log(f"Worker health check OK ({detail})") + return True + log(f"Worker health check failed: {detail}") + if attempt < attempts: + log("Restarting server to clear poisoned GPU state") + server.force_kill() + server.wait_exited(timeout=10) + time.sleep(1.0) + server.start() + return False + + +def trial_result( + server: ServerHandle, + datasets: Path, + traffic_gate: threading.Event, +) -> TrialResult: + """ + Positive GetResult coverage for both LP and MIP via Client.result(). + + The recent API removed client-side is_mip; LP vs MIP is taken from the + server response. This trial fails if either flavor cannot be fetched. + """ + traffic_gate.clear() + t0 = time.monotonic() + notes = [] + if not ensure_workers_healthy(server, []): + return TrialResult( + action="result", + ok=False, + detail="workers unhealthy after cancel/delete SIGKILL churn; " + "server restart did not recover", + elapsed_s=time.monotonic() - t0, + client_status_ok=False, + ) + + client = make_client(server.port) + server.mark_log() + + lp_ok, lp_detail = fetch_and_check_result( + client, + _synth_tiny_lp(), + make_settings(10.0), + ["x", "y"], + expected_obj=0.36, + label="LP(synth_tiny_lp)", + server=server, + ) + notes.append(lp_detail) + + mip_ok, mip_detail = fetch_and_check_result( + client, + _synth_known_mip(), + make_settings(30.0), + ["x", "y"], + expected_obj=15.0, + label="MIP(known)", + server=server, + ) + notes.append(mip_detail) + + afiro_path = datasets / "linear_programming" / "afiro_original.mps" + if afiro_path.is_file(): + afiro_ok, afiro_detail = fetch_and_check_result( + client, + Read(str(afiro_path)), + make_settings(30.0), + [], + expected_obj=-464.753, + label="LP(afiro)", + timeout=60.0, + server=server, + ) + notes.append(afiro_detail) + else: + afiro_ok = False + notes.append("LP(afiro): mps not found") + + ok = lp_ok and mip_ok and afiro_ok + return TrialResult( + action="result", + ok=ok, + detail="; ".join(notes), + elapsed_s=time.monotonic() - t0, + client_status_ok=ok, + ) + + +# --------------------------------------------------------------------------- +# Job helpers +# --------------------------------------------------------------------------- + + +def make_client(port: int) -> Client: + return Client("localhost", port) + + +def wait_until_processing( + client: Client, job_id: str, timeout: float = 20.0 +) -> JobStatus: + deadline = time.monotonic() + timeout + last = JobStatus.QUEUED + while time.monotonic() < deadline: + last = client.status(job_id) + if last == JobStatus.PROCESSING: + return last + if last in ( + JobStatus.COMPLETED, + JobStatus.FAILED, + JobStatus.CANCELLED, + JobStatus.NOT_FOUND, + ): + return last + time.sleep(0.05) + return last + + +def submit_spec(client: Client, spec: ProblemSpec) -> str: + problem = load_problem(spec) + job_id = client.submit(problem, make_settings(spec.time_limit)) + log( + f"Submitted {spec.name} ({spec.kind}) job_id={job_id} tl={spec.time_limit}s" + ) + return job_id + + +def submit_long_jobs( + client: Client, + pool: list[ProblemSpec], + n: int, + rng: random.Random, + min_time_limit: float = 60.0, +) -> list[str]: + """Submit n long-running MIP-ish jobs (prefer large time_limit).""" + long_specs = [s for s in pool if s.time_limit >= min_time_limit] or pool + # Prefer real MIP datasets over tiny synth for cancel/shutdown occupancy. + mip_specs = [s for s in long_specs if s.kind == "mip"] or long_specs + job_ids = [] + for _ in range(n): + spec = rng.choice(mip_specs) + job_ids.append(submit_spec(client, spec)) + time.sleep(rng.uniform(0.05, 0.4)) + return job_ids + + +def _tiny_probe_spec(pool: list[ProblemSpec]) -> ProblemSpec: + for s in pool: + if s.synth_factory == "tiny_lp": + return s + short = [s for s in pool if s.time_limit <= 30] or pool + return short[0] + + +def verify_restarted_worker_executes( + server: ServerHandle, + client: Client, + pool: list[ProblemSpec], + num_workers: int, + restart_match: Optional[re.Match], + timeout: float = 60.0, +) -> tuple[bool, str]: + """ + Prove the *specific* restarted worker dequeues and runs work. + + A single probe is not enough with N>1 workers — the surviving worker can + claim it. Submit a burst of tiny jobs (and keep topping up) until the + server log shows ``[Worker ] Processing job: ...``. + """ + if restart_match is None: + return False, "no Restarted worker log line to identify replacement" + + worker_index = int(restart_match.group(1)) + new_pid = int(restart_match.group(2)) + log( + f"Verifying restarted worker index={worker_index} pid={new_pid} " + f"actually dequeues jobs" + ) + + # Scan only log written after the restart line. + server.mark_log() + processing_pat = re.compile( + rf"\[Worker\s+{worker_index}\]\s+Processing job:\s+(\S+)" + ) + completed_pat = re.compile( + rf"\[Worker\s+{worker_index}\]\s+Completed job:\s+(\S+)" + ) + + spec = _tiny_probe_spec(pool) + # Enough jobs that every worker slot can be busy and the new one still + # gets at least one — then keep topping up until we see it or timeout. + batch = max(4, num_workers * 3) + probe_ids: list[str] = [] + deadline = time.monotonic() + timeout + saw_processing_job: Optional[str] = None + saw_completed_job: Optional[str] = None + + def _scan() -> None: + nonlocal saw_processing_job, saw_completed_job + text = server.read_log_delta() + probe_set = set(probe_ids) + for m in processing_pat.finditer(text): + jid = m.group(1) + if jid in probe_set and saw_processing_job is None: + saw_processing_job = jid + log( + f"[Worker {worker_index}] Processing probe job observed: " + f"{saw_processing_job}" + ) + break + for m2 in completed_pat.finditer(text): + jid = m2.group(1) + if jid in probe_set and saw_completed_job is None: + saw_completed_job = jid + log( + f"[Worker {worker_index}] Completed probe job observed: " + f"{saw_completed_job}" + ) + break + + try: + # Initial burst. + for _ in range(batch): + probe_ids.append(submit_spec(client, spec)) + time.sleep(0.02) + + while time.monotonic() < deadline: + _scan() + if saw_processing_job is not None: + # Prefer also seeing completion, but Processing is the + # dequeue/execute proof the user asked for. + # Give a short extra window for Completed on tiny LPs. + extra = time.monotonic() + 5.0 + while time.monotonic() < extra and saw_completed_job is None: + _scan() + time.sleep(0.05) + break + + # Top up so the queue stays non-empty if survivors are fast. + probe_ids.append(submit_spec(client, spec)) + time.sleep(0.05) + else: + _scan() + finally: + # Prefer waiting for tiny probes to finish so we do not SIGKILL the + # replacement worker we just proved is healthy. + deadline_cleanup = time.monotonic() + 10.0 + for jid in probe_ids: + while time.monotonic() < deadline_cleanup: + try: + st = client.status(jid) + except GrpcError: + break + if st not in (JobStatus.QUEUED, JobStatus.PROCESSING): + break + time.sleep(0.05) + try: + client.delete(jid) + except GrpcError: + pass + + if saw_processing_job is None: + # Also dump which workers *did* process probes, for diagnosis. + text = server.read_log_delta() + others = WORKER_PROCESSING_RE.findall(text) + return ( + False, + f"restarted worker {worker_index} (pid={new_pid}) never logged " + f"Processing for a probe job_id; other Processing lines={others!r}; " + f"probes_submitted={len(probe_ids)} probe_ids={probe_ids[:8]}", + ) + + if saw_processing_job not in probe_ids: + return ( + False, + f"restarted worker {worker_index} processed non-probe job " + f"{saw_processing_job}; probes={probe_ids[:8]}", + ) + + detail = ( + f"restarted_worker={worker_index} pid={new_pid} " + f"processed={saw_processing_job} completed={saw_completed_job} " + f"probes_submitted={len(probe_ids)}" + ) + return True, detail + + +def wait_for_job_cancelled_in_log( + server: ServerHandle, job_id: str, timeout: float = 10.0 +) -> tuple[bool, str]: + """Require server log proof that job_id was marked CANCELLED in the tracker.""" + pat = re.compile( + rf"Marked job CANCELLED in job_tracker:\s+{re.escape(job_id)}\b" + ) + m = server.wait_log_match(pat, timeout=timeout) + if m is None: + return False, f"no 'Marked job CANCELLED' log for {job_id}" + return True, f"tracker CANCELLED logged for {job_id}" + + +def cancel_log_names_job( + text: str, job_id: str +) -> tuple[Optional[str], Optional[re.Match]]: + """Return (mode, match) if a cancel-start / kill log names this job_id.""" + for mode, pat in ( + ("cooperative", CANCEL_COOP_RE), + ("legacy_kill", CANCEL_KILL_RE), + ("kill_fallback", CANCEL_FALLBACK_KILL_RE), + ): + for m in pat.finditer(text): + if m.group(1) == job_id: + return mode, m + return None, None + + +def check_cancelled_client_view( + client: Client, job_id: str +) -> tuple[bool, str]: + """After cancel: status CANCELLED; result raises; wait returns CANCELLED.""" + notes = [] + try: + st = client.status(job_id) + except GrpcError as e: + return False, f"status() raised: {e}" + notes.append(f"status={st.name}") + if st != JobStatus.CANCELLED: + return False, f"expected CANCELLED, got {st.name}" + + try: + wait_st = client.wait(job_id, timeout=10) + notes.append(f"wait={wait_st.name}") + if wait_st != JobStatus.CANCELLED: + return False, f"wait expected CANCELLED, got {wait_st.name}" + except GrpcError as e: + notes.append(f"wait raised (acceptable if racing): {e}") + + try: + client.result(job_id) + return ( + False, + "result() unexpectedly succeeded after cancel; " + + "; ".join(notes), + ) + except GrpcError as e: + notes.append(f"result raised as expected: {e}") + return True, "; ".join(notes) + + +def check_deleted_client_view(client: Client, job_id: str) -> tuple[bool, str]: + """After delete: status NOT_FOUND; result raises; cancel raises / not-found.""" + notes = [] + try: + st = client.status(job_id) + except GrpcError as e: + return False, f"status() raised: {e}" + notes.append(f"status={st.name}") + if st != JobStatus.NOT_FOUND: + return False, f"expected NOT_FOUND, got {st.name}" + + try: + client.result(job_id) + return ( + False, + "result() unexpectedly succeeded after delete; " + + "; ".join(notes), + ) + except GrpcError as e: + notes.append(f"result raised as expected: {e}") + + try: + client.cancel(job_id) + notes.append("cancel() returned without error (unexpected)") + return False, "; ".join(notes) + except GrpcError as e: + notes.append(f"cancel raised as expected: {e}") + return True, "; ".join(notes) + + +# --------------------------------------------------------------------------- +# Background traffic (random sizes / timing) +# --------------------------------------------------------------------------- + + +class TrafficGenerator: + """Daemon thread that keeps submitting mixed jobs with jitter.""" + + def __init__( + self, + port: int, + pool: list[ProblemSpec], + rng: random.Random, + enabled: threading.Event, + ): + self.port = port + self.pool = pool + self.rng = rng + self.enabled = enabled + self._stop = threading.Event() + self._idle = threading.Event() + self._idle.set() + self._lock = threading.Lock() + self._outstanding: set[str] = set() + self._thread = threading.Thread( + target=self._run, name="traffic", daemon=True + ) + self.submitted = 0 + self.errors = 0 + + def start(self) -> None: + self._thread.start() + + def stop(self) -> None: + self._stop.set() + self.enabled.set() # unblock waiters + self._thread.join(timeout=5) + + def wait_idle(self, timeout: float = 10.0) -> bool: + return self._idle.wait(timeout=timeout) + + def snapshot_outstanding(self) -> list[str]: + with self._lock: + return list(self._outstanding) + + def pause_and_drain(self, client: Client, settle_s: float = 1.0) -> None: + """Stop submitting, wait idle, prefer letting short jobs finish over SIGKILL.""" + self.enabled.clear() + if not self.wait_idle(timeout=15.0): + log("WARNING: traffic generator did not go idle within 15s") + leftover = self.snapshot_outstanding() + if leftover: + log(f"Draining {len(leftover)} traffic job(s) before trial") + for jid in leftover: + # Prefer waiting for completion so we do not SIGKILL mid-CUDA + # right before the next trial (especially result/shutdown). + try: + st = client.status(jid) + if st in (JobStatus.QUEUED, JobStatus.PROCESSING): + try: + client.wait(jid, timeout=8) + except GrpcError: + pass + except GrpcError: + pass + try: + client.delete(jid) + except GrpcError: + pass + with self._lock: + self._outstanding.discard(jid) + if leftover and settle_s > 0: + time.sleep(settle_s) + + def _run(self) -> None: + while not self._stop.is_set(): + if not self.enabled.is_set(): + self._idle.set() + time.sleep(0.05) + continue + self._idle.clear() + try: + client = make_client(self.port) + # Keep traffic short so pause_and_drain rarely needs SIGKILL. + short = [ + s + for s in self.pool + if s.time_limit <= 30 + or s.synth_factory in ("tiny_lp", "small_mip") + ] or self.pool + spec = self.rng.choice(short) + job_id = submit_spec(client, spec) + with self._lock: + self._outstanding.add(job_id) + self.submitted += 1 + # Sometimes leave them running; sometimes cancel/delete soon. + action = self.rng.choice( + ["leave", "leave", "cancel", "delete", "wait"] + ) + delay = self.rng.uniform(0.1, 1.5) + # If paused mid-delay, bail without further mutations. + deadline = time.monotonic() + delay + while time.monotonic() < deadline and self.enabled.is_set(): + time.sleep(0.05) + if not self.enabled.is_set(): + continue + if action == "cancel": + try: + client.cancel(job_id) + except GrpcError: + pass + elif action == "delete": + try: + client.delete(job_id) + except GrpcError: + pass + with self._lock: + self._outstanding.discard(job_id) + elif action == "wait": + try: + client.wait(job_id, timeout=5) + client.delete(job_id) + except GrpcError: + pass + with self._lock: + self._outstanding.discard(job_id) + except Exception as e: + self.errors += 1 + log(f"Traffic generator error: {e}") + time.sleep(0.5) + self._idle.set() + # Inter-job pause also respects the gate. + gap_deadline = time.monotonic() + self.rng.uniform(0.05, 0.8) + while time.monotonic() < gap_deadline and self.enabled.is_set(): + time.sleep(0.05) + + +def delete_jobs_quiet(client: Client, job_ids: list[str]) -> None: + for jid in job_ids: + try: + client.delete(jid) + except GrpcError: + pass + + +def settle_after_worker_kills( + server: ServerHandle, seconds: float = 1.5 +) -> None: + """Brief pause so replacement workers finish CUDA init after SIGKILL churn.""" + time.sleep(seconds) + # Touch the log so a subsequent mark_log starts after settle. + server.mark_log() + + +# --------------------------------------------------------------------------- +# Trial implementations +# --------------------------------------------------------------------------- + + +def trial_cancel( + server: ServerHandle, + pool: list[ProblemSpec], + rng: random.Random, + traffic_gate: threading.Event, +) -> TrialResult: + traffic_gate.clear() + client = make_client(server.port) + # Occupy both workers with long solves. + jobs = submit_long_jobs(client, pool, n=max(2, server.workers), rng=rng) + target = None + for jid in jobs: + st = wait_until_processing(client, jid, timeout=25) + if st == JobStatus.PROCESSING: + target = jid + break + if target is None: + for jid in jobs: + try: + client.delete(jid) + except GrpcError: + pass + return TrialResult( + action="cancel", + ok=False, + detail="no job reached PROCESSING before cancel", + ) + + server.mark_log() + t0 = time.monotonic() + try: + client.cancel(target) + except GrpcError as e: + return TrialResult( + action="cancel", + ok=False, + detail=f"cancel() raised: {e}", + elapsed_s=time.monotonic() - t0, + ) + cancel_elapsed = time.monotonic() - t0 + + client_ok, client_detail = check_cancelled_client_view(client, target) + + # Prefer cooperative cancel (no worker kill). Kill+restart is still OK + # as a fallback when the solver does not unwind within the grace period. + coop_mode = None + coop_match = None + kill_match = None + fallback_match = None + deadline = time.monotonic() + 7.0 + while time.monotonic() < deadline: + text = server.read_log_delta() + if coop_mode is None: + coop_mode, coop_match = cancel_log_names_job(text, target) + if kill_match is None: + kill_match = CANCEL_KILL_RE.search(text) + if kill_match is not None and kill_match.group(1) != target: + kill_match = None + if fallback_match is None: + fallback_match = CANCEL_FALLBACK_KILL_RE.search(text) + if ( + fallback_match is not None + and fallback_match.group(1) != target + ): + fallback_match = None + if ( + coop_mode in ("legacy_kill", "kill_fallback") + or kill_match + or fallback_match + ): + break + if ( + coop_mode == "cooperative" + and time.monotonic() > t0 + cancel_elapsed + 0.3 + ): + break + time.sleep(0.1) + + # Must see server log that this job was recorded CANCELLED (result path). + log_cancelled_ok, log_cancelled_detail = wait_for_job_cancelled_in_log( + server, target, timeout=10.0 + ) + # Cooperative cancel returns to the client before the worker finishes + # unwinding; on slow MIPs the Marked CANCELLED line can lag past the + # wait window (or arrive only after SIGKILL). Cancel-start log + client + # CANCELLED is sufficient in that race. + if not log_cancelled_ok and coop_mode is not None: + log_cancelled_detail = ( + f"no tracker CANCELLED line yet (ok for slow cooperative unwind); " + f"cancel-start log present mode={coop_mode}" + ) + log_cancelled_ok = True + + restart_match = None + worker_restarted = False + probe_ok = False + probe_detail = "pending" + mode = coop_mode or "unknown" + if fallback_match is not None: + mode = "kill_fallback" + elif kill_match is not None: + mode = "legacy_kill" + + if mode in ("kill_fallback", "legacy_kill"): + restart_match = server.wait_log_match( + RESTARTED_WORKER_RE, timeout=15.0 + ) + worker_restarted = restart_match is not None + if worker_restarted: + probe_ok, probe_detail = verify_restarted_worker_executes( + server, client, pool, server.workers, restart_match, timeout=45 + ) + else: + probe_detail = "worker kill logged but no Restarted worker line" + elif mode == "cooperative": + health_ok, health_detail = fetch_and_check_result( + client, + _synth_tiny_lp(), + make_settings(10.0), + ["x", "y"], + expected_obj=0.36, + label="post-cancel-health", + server=server, + ) + probe_ok = health_ok + probe_detail = f"cooperative health: {health_detail}" + seg_note = note_unexpected_worker_segfault(server.read_log_delta()) + if seg_note: + probe_detail = f"{probe_detail}; {seg_note}" + else: + probe_detail = f"no cancel log naming job {target}" + + delete_jobs_quiet(client, jobs) + settle_after_worker_kills(server) + + ok = ( + client_ok + and probe_ok + and log_cancelled_ok + and mode != "unknown" + and cancel_elapsed < 75 + ) + detail = ( + f"job={target}; mode={mode}; cancel_in={cancel_elapsed:.2f}s; " + f"cancel_log_for_job={mode != 'unknown'}; " + f"restart_logged={restart_match is not None}; " + f"log_cancelled=({log_cancelled_detail}); " + f"probe=({probe_detail}); {client_detail}" + ) + return TrialResult( + action="cancel", + ok=ok, + detail=detail, + elapsed_s=time.monotonic() - t0, + worker_restarted=worker_restarted, + client_status_ok=client_ok, + ) + + +def trial_delete( + server: ServerHandle, + pool: list[ProblemSpec], + rng: random.Random, + traffic_gate: threading.Event, +) -> TrialResult: + traffic_gate.clear() + client = make_client(server.port) + jobs = submit_long_jobs(client, pool, n=max(2, server.workers), rng=rng) + target = None + for jid in jobs: + st = wait_until_processing(client, jid, timeout=25) + if st == JobStatus.PROCESSING: + target = jid + break + if target is None: + for jid in jobs: + try: + client.delete(jid) + except GrpcError: + pass + return TrialResult( + action="delete", + ok=False, + detail="no job reached PROCESSING before delete", + ) + + server.mark_log() + t0 = time.monotonic() + try: + client.delete(target) + except GrpcError as e: + return TrialResult( + action="delete", + ok=False, + detail=f"delete() raised: {e}", + elapsed_s=time.monotonic() - t0, + ) + delete_elapsed = time.monotonic() - t0 + + client_ok, client_detail = check_deleted_client_view(client, target) + + coop_mode = None + deadline = time.monotonic() + 7.0 + while time.monotonic() < deadline: + text = server.read_log_delta() + if coop_mode is None: + coop_mode, _ = cancel_log_names_job(text, target) + if coop_mode in ("legacy_kill", "kill_fallback"): + break + if ( + coop_mode == "cooperative" + and time.monotonic() > t0 + delete_elapsed + 0.3 + ): + break + time.sleep(0.1) + + # delete() cancels then erases the tracker entry, so the result-thread + # "Marked job CANCELLED" line can race and be skipped. Prefer it when + # present; always require a cancel-start log naming this job_id. + log_cancelled_ok, log_cancelled_detail = wait_for_job_cancelled_in_log( + server, target, timeout=3.0 + ) + cancel_log_ok = coop_mode is not None + if not log_cancelled_ok and cancel_log_ok: + log_cancelled_detail = ( + f"no tracker CANCELLED line (ok for delete race); " + f"cancel-start log present mode={coop_mode}" + ) + log_cancelled_ok = True + + restart_match = None + worker_restarted = False + probe_ok = False + probe_detail = "pending" + mode = coop_mode or "unknown" + + if mode in ("kill_fallback", "legacy_kill"): + restart_match = server.wait_log_match( + RESTARTED_WORKER_RE, timeout=15.0 + ) + worker_restarted = restart_match is not None + if worker_restarted: + probe_ok, probe_detail = verify_restarted_worker_executes( + server, client, pool, server.workers, restart_match, timeout=45 + ) + else: + probe_detail = "worker kill logged but no Restarted worker line" + elif mode == "cooperative": + health_ok, health_detail = fetch_and_check_result( + client, + _synth_tiny_lp(), + make_settings(10.0), + ["x", "y"], + expected_obj=0.36, + label="post-delete-health", + server=server, + ) + probe_ok = health_ok + probe_detail = f"cooperative health: {health_detail}" + seg_note = note_unexpected_worker_segfault(server.read_log_delta()) + if seg_note: + probe_detail = f"{probe_detail}; {seg_note}" + else: + probe_detail = f"no cancel log naming job {target}" + + delete_jobs_quiet(client, [j for j in jobs if j != target]) + settle_after_worker_kills(server) + + ok = ( + client_ok + and probe_ok + and log_cancelled_ok + and cancel_log_ok + and delete_elapsed < 75 + ) + detail = ( + f"job={target}; mode={mode}; delete_in={delete_elapsed:.2f}s; " + f"cancel_log_for_job={cancel_log_ok}; " + f"restart_logged={restart_match is not None}; " + f"log_cancelled=({log_cancelled_detail}); " + f"probe=({probe_detail}); {client_detail}" + ) + return TrialResult( + action="delete", + ok=ok, + detail=detail, + elapsed_s=time.monotonic() - t0, + worker_restarted=worker_restarted, + client_status_ok=client_ok, + ) + + +def wait_for_solver_intermediates( + client: Client, + job_id: str, + min_intermediate_lines: int = 2, + timeout: float = 60.0, +) -> tuple[bool, list[str]]: + """ + Stream solver logs until we have seen enough intermediate progress lines + (e.g. MIP ``New solution ... Objective ...``), proving the solve is well + underway — not merely queued or in the banner phase. + """ + collected: list[str] = [] + intermediates: list[str] = [] + lock = threading.Lock() + + def _on_line(line, job_complete=False): + text = line if isinstance(line, str) else str(line) + with lock: + collected.append(text) + if INTERMEDIATE_LOG_RE.search(text): + intermediates.append(text) + if len(intermediates) <= 5: + log(f"[solver-log] {text}") + # Keep streaming until the harness stops / server dies. + return True + + try: + client.start_log_stream(job_id, callback=_on_line) + except GrpcError as e: + log(f"WARNING: could not start log stream for {job_id}: {e}") + return False, [] + + deadline = time.monotonic() + timeout + try: + while time.monotonic() < deadline: + with lock: + n_inter = len(intermediates) + if n_inter >= min_intermediate_lines: + log( + f"Saw {n_inter} intermediate solver log lines on {job_id}; " + "ready for SIGINT" + ) + return True, list(intermediates) + + try: + st = client.status(job_id) + except GrpcError: + break + if st not in (JobStatus.QUEUED, JobStatus.PROCESSING): + log( + f"Job {job_id} left PROCESSING before intermediates " + f"(status={st.name}, intermediates={n_inter})" + ) + break + time.sleep(0.1) + finally: + # Best-effort join; stream may die with the upcoming SIGINT anyway. + try: + client.join_log_stream(job_id, timeout=1.0) + except Exception: + pass + + with lock: + return len(intermediates) >= min_intermediate_lines, list( + intermediates + ) + + +def trial_shutdown( + server: ServerHandle, + pool: list[ProblemSpec], + rng: random.Random, + traffic_gate: threading.Event, + shutdown_timeout: float, + intermediate_timeout: float = 60.0, + min_intermediate_lines: int = 2, +) -> TrialResult: + """SIGINT while work is in flight; pass if the server process exits.""" + del ( + intermediate_timeout, + min_intermediate_lines, + ) # kept for call-site compat + traffic_gate.clear() + if not server.is_running(): + server.start() + if not ensure_workers_healthy(server, pool): + return TrialResult( + action="shutdown", + ok=False, + detail="workers unhealthy before shutdown trial; " + "server restart did not recover", + elapsed_s=0.0, + server_exited=False, + mid_job_at_signal=False, + ) + client = make_client(server.port) + jobs = submit_long_jobs( + client, pool, n=max(2, server.workers), rng=rng, min_time_limit=60.0 + ) + processing_ids = [] + for jid in jobs: + st = wait_until_processing(client, jid, timeout=45) + if st == JobStatus.PROCESSING: + processing_ids.append(jid) + + mid_job_at_signal = len(processing_ids) > 0 + if not mid_job_at_signal: + log("WARNING: no PROCESSING jobs at SIGINT time; still testing exit") + + server.mark_log() + t0 = time.monotonic() + server.send_sigint() + exited = server.wait_exited(timeout=shutdown_timeout) + elapsed = time.monotonic() - t0 + shutdown_logged = ( + server.wait_log_match(SHUTDOWN_SIGNAL_RE, timeout=0.5) is not None + ) + + if not exited: + server.force_kill() + server.wait_exited(timeout=5) + + ok = exited and elapsed < shutdown_timeout + detail = ( + f"processing_at_signal={processing_ids}; " + f"exited={exited}; elapsed={elapsed:.2f}s; " + f"shutdown_logged={shutdown_logged}; " + f"timeout={shutdown_timeout}s" + ) + + try: + server.start() + except Exception as e: + ok = False + detail += f"; restart_after_shutdown_failed: {e}" + return TrialResult( + action="shutdown", + ok=ok, + detail=detail, + elapsed_s=elapsed, + server_exited=exited, + mid_job_at_signal=mid_job_at_signal, + ) + + +# --------------------------------------------------------------------------- +# Main loop +# --------------------------------------------------------------------------- + + +def parse_args(argv: Optional[list[str]] = None) -> argparse.Namespace: + p = argparse.ArgumentParser( + description=__doc__, + formatter_class=argparse.RawDescriptionHelpFormatter, + ) + p.add_argument( + "--workers", type=int, default=2, help="gRPC worker processes" + ) + p.add_argument("--port", type=int, default=15051, help="listen port") + p.add_argument( + "--loops", type=int, default=12, help="number of trial iterations" + ) + p.add_argument( + "--actions", + default="cancel,delete,result,shutdown", + help="comma-separated actions to cycle (cancel,delete,result,shutdown)", + ) + p.add_argument( + "--shutdown-timeout", + type=float, + default=15.0, + help="max seconds to wait for server exit after SIGINT", + ) + p.add_argument( + "--intermediate-timeout", + type=float, + default=60.0, + help="max seconds to wait for solver intermediate log lines before SIGINT", + ) + p.add_argument( + "--min-intermediate-lines", + type=int, + default=1, + help="require this many intermediate solver log lines before SIGINT", + ) + p.add_argument( + "--log-dir", + type=Path, + default=Path("/tmp/cuopt_grpc_lifecycle_stress"), + help="directory for server logs and harness output", + ) + p.add_argument( + "--datasets", + type=Path, + default=Path( + os.environ.get("RAPIDS_DATASET_ROOT_DIR", str(_DEFAULT_DATASETS)) + ), + help="dataset root containing mip/ and linear_programming/", + ) + p.add_argument("--seed", type=int, default=None, help="RNG seed") + p.add_argument( + "--no-traffic", + action="store_true", + help="disable background random job traffic", + ) + p.add_argument( + "--fixed-order", + action="store_true", + help="cycle actions in listed order instead of random choice", + ) + return p.parse_args(argv) + + +def main(argv: Optional[list[str]] = None) -> int: + args = parse_args(argv) + rng = random.Random(args.seed if args.seed is not None else time.time_ns()) + actions = [a.strip() for a in args.actions.split(",") if a.strip()] + for a in actions: + if a not in ("cancel", "delete", "result", "shutdown"): + log(f"Unknown action: {a}") + return 2 + + args.log_dir.mkdir(parents=True, exist_ok=True) + server_log = args.log_dir / "server.log" + summary = Summary() + + pool = build_problem_pool(args.datasets) + log(f"Problem pool ({len(pool)}): {[s.name for s in pool]}") + log(f"Datasets root: {args.datasets}") + log(f"Log dir: {args.log_dir}") + + server = ServerHandle( + workers=args.workers, + port=args.port, + server_log=server_log, + ) + traffic_gate = threading.Event() + # Stay closed unless we explicitly open it for a brief inter-trial burst. + traffic = None + + try: + server.start() + if not args.no_traffic: + traffic = TrafficGenerator( + args.port, pool, random.Random(rng.random()), traffic_gate + ) + traffic.start() + for i in range(args.loops): + if args.fixed_order: + action = actions[i % len(actions)] + else: + action = rng.choice(actions) + log("-" * 72) + log(f"LOOP {i + 1}/{args.loops} action={action}") + if not server.is_running(): + log("Server not running; restarting before trial") + server.start() + + # Pause traffic and drain leftovers so trials are not racing + # cancel/delete SIGKILLs against background jobs. + traffic_gate.clear() + client = make_client(args.port) + if traffic is not None: + traffic.pause_and_drain(client, settle_s=1.0) + + if action == "cancel": + trial = trial_cancel(server, pool, rng, traffic_gate) + elif action == "delete": + trial = trial_delete(server, pool, rng, traffic_gate) + elif action == "result": + trial = trial_result(server, args.datasets, traffic_gate) + else: + trial = trial_shutdown( + server, + pool, + rng, + traffic_gate, + args.shutdown_timeout, + intermediate_timeout=args.intermediate_timeout, + min_intermediate_lines=args.min_intermediate_lines, + ) + summary.add(trial) + + # Optional brief traffic burst *between* trials — skip before + # result/shutdown so we do not SIGKILL mid-CUDA into a poisoned + # GPU right before those checks. + next_action = None + if i + 1 < args.loops: + next_action = ( + actions[(i + 1) % len(actions)] + if args.fixed_order + else None + ) + allow_burst = traffic is not None and i + 1 < args.loops + if allow_burst and next_action in ("result", "shutdown"): + allow_burst = False + log(f"Skipping traffic burst before next action={next_action}") + if allow_burst: + traffic_gate.set() + time.sleep(rng.uniform(0.3, 1.2)) + traffic_gate.clear() + traffic.pause_and_drain(make_client(args.port), settle_s=1.0) + else: + time.sleep(rng.uniform(0.2, 0.5)) + except KeyboardInterrupt: + log("Interrupted by user") + finally: + if traffic is not None: + traffic.stop() + log( + f"Traffic submitted={traffic.submitted} errors={traffic.errors}" + ) + server.force_kill() + + return summary.print_report() + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/python/cuopt/cuopt/tests/linear_programming/test_grpc_client.py b/python/cuopt/cuopt/tests/linear_programming/test_grpc_client.py index 7807c757ff..4321194447 100644 --- a/python/cuopt/cuopt/tests/linear_programming/test_grpc_client.py +++ b/python/cuopt/cuopt/tests/linear_programming/test_grpc_client.py @@ -2,7 +2,10 @@ # SPDX-License-Identifier: Apache-2.0 import os +import re +import tempfile import time +from pathlib import Path import pytest @@ -18,7 +21,12 @@ from cuopt.linear_programming.problem import INTEGER, MAXIMIZE, Problem from cuopt.linear_programming.solver.solver_parameters import CUOPT_TIME_LIMIT -from grpc_server_fixtures import GRPC_PORT_OFFSET_CLIENT +from grpc_server_fixtures import ( + GRPC_PORT_OFFSET_CLIENT, + GRPC_PORT_OFFSET_COOP_CANCEL, + start_grpc_server, + stop_grpc_server, +) RAPIDS_DATASET_ROOT_DIR = os.getenv("RAPIDS_DATASET_ROOT_DIR") if RAPIDS_DATASET_ROOT_DIR is None: @@ -26,10 +34,26 @@ RAPIDS_DATASET_ROOT_DIR = os.path.join(RAPIDS_DATASET_ROOT_DIR, "datasets") _SWATH1_MPS = os.path.join(RAPIDS_DATASET_ROOT_DIR, "mip", "swath1.mps") +_SEYMOUR1_MPS = os.path.join(RAPIDS_DATASET_ROOT_DIR, "mip", "seymour1.mps") _DEMO_LP_NAMES = ["x", "y"] _MIP_NAMES = ["x", "y"] +_CANCEL_COOP_RE = re.compile( + r"Cancelling running job\s+(\S+)\s+\(cooperative cancel; worker\s+(\d+)\)" +) +_CANCEL_KILL_RE = re.compile( + r"Cancelling running job\s+(\S+)\s+\(killing worker\s+(\d+)\)" +) +_CANCEL_FALLBACK_KILL_RE = re.compile( + r"Job\s+(\S+)\s+still running after cooperative cancel grace;.*" + r"SIGKILL to worker\s+(\d+)", + re.DOTALL, +) +_RESTARTED_WORKER_RE = re.compile( + r"Restarted worker\s+(\d+)\s+with PID\s+(\d+)" +) + def _demo_lp_problem(): problem = Problem("grpc_demo") @@ -55,6 +79,18 @@ def _poll_until_complete( return client.status(job_id) +def _wait_until_processing(client, job_id, timeout=30.0): + deadline = time.monotonic() + timeout + while time.monotonic() < deadline: + status = client.status(job_id) + if status == JobStatus.PROCESSING: + return status + if status not in (JobStatus.QUEUED, JobStatus.PROCESSING): + return status + time.sleep(0.05) + return client.status(job_id) + + def _infeasible_lp_problem(): problem = Problem("grpc_infeasible") x = problem.addVariable(lb=0.0, name="x") @@ -75,6 +111,25 @@ def _assert_demo_lp_solution(client): client.delete(job_id) +def _read_log(path: Path) -> str: + try: + return path.read_text(encoding="utf-8", errors="replace") + except OSError: + return "" + + +def _wait_log_match(path: Path, pattern: re.Pattern, timeout: float = 30.0): + deadline = time.monotonic() + timeout + while time.monotonic() < deadline: + text = _read_log(path) + match = pattern.search(text) + if match is not None: + return match, text + time.sleep(0.1) + text = _read_log(path) + return pattern.search(text), text + + class TestTlsConfig: def test_mtls_requires_both_client_materials(self): pem = "-----BEGIN CERTIFICATE-----\nabc\n-----END CERTIFICATE-----" @@ -224,18 +279,20 @@ def test_cancel_job(self, grpc_server): problem = Read(_SWATH1_MPS) settings = SolverSettings() - settings.set_parameter(CUOPT_TIME_LIMIT, 10) + settings.set_parameter(CUOPT_TIME_LIMIT, 120) client = Client("localhost", grpc_server) job_id = client.submit(problem, settings) - status = client.status(job_id) - if status not in (JobStatus.QUEUED, JobStatus.PROCESSING): + status = _wait_until_processing(client, job_id, timeout=30.0) + if status != JobStatus.PROCESSING: client.delete(job_id) - pytest.skip("Job completed before cancellation could be observed") + pytest.skip( + f"Job never reached PROCESSING before cancel (status={status})" + ) client.cancel(job_id) - assert client.wait(job_id, timeout=30) == JobStatus.CANCELLED + assert client.wait(job_id, timeout=90) == JobStatus.CANCELLED with pytest.raises(GrpcError): client.result(job_id) client.delete(job_id) @@ -332,3 +389,78 @@ def test_mtls_server_rejects_missing_client_cert(self, mtls_server_info): mtls_server_info["port"], tls=TlsConfig(os.path.join(cert_dir, "ca.crt")), ) + + +@pytest.mark.xdist_group(name="grpc_coop_cancel") +@pytest.mark.filterwarnings("ignore::DeprecationWarning") +class TestGrpcCooperativeCancel: + """Cancel must unwind cooperatively without killing/restarting the worker.""" + + def test_cancel_is_cooperative_without_worker_restart(self): + mps = _SEYMOUR1_MPS if os.path.isfile(_SEYMOUR1_MPS) else _SWATH1_MPS + if not os.path.isfile(mps): + pytest.skip(f"dataset not found: {mps}") + + with tempfile.TemporaryDirectory() as tmp: + server_log = Path(tmp) / "cuopt_grpc_server.log" + server_log.write_text("") + proc = None + try: + proc, client_env = start_grpc_server( + GRPC_PORT_OFFSET_COOP_CANCEL, + server_log_path=server_log, + workers=1, + ) + port = int(client_env["CUOPT_REMOTE_PORT"]) + client = Client("localhost", port) + + settings = SolverSettings() + settings.set_parameter(CUOPT_TIME_LIMIT, 300) + job_id = client.submit(Read(mps), settings) + + status = _wait_until_processing(client, job_id, timeout=45.0) + if status != JobStatus.PROCESSING: + client.delete(job_id) + pytest.skip( + f"Job never reached PROCESSING (status={status})" + ) + + # Snapshot log size before cancel so post-cancel restart lines + # are attributable to this trial. + pre_cancel_log = _read_log(server_log) + client.cancel(job_id) + assert client.wait(job_id, timeout=120) == JobStatus.CANCELLED + with pytest.raises(GrpcError): + client.result(job_id) + + coop_match, log_text = _wait_log_match( + server_log, _CANCEL_COOP_RE, timeout=15.0 + ) + assert coop_match is not None, ( + "expected cooperative cancel log line; " + f"log tail:\n{log_text[-2000:]}" + ) + assert coop_match.group(1) == job_id + + post_cancel = log_text[len(pre_cancel_log) :] + assert _CANCEL_KILL_RE.search(post_cancel) is None, ( + "legacy immediate-kill cancel path used" + ) + assert _CANCEL_FALLBACK_KILL_RE.search(post_cancel) is None, ( + "cooperative cancel fell back to SIGKILL" + ) + assert _RESTARTED_WORKER_RE.search(post_cancel) is None, ( + "worker restarted after cancel" + ) + + # Same worker must still accept work without a restart. + _assert_demo_lp_solution(client) + health_log = _read_log(server_log)[len(pre_cancel_log) :] + assert _RESTARTED_WORKER_RE.search(health_log) is None, ( + "worker restarted during post-cancel health solve" + ) + + client.delete(job_id) + finally: + if proc is not None: + stop_grpc_server(proc)