From 09eb2c976a23e6cc6b1d74f5cba2ca04bef8ad7e Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Wed, 22 Apr 2026 16:56:57 +0200 Subject: [PATCH 01/30] reduced memory footprint of problem_t inside pdlp. reduced memory footprint of duplicated vectors in op_problem_scaled passes all tests --- cpp/src/mip_heuristics/problem/problem.cu | 16 +++++++++++ cpp/src/mip_heuristics/problem/problem.cuh | 1 + cpp/src/pdlp/cusparse_view.cu | 32 ++++++++++++++++++++++ cpp/src/pdlp/cusparse_view.hpp | 3 ++ cpp/src/pdlp/pdlp.cu | 21 +++++++++++++- cpp/src/pdlp/solve.cu | 5 +++- 6 files changed, 76 insertions(+), 2 deletions(-) diff --git a/cpp/src/mip_heuristics/problem/problem.cu b/cpp/src/mip_heuristics/problem/problem.cu index 10d80586b4..f017fc3a46 100644 --- a/cpp/src/mip_heuristics/problem/problem.cu +++ b/cpp/src/mip_heuristics/problem/problem.cu @@ -2391,6 +2391,22 @@ void problem_t::update_variable_bounds(const std::vector& var_ind RAFT_CHECK_CUDA(handle_ptr->get_stream()); } +template +void poison_and_free(rmm::device_uvector& uvec, rmm::cuda_stream_view stream) { + if (uvec.size() > 0) { + thrust::fill(thrust::cuda::par_nosync.on(stream), uvec.begin(), uvec.end(), 5); + } + uvec.resize(0, stream); +} + +template +void problem_t::pdlp_lighten(){ + variable_types.resize(0, handle_ptr->get_stream()); + poison_and_free(related_variables_offsets, handle_ptr->get_stream()); + poison_and_free(integer_fixed_variable_map, handle_ptr->get_stream()); + poison_and_free(combined_bounds, handle_ptr->get_stream()); +} + #if MIP_INSTANTIATE_FLOAT || PDLP_INSTANTIATE_FLOAT template class problem_t; #endif diff --git a/cpp/src/mip_heuristics/problem/problem.cuh b/cpp/src/mip_heuristics/problem/problem.cuh index a801cc4067..0a25bc63b0 100644 --- a/cpp/src/mip_heuristics/problem/problem.cuh +++ b/cpp/src/mip_heuristics/problem/problem.cuh @@ -151,6 +151,7 @@ class problem_t { const std::vector& offset_values, const std::vector& coefficient_values); void sort_rows_by_variables(const raft::handle_t* handle_ptr); + void pdlp_lighten(); enum var_flags_t : i_t { VAR_IMPLIED_INTEGER = 1 << 0, }; diff --git a/cpp/src/pdlp/cusparse_view.cu b/cpp/src/pdlp/cusparse_view.cu index 64ec44f5ef..22a3ef5b2d 100644 --- a/cpp/src/pdlp/cusparse_view.cu +++ b/cpp/src/pdlp/cusparse_view.cu @@ -1082,6 +1082,38 @@ void cusparse_view_t::update_mixed_precision_matrices() } } +template +void cusparse_view_t::redirect_rows_and_cols( + const problem_t& original_problem) +{ + RAFT_CUSPARSE_TRY(cusparseCsrSetPointers(A, + const_cast(original_problem.offsets.data()), + const_cast(original_problem.variables.data()), + const_cast(A_.data()))); + + RAFT_CUSPARSE_TRY( + cusparseCsrSetPointers(A_T, + const_cast(original_problem.reverse_offsets.data()), + const_cast(original_problem.reverse_constraints.data()), + const_cast(A_T_.data()))); + + if constexpr (std::is_same_v) { + if (mixed_precision_enabled_) { + RAFT_CUSPARSE_TRY( + cusparseCsrSetPointers(A_mixed_, + const_cast(original_problem.offsets.data()), + const_cast(original_problem.variables.data()), + A_float_.data())); + + RAFT_CUSPARSE_TRY( + cusparseCsrSetPointers(A_T_mixed_, + const_cast(original_problem.reverse_offsets.data()), + const_cast(original_problem.reverse_constraints.data()), + A_T_float_.data())); + } + } +} + // Mixed precision SpMV implementation: FP32 matrix with FP64 vectors and FP64 compute type size_t mixed_precision_spmv_buffersize(cusparseHandle_t handle, cusparseOperation_t opA, diff --git a/cpp/src/pdlp/cusparse_view.hpp b/cpp/src/pdlp/cusparse_view.hpp index 416a0b1e5f..81fcbd9ac3 100644 --- a/cpp/src/pdlp/cusparse_view.hpp +++ b/cpp/src/pdlp/cusparse_view.hpp @@ -208,6 +208,9 @@ class cusparse_view_t { // Update FP32 matrix copies after scaling (must be called after scale_problem()) void update_mixed_precision_matrices(); + + // Redirects the vectors of rows and columns from op_problem_scaled_ to the original problem so that they can be freed + void redirect_rows_and_cols(const problem_t& original_problem); }; // Mixed precision SpMV: FP32 matrix with FP64 vectors and FP64 compute type diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index a759887fc5..1377232fa7 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -292,7 +292,6 @@ pdlp_solver_t::pdlp_solver_t(problem_t& op_problem, ? -std::numeric_limits::infinity() : std::numeric_limits::infinity(); op_problem.check_problem_representation(true, false); - op_problem_scaled_.check_problem_representation(true, false); if (settings_.new_bounds.size() > 0) { batch_solution_to_return_.get_additional_termination_informations().resize( @@ -2158,6 +2157,14 @@ void pdlp_solver_t::transpose_primal_dual_to_row( stream_view_); } +template +void poison_and_free(rmm::device_uvector& uvec, rmm::cuda_stream_view stream) { + if (uvec.size() > 0) { + thrust::fill(thrust::cuda::par_nosync.on(stream), uvec.begin(), uvec.end(), 5); + } + uvec.resize(0, stream); +} + template void pdlp_solver_t::transpose_primal_dual_back_to_col( rmm::device_uvector& primal_to_transpose, @@ -2263,6 +2270,18 @@ optimization_problem_solution_t pdlp_solver_t::run_solver(co // Update FP32 matrix copies for mixed precision SpMV after scaling pdhg_solver_.get_cusparse_view().update_mixed_precision_matrices(); + // Redirect cuSPARSE descriptors to use the original problem's structural data (offsets, indices), + // then free the duplicated structural vectors from the scaled copy to save device memory. + pdhg_solver_.get_cusparse_view().redirect_rows_and_cols(*problem_ptr); + poison_and_free(op_problem_scaled_.variables, stream_view_); + poison_and_free(op_problem_scaled_.offsets, stream_view_); + poison_and_free(op_problem_scaled_.reverse_constraints, stream_view_); + poison_and_free(op_problem_scaled_.reverse_offsets, stream_view_); + + // Remove unused device vectors + op_problem_scaled_.pdlp_lighten(); + problem_ptr->pdlp_lighten(); + if (!settings_.hyper_params.compute_initial_step_size_before_scaling && !settings_.get_initial_step_size().has_value()) compute_initial_step_size(); diff --git a/cpp/src/pdlp/solve.cu b/cpp/src/pdlp/solve.cu index f0345368a7..64acf4a04c 100644 --- a/cpp/src/pdlp/solve.cu +++ b/cpp/src/pdlp/solve.cu @@ -315,8 +315,10 @@ void set_pdlp_solver_mode(pdlp_solver_settings_t& settings) set_Methodical1(settings.hyper_params); else if (settings.pdlp_solver_mode == pdlp_solver_mode_t::Fast1) set_Fast1(settings.hyper_params); - else if (settings.pdlp_solver_mode == pdlp_solver_mode_t::Stable3) + else if (settings.pdlp_solver_mode == pdlp_solver_mode_t::Stable3) { set_Stable3(settings.hyper_params); + settings.detect_infeasibility = false; + } } std::atomic global_concurrent_halt{0}; @@ -1678,6 +1680,7 @@ cuopt::linear_programming::optimization_problem_t mps_data_model_to_op return op_problem; } +//HERE template optimization_problem_solution_t solve_lp( raft::handle_t const* handle_ptr, From e9f26a1cdc7c6fe0217ee74413b618bb4aeed7ed Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Wed, 22 Apr 2026 17:10:16 +0200 Subject: [PATCH 02/30] added is_cupdlpx bool and removed unscaled_primal_dual_avg_solution from cupdlpx. passes all tests --- cpp/src/pdlp/pdlp.cu | 5 +++-- cpp/src/pdlp/pdlp.cuh | 1 + 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index 1377232fa7..1a631ffb34 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -118,10 +118,11 @@ pdlp_solver_t::pdlp_solver_t(problem_t& op_problem, stream_view_(handle_ptr_->get_stream()), settings_(settings), problem_ptr(&op_problem), + is_cupdlpx_(is_cupdlpx_restart(settings.hyper_params)), op_problem_scaled_( op_problem, false), // False to call the PDLP custom version of the problem copy constructor - unscaled_primal_avg_solution_{static_cast(op_problem.n_variables), stream_view_}, - unscaled_dual_avg_solution_{static_cast(op_problem.n_constraints), stream_view_}, + unscaled_primal_avg_solution_{is_cupdlpx_ ? 0 : static_cast(op_problem.n_variables), stream_view_}, + unscaled_dual_avg_solution_{is_cupdlpx_ ? 0 : static_cast(op_problem.n_constraints), stream_view_}, primal_size_h_(op_problem.n_variables), dual_size_h_(op_problem.n_constraints), primal_step_size_{climber_strategies_.size(), stream_view_}, diff --git a/cpp/src/pdlp/pdlp.cuh b/cpp/src/pdlp/pdlp.cuh index d03430f150..9251b268e1 100644 --- a/cpp/src/pdlp/pdlp.cuh +++ b/cpp/src/pdlp/pdlp.cuh @@ -143,6 +143,7 @@ class pdlp_solver_t { settings_.shared_sb_solved}; problem_t* problem_ptr; + bool is_cupdlpx_; // Combined bounds in op_problem_scaled_ will only be scaled if // compute_initial_primal_weight_before_scaling is false because of compute_initial_primal_weight problem_t op_problem_scaled_; From e134ce883d66fc49f1b8ad8691c1bf8a97301b18 Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Thu, 23 Apr 2026 13:55:44 +0200 Subject: [PATCH 03/30] free iteration_constraint_matrix_scaling_ and variables. passes all tests --- .../pdlp/initial_scaling_strategy/initial_scaling.cu | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/cpp/src/pdlp/initial_scaling_strategy/initial_scaling.cu b/cpp/src/pdlp/initial_scaling_strategy/initial_scaling.cu index a76b1773f9..1a5108925d 100644 --- a/cpp/src/pdlp/initial_scaling_strategy/initial_scaling.cu +++ b/cpp/src/pdlp/initial_scaling_strategy/initial_scaling.cu @@ -28,6 +28,13 @@ namespace cuopt::linear_programming::detail { +template +void poison_and_free(rmm::device_uvector& uvec, rmm::cuda_stream_view stream) { + if (uvec.size() > 0) { + thrust::fill(thrust::cuda::par_nosync.on(stream), uvec.begin(), uvec.end(), 5); + } + uvec.resize(0, stream); +} template pdlp_initial_scaling_strategy_t::pdlp_initial_scaling_strategy_t( raft::handle_t const* handle_ptr, @@ -81,6 +88,9 @@ pdlp_initial_scaling_strategy_t::pdlp_initial_scaling_strategy_t( f_t(1)); compute_scaling_vectors(number_of_ruiz_iterations, alpha); + + poison_and_free(iteration_constraint_matrix_scaling_, stream_view_); + poison_and_free(iteration_variable_scaling_, stream_view_); } template From 9c83d56948697863f0af717c483df0873d818414 Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Thu, 23 Apr 2026 14:13:06 +0200 Subject: [PATCH 04/30] removed useless data in pdlp_restart_strategy_t::weighted_average_solution_ for cupdlpx --- cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu index 821238fe84..b7c767d382 100644 --- a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu +++ b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu @@ -76,7 +76,10 @@ pdlp_restart_strategy_t::pdlp_restart_strategy_t( : handle_ptr_(handle_ptr), stream_view_(handle_ptr_->get_stream()), batch_mode_(climber_strategies.size() > 1), - weighted_average_solution_{handle_ptr_, primal_size, dual_size, is_legacy_batch_mode}, + weighted_average_solution_{handle_ptr_, + is_cupdlpx_restart(hyper_params) ? 0 : primal_size, + is_cupdlpx_restart(hyper_params) ? 0 : dual_size, + is_legacy_batch_mode}, primal_size_h_(primal_size), dual_size_h_(dual_size), problem_ptr(&op_problem), From f9e5c5f1571bd07fb73deb1a15f8ec6c4be84a6e Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Thu, 23 Apr 2026 14:27:43 +0200 Subject: [PATCH 05/30] only allocate infeasibility_information_t.primal/dual_slack if set to true (default to false) passes all tests --- .../pdlp/termination_strategy/infeasibility_information.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/src/pdlp/termination_strategy/infeasibility_information.cu b/cpp/src/pdlp/termination_strategy/infeasibility_information.cu index f795d2c4ca..9268e17910 100644 --- a/cpp/src/pdlp/termination_strategy/infeasibility_information.cu +++ b/cpp/src/pdlp/termination_strategy/infeasibility_information.cu @@ -81,11 +81,11 @@ infeasibility_information_t::infeasibility_information_t( (!infeasibility_detection) ? 0 : static_cast(dual_size_h_), stream_view_}, homogenous_dual_upper_bounds_{ (!infeasibility_detection) ? 0 : static_cast(dual_size_h_), stream_view_}, - primal_slack_{(is_cupdlpx_restart(hyper_params)) + primal_slack_{(is_cupdlpx_restart(hyper_params) && infeasibility_detection) ? static_cast(dual_size_h_ * climber_strategies.size()) : 0, stream_view_}, - dual_slack_{(is_cupdlpx_restart(hyper_params)) + dual_slack_{(is_cupdlpx_restart(hyper_params) && infeasibility_detection) ? static_cast(primal_size_h_ * climber_strategies.size()) : 0, stream_view_}, From 2ea5e4c7851c3f7ad734a0c378819263255fe729 Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Thu, 23 Apr 2026 15:07:53 +0200 Subject: [PATCH 06/30] dont allocate device vectors in average_termination_strategy passes all tests --- cpp/src/pdlp/pdlp.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index 1a631ffb34..448b419202 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -193,8 +193,8 @@ pdlp_solver_t::pdlp_solver_t(problem_t& op_problem, op_problem_scaled_, average_op_problem_evaluation_cusparse_view_, pdhg_solver_.get_cusparse_view(), - primal_size_h_, - dual_size_h_, + is_cupdlpx_ ? 0 : primal_size_h_, + is_cupdlpx_ ? 0 : dual_size_h_, initial_scaling_strategy_, settings_, climber_strategies_}, From a77eed83ec38363b1e48ad316f92b6bd96a6f112 Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Thu, 23 Apr 2026 15:27:48 +0200 Subject: [PATCH 07/30] removed primal_grqadient from saddlepoint in cupdlpx --- cpp/src/pdlp/pdhg.cu | 2 +- cpp/src/pdlp/saddle_point.cu | 2 +- cpp/src/pdlp/saddle_point.hpp | 3 ++- 3 files changed, 4 insertions(+), 3 deletions(-) diff --git a/cpp/src/pdlp/pdhg.cu b/cpp/src/pdlp/pdhg.cu index cb16c9d662..910875ea03 100644 --- a/cpp/src/pdlp/pdhg.cu +++ b/cpp/src/pdlp/pdhg.cu @@ -52,7 +52,7 @@ pdhg_solver_t::pdhg_solver_t( primal_size_h_(problem_ptr->n_variables), dual_size_h_(problem_ptr->n_constraints), current_saddle_point_state_{ - handle_ptr_, problem_ptr->n_variables, problem_ptr->n_constraints, climber_strategies.size()}, + handle_ptr_, problem_ptr->n_variables, problem_ptr->n_constraints, climber_strategies.size(), is_cupdlpx_restart(hyper_params)}, tmp_primal_{(climber_strategies.size() * problem_ptr->n_variables), stream_view_}, tmp_dual_{(climber_strategies.size() * problem_ptr->n_constraints), stream_view_}, potential_next_primal_solution_{(climber_strategies.size() * problem_ptr->n_variables), diff --git a/cpp/src/pdlp/saddle_point.cu b/cpp/src/pdlp/saddle_point.cu index 157e7fa389..e5f7a13da0 100644 --- a/cpp/src/pdlp/saddle_point.cu +++ b/cpp/src/pdlp/saddle_point.cu @@ -28,7 +28,7 @@ saddle_point_state_t::saddle_point_state_t(raft::handle_t const* handl delta_primal_{batch_size * primal_size, handle_ptr->get_stream()}, delta_dual_{batch_size * dual_size, handle_ptr->get_stream()}, // Primal gradient is only used in trust region restart mode which does not support batch mode - primal_gradient_{static_cast(primal_size), handle_ptr->get_stream()}, + primal_gradient_{is_cupdlpx ? 0 : static_cast(primal_size), handle_ptr->get_stream()}, dual_gradient_{batch_size * dual_size, handle_ptr->get_stream()}, current_AtY_{batch_size * primal_size, handle_ptr->get_stream()}, next_AtY_{batch_size * primal_size, handle_ptr->get_stream()} diff --git a/cpp/src/pdlp/saddle_point.hpp b/cpp/src/pdlp/saddle_point.hpp index 7e8f87fa25..bca633912b 100644 --- a/cpp/src/pdlp/saddle_point.hpp +++ b/cpp/src/pdlp/saddle_point.hpp @@ -64,7 +64,8 @@ class saddle_point_state_t { saddle_point_state_t(raft::handle_t const* handle_ptr, i_t primal_size, i_t dual_size, - size_t batch_size); + size_t batch_size, + const bool is_cupdlpx = false); /** * @brief Copies the values of the solutions in another saddle_point_state_t From 660a4ae2b3f4638cac8c7ae65d37f8195250a523 Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Thu, 23 Apr 2026 15:42:33 +0200 Subject: [PATCH 08/30] forgot to push the last change >_< --- cpp/src/pdlp/saddle_point.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/cpp/src/pdlp/saddle_point.cu b/cpp/src/pdlp/saddle_point.cu index e5f7a13da0..30077d8512 100644 --- a/cpp/src/pdlp/saddle_point.cu +++ b/cpp/src/pdlp/saddle_point.cu @@ -20,7 +20,8 @@ template saddle_point_state_t::saddle_point_state_t(raft::handle_t const* handle_ptr, const i_t primal_size, const i_t dual_size, - const size_t batch_size) + const size_t batch_size, + const bool is_cupdlpx) : primal_size_{primal_size}, dual_size_{dual_size}, primal_solution_{batch_size * primal_size, handle_ptr->get_stream()}, From abd82f37b704d54571de9fc4975d0e82728ca1b9 Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Tue, 28 Apr 2026 11:25:15 +0200 Subject: [PATCH 09/30] cleaned --- cpp/src/mip_heuristics/problem/problem.cu | 17 ++++-------- cpp/src/pdlp/cusparse_view.cu | 18 ++++++------- cpp/src/pdlp/cusparse_view.hpp | 3 ++- .../initial_scaling.cu | 12 ++------- cpp/src/pdlp/pdhg.cu | 7 +++-- cpp/src/pdlp/pdlp.cu | 26 +++++++------------ .../restart_strategy/pdlp_restart_strategy.cu | 8 +++--- cpp/src/pdlp/solve.cu | 2 -- 8 files changed, 36 insertions(+), 57 deletions(-) diff --git a/cpp/src/mip_heuristics/problem/problem.cu b/cpp/src/mip_heuristics/problem/problem.cu index f017fc3a46..7d55655d4c 100644 --- a/cpp/src/mip_heuristics/problem/problem.cu +++ b/cpp/src/mip_heuristics/problem/problem.cu @@ -2391,20 +2391,13 @@ void problem_t::update_variable_bounds(const std::vector& var_ind RAFT_CHECK_CUDA(handle_ptr->get_stream()); } -template -void poison_and_free(rmm::device_uvector& uvec, rmm::cuda_stream_view stream) { - if (uvec.size() > 0) { - thrust::fill(thrust::cuda::par_nosync.on(stream), uvec.begin(), uvec.end(), 5); - } - uvec.resize(0, stream); -} - template -void problem_t::pdlp_lighten(){ +void problem_t::pdlp_lighten() +{ variable_types.resize(0, handle_ptr->get_stream()); - poison_and_free(related_variables_offsets, handle_ptr->get_stream()); - poison_and_free(integer_fixed_variable_map, handle_ptr->get_stream()); - poison_and_free(combined_bounds, handle_ptr->get_stream()); + related_variables_offsets.resize(0, handle_ptr->get_stream()); + integer_fixed_variable_map.resize(0, handle_ptr->get_stream()); + combined_bounds.resize(0, handle_ptr->get_stream()); } #if MIP_INSTANTIATE_FLOAT || PDLP_INSTANTIATE_FLOAT diff --git a/cpp/src/pdlp/cusparse_view.cu b/cpp/src/pdlp/cusparse_view.cu index 22a3ef5b2d..e816d61a44 100644 --- a/cpp/src/pdlp/cusparse_view.cu +++ b/cpp/src/pdlp/cusparse_view.cu @@ -1083,13 +1083,12 @@ void cusparse_view_t::update_mixed_precision_matrices() } template -void cusparse_view_t::redirect_rows_and_cols( - const problem_t& original_problem) +void cusparse_view_t::redirect_rows_and_cols(const problem_t& original_problem) { RAFT_CUSPARSE_TRY(cusparseCsrSetPointers(A, - const_cast(original_problem.offsets.data()), - const_cast(original_problem.variables.data()), - const_cast(A_.data()))); + const_cast(original_problem.offsets.data()), + const_cast(original_problem.variables.data()), + const_cast(A_.data()))); RAFT_CUSPARSE_TRY( cusparseCsrSetPointers(A_T, @@ -1099,11 +1098,10 @@ void cusparse_view_t::redirect_rows_and_cols( if constexpr (std::is_same_v) { if (mixed_precision_enabled_) { - RAFT_CUSPARSE_TRY( - cusparseCsrSetPointers(A_mixed_, - const_cast(original_problem.offsets.data()), - const_cast(original_problem.variables.data()), - A_float_.data())); + RAFT_CUSPARSE_TRY(cusparseCsrSetPointers(A_mixed_, + const_cast(original_problem.offsets.data()), + const_cast(original_problem.variables.data()), + A_float_.data())); RAFT_CUSPARSE_TRY( cusparseCsrSetPointers(A_T_mixed_, diff --git a/cpp/src/pdlp/cusparse_view.hpp b/cpp/src/pdlp/cusparse_view.hpp index 81fcbd9ac3..eb919bbea9 100644 --- a/cpp/src/pdlp/cusparse_view.hpp +++ b/cpp/src/pdlp/cusparse_view.hpp @@ -209,7 +209,8 @@ class cusparse_view_t { // Update FP32 matrix copies after scaling (must be called after scale_problem()) void update_mixed_precision_matrices(); - // Redirects the vectors of rows and columns from op_problem_scaled_ to the original problem so that they can be freed + // Redirects the vectors of rows and columns from op_problem_scaled_ to the original problem so + // that they can be freed void redirect_rows_and_cols(const problem_t& original_problem); }; diff --git a/cpp/src/pdlp/initial_scaling_strategy/initial_scaling.cu b/cpp/src/pdlp/initial_scaling_strategy/initial_scaling.cu index 1a5108925d..d3a76ef2f1 100644 --- a/cpp/src/pdlp/initial_scaling_strategy/initial_scaling.cu +++ b/cpp/src/pdlp/initial_scaling_strategy/initial_scaling.cu @@ -27,14 +27,6 @@ #include namespace cuopt::linear_programming::detail { - -template -void poison_and_free(rmm::device_uvector& uvec, rmm::cuda_stream_view stream) { - if (uvec.size() > 0) { - thrust::fill(thrust::cuda::par_nosync.on(stream), uvec.begin(), uvec.end(), 5); - } - uvec.resize(0, stream); -} template pdlp_initial_scaling_strategy_t::pdlp_initial_scaling_strategy_t( raft::handle_t const* handle_ptr, @@ -89,8 +81,8 @@ pdlp_initial_scaling_strategy_t::pdlp_initial_scaling_strategy_t( compute_scaling_vectors(number_of_ruiz_iterations, alpha); - poison_and_free(iteration_constraint_matrix_scaling_, stream_view_); - poison_and_free(iteration_variable_scaling_, stream_view_); + iteration_constraint_matrix_scaling_.resize(0, stream_view_); + iteration_variable_scaling_.resize(0, stream_view_); } template diff --git a/cpp/src/pdlp/pdhg.cu b/cpp/src/pdlp/pdhg.cu index 910875ea03..0e3bf30edb 100644 --- a/cpp/src/pdlp/pdhg.cu +++ b/cpp/src/pdlp/pdhg.cu @@ -51,8 +51,11 @@ pdhg_solver_t::pdhg_solver_t( problem_ptr(&op_problem_scaled), primal_size_h_(problem_ptr->n_variables), dual_size_h_(problem_ptr->n_constraints), - current_saddle_point_state_{ - handle_ptr_, problem_ptr->n_variables, problem_ptr->n_constraints, climber_strategies.size(), is_cupdlpx_restart(hyper_params)}, + current_saddle_point_state_{handle_ptr_, + problem_ptr->n_variables, + problem_ptr->n_constraints, + climber_strategies.size(), + is_cupdlpx_restart(hyper_params)}, tmp_primal_{(climber_strategies.size() * problem_ptr->n_variables), stream_view_}, tmp_dual_{(climber_strategies.size() * problem_ptr->n_constraints), stream_view_}, potential_next_primal_solution_{(climber_strategies.size() * problem_ptr->n_variables), diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index 448b419202..f3d556b990 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -121,8 +121,10 @@ pdlp_solver_t::pdlp_solver_t(problem_t& op_problem, is_cupdlpx_(is_cupdlpx_restart(settings.hyper_params)), op_problem_scaled_( op_problem, false), // False to call the PDLP custom version of the problem copy constructor - unscaled_primal_avg_solution_{is_cupdlpx_ ? 0 : static_cast(op_problem.n_variables), stream_view_}, - unscaled_dual_avg_solution_{is_cupdlpx_ ? 0 : static_cast(op_problem.n_constraints), stream_view_}, + unscaled_primal_avg_solution_{is_cupdlpx_ ? 0 : static_cast(op_problem.n_variables), + stream_view_}, + unscaled_dual_avg_solution_{is_cupdlpx_ ? 0 : static_cast(op_problem.n_constraints), + stream_view_}, primal_size_h_(op_problem.n_variables), dual_size_h_(op_problem.n_constraints), primal_step_size_{climber_strategies_.size(), stream_view_}, @@ -2158,14 +2160,6 @@ void pdlp_solver_t::transpose_primal_dual_to_row( stream_view_); } -template -void poison_and_free(rmm::device_uvector& uvec, rmm::cuda_stream_view stream) { - if (uvec.size() > 0) { - thrust::fill(thrust::cuda::par_nosync.on(stream), uvec.begin(), uvec.end(), 5); - } - uvec.resize(0, stream); -} - template void pdlp_solver_t::transpose_primal_dual_back_to_col( rmm::device_uvector& primal_to_transpose, @@ -2274,12 +2268,12 @@ optimization_problem_solution_t pdlp_solver_t::run_solver(co // Redirect cuSPARSE descriptors to use the original problem's structural data (offsets, indices), // then free the duplicated structural vectors from the scaled copy to save device memory. pdhg_solver_.get_cusparse_view().redirect_rows_and_cols(*problem_ptr); - poison_and_free(op_problem_scaled_.variables, stream_view_); - poison_and_free(op_problem_scaled_.offsets, stream_view_); - poison_and_free(op_problem_scaled_.reverse_constraints, stream_view_); - poison_and_free(op_problem_scaled_.reverse_offsets, stream_view_); - - // Remove unused device vectors + op_problem_scaled_.variables.resize(0, stream_view_); + op_problem_scaled_.offsets.resize(0, stream_view_); + op_problem_scaled_.reverse_constraints.resize(0, stream_view_); + op_problem_scaled_.reverse_offsets.resize(0, stream_view_); + + // Remove unused device vectors in PDLP op_problem_scaled_.pdlp_lighten(); problem_ptr->pdlp_lighten(); diff --git a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu index b7c767d382..b09e87f36b 100644 --- a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu +++ b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu @@ -76,10 +76,10 @@ pdlp_restart_strategy_t::pdlp_restart_strategy_t( : handle_ptr_(handle_ptr), stream_view_(handle_ptr_->get_stream()), batch_mode_(climber_strategies.size() > 1), - weighted_average_solution_{handle_ptr_, - is_cupdlpx_restart(hyper_params) ? 0 : primal_size, - is_cupdlpx_restart(hyper_params) ? 0 : dual_size, - is_legacy_batch_mode}, + weighted_average_solution_{handle_ptr_, + is_cupdlpx_restart(hyper_params) ? 0 : primal_size, + is_cupdlpx_restart(hyper_params) ? 0 : dual_size, + is_legacy_batch_mode}, primal_size_h_(primal_size), dual_size_h_(dual_size), problem_ptr(&op_problem), diff --git a/cpp/src/pdlp/solve.cu b/cpp/src/pdlp/solve.cu index 64acf4a04c..030088e9ab 100644 --- a/cpp/src/pdlp/solve.cu +++ b/cpp/src/pdlp/solve.cu @@ -317,7 +317,6 @@ void set_pdlp_solver_mode(pdlp_solver_settings_t& settings) set_Fast1(settings.hyper_params); else if (settings.pdlp_solver_mode == pdlp_solver_mode_t::Stable3) { set_Stable3(settings.hyper_params); - settings.detect_infeasibility = false; } } @@ -1680,7 +1679,6 @@ cuopt::linear_programming::optimization_problem_t mps_data_model_to_op return op_problem; } -//HERE template optimization_problem_solution_t solve_lp( raft::handle_t const* handle_ptr, From 5e9d12b3cba418273b9834b2b8368e9f70130275 Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Tue, 28 Apr 2026 11:32:17 +0200 Subject: [PATCH 10/30] better --- cpp/src/pdlp/solve.cu | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/cpp/src/pdlp/solve.cu b/cpp/src/pdlp/solve.cu index 030088e9ab..f0345368a7 100644 --- a/cpp/src/pdlp/solve.cu +++ b/cpp/src/pdlp/solve.cu @@ -315,9 +315,8 @@ void set_pdlp_solver_mode(pdlp_solver_settings_t& settings) set_Methodical1(settings.hyper_params); else if (settings.pdlp_solver_mode == pdlp_solver_mode_t::Fast1) set_Fast1(settings.hyper_params); - else if (settings.pdlp_solver_mode == pdlp_solver_mode_t::Stable3) { + else if (settings.pdlp_solver_mode == pdlp_solver_mode_t::Stable3) set_Stable3(settings.hyper_params); - } } std::atomic global_concurrent_halt{0}; From f948df5fb354866587250baf2ca78f4bab260563 Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Tue, 28 Apr 2026 13:48:13 +0200 Subject: [PATCH 11/30] updated unscaled_primal/dual_avg_solution_ construction to also check for warm_start_data --- cpp/src/pdlp/pdlp.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index f3d556b990..9fa40d6963 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -121,9 +121,9 @@ pdlp_solver_t::pdlp_solver_t(problem_t& op_problem, is_cupdlpx_(is_cupdlpx_restart(settings.hyper_params)), op_problem_scaled_( op_problem, false), // False to call the PDLP custom version of the problem copy constructor - unscaled_primal_avg_solution_{is_cupdlpx_ ? 0 : static_cast(op_problem.n_variables), + unscaled_primal_avg_solution_{is_cupdlpx && !settings.get_pdlp_warm_start_data().is_populated() ? 0 : problem_size, stream_view_}, - unscaled_dual_avg_solution_{is_cupdlpx_ ? 0 : static_cast(op_problem.n_constraints), + unscaled_dual_avg_solution_{is_cupdlpx && !settings.get_pdlp_warm_start_data().is_populated() ? 0 : problem_size, stream_view_}, primal_size_h_(op_problem.n_variables), dual_size_h_(op_problem.n_constraints), From 798b51ee22de46f085aa3ef65d756c327f3f012e Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Tue, 28 Apr 2026 14:03:05 +0200 Subject: [PATCH 12/30] fixed use after free of combined bounds and ensure it is freed only if pre_constraint_residual is false --- cpp/src/mip_heuristics/problem/problem.cu | 1 - cpp/src/pdlp/pdlp.cu | 4 ++++ 2 files changed, 4 insertions(+), 1 deletion(-) diff --git a/cpp/src/mip_heuristics/problem/problem.cu b/cpp/src/mip_heuristics/problem/problem.cu index 7d55655d4c..789251188d 100644 --- a/cpp/src/mip_heuristics/problem/problem.cu +++ b/cpp/src/mip_heuristics/problem/problem.cu @@ -2397,7 +2397,6 @@ void problem_t::pdlp_lighten() variable_types.resize(0, handle_ptr->get_stream()); related_variables_offsets.resize(0, handle_ptr->get_stream()); integer_fixed_variable_map.resize(0, handle_ptr->get_stream()); - combined_bounds.resize(0, handle_ptr->get_stream()); } #if MIP_INSTANTIATE_FLOAT || PDLP_INSTANTIATE_FLOAT diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index 9fa40d6963..4af3efd0d7 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -2458,6 +2458,10 @@ optimization_problem_solution_t pdlp_solver_t::run_solver(co raft::print_device_vector("Initial primal_step_size", primal_step_size_.data(), 1, std::cout); raft::print_device_vector("Initial dual_step_size", dual_step_size_.data(), 1, std::cout); } + if (!per_constraint_residual){ + op_problem_scaled_.combined_bounds.resize(0, handle_ptr->get_stream()); + problem_ptr->combined_bounds.resize(0, handle_ptr->get_stream()); + } #ifdef CUPDLP_DEBUG_MODE raft::print_device_vector("Initial step_size", step_size_.data(), step_size_.size(), std::cout); raft::print_device_vector( From c1bbb42cda495328dee676b213396b0948d24c97 Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Tue, 28 Apr 2026 14:04:29 +0200 Subject: [PATCH 13/30] made pdlp_lighten internal --- cpp/src/mip_heuristics/problem/problem.cu | 8 -------- cpp/src/mip_heuristics/problem/problem.cuh | 1 - cpp/src/pdlp/pdlp.cu | 12 ++++++++++-- 3 files changed, 10 insertions(+), 11 deletions(-) diff --git a/cpp/src/mip_heuristics/problem/problem.cu b/cpp/src/mip_heuristics/problem/problem.cu index 789251188d..10d80586b4 100644 --- a/cpp/src/mip_heuristics/problem/problem.cu +++ b/cpp/src/mip_heuristics/problem/problem.cu @@ -2391,14 +2391,6 @@ void problem_t::update_variable_bounds(const std::vector& var_ind RAFT_CHECK_CUDA(handle_ptr->get_stream()); } -template -void problem_t::pdlp_lighten() -{ - variable_types.resize(0, handle_ptr->get_stream()); - related_variables_offsets.resize(0, handle_ptr->get_stream()); - integer_fixed_variable_map.resize(0, handle_ptr->get_stream()); -} - #if MIP_INSTANTIATE_FLOAT || PDLP_INSTANTIATE_FLOAT template class problem_t; #endif diff --git a/cpp/src/mip_heuristics/problem/problem.cuh b/cpp/src/mip_heuristics/problem/problem.cuh index 0a25bc63b0..a801cc4067 100644 --- a/cpp/src/mip_heuristics/problem/problem.cuh +++ b/cpp/src/mip_heuristics/problem/problem.cuh @@ -151,7 +151,6 @@ class problem_t { const std::vector& offset_values, const std::vector& coefficient_values); void sort_rows_by_variables(const raft::handle_t* handle_ptr); - void pdlp_lighten(); enum var_flags_t : i_t { VAR_IMPLIED_INTEGER = 1 << 0, }; diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index 4af3efd0d7..0ae5508565 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -108,6 +108,14 @@ static size_t batch_size_handler(const problem_t& op_problem, return settings.new_bounds.size(); } +template +static void lighten_problem_for_pdlp(problem_t& problem, rmm::cuda_stream_view stream) +{ + problem.variable_types.resize(0, stream); + problem.related_variables_offsets.resize(0, stream); + problem.integer_fixed_variable_map.resize(0, stream); +} + template pdlp_solver_t::pdlp_solver_t(problem_t& op_problem, pdlp_solver_settings_t const& settings, @@ -2274,8 +2282,8 @@ optimization_problem_solution_t pdlp_solver_t::run_solver(co op_problem_scaled_.reverse_offsets.resize(0, stream_view_); // Remove unused device vectors in PDLP - op_problem_scaled_.pdlp_lighten(); - problem_ptr->pdlp_lighten(); + lighten_problem_for_pdlp(op_problem_scaled_, stream_view_); + lighten_problem_for_pdlp(*problem_ptr, stream_view_); if (!settings_.hyper_params.compute_initial_step_size_before_scaling && !settings_.get_initial_step_size().has_value()) From 7c2e4148c0292c2e99e4469615b6b94042d2e192 Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Tue, 28 Apr 2026 14:06:01 +0200 Subject: [PATCH 14/30] minor fix --- cpp/src/pdlp/pdlp.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index 0ae5508565..d369ebf78f 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -129,9 +129,9 @@ pdlp_solver_t::pdlp_solver_t(problem_t& op_problem, is_cupdlpx_(is_cupdlpx_restart(settings.hyper_params)), op_problem_scaled_( op_problem, false), // False to call the PDLP custom version of the problem copy constructor - unscaled_primal_avg_solution_{is_cupdlpx && !settings.get_pdlp_warm_start_data().is_populated() ? 0 : problem_size, + unscaled_primal_avg_solution_{is_cupdlpx_ && !settings.get_pdlp_warm_start_data().is_populated() ? 0 : problem_size, stream_view_}, - unscaled_dual_avg_solution_{is_cupdlpx && !settings.get_pdlp_warm_start_data().is_populated() ? 0 : problem_size, + unscaled_dual_avg_solution_{is_cupdlpx_ && !settings.get_pdlp_warm_start_data().is_populated() ? 0 : problem_size, stream_view_}, primal_size_h_(op_problem.n_variables), dual_size_h_(op_problem.n_constraints), From 59b385f993573a8f376b783df849ff29ddf05e0a Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Tue, 28 Apr 2026 14:26:42 +0200 Subject: [PATCH 15/30] fixes again --- cpp/src/pdlp/cusparse_view.cu | 2 +- cpp/src/pdlp/pdlp.cu | 10 +++++----- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/cpp/src/pdlp/cusparse_view.cu b/cpp/src/pdlp/cusparse_view.cu index e816d61a44..3c5447def9 100644 --- a/cpp/src/pdlp/cusparse_view.cu +++ b/cpp/src/pdlp/cusparse_view.cu @@ -407,7 +407,7 @@ cusparse_view_t::cusparse_view_t( _tmp_primal.data(), CUSPARSE_ORDER_COL); - primal_gradient.create(op_problem_scaled.n_variables, + primal_gradient.create(current_saddle_point_state.get_primal_gradient().size(), // It is 0 in cupdlpx current_saddle_point_state.get_primal_gradient().data()); dual_gradient.create(op_problem_scaled.n_constraints, current_saddle_point_state.get_dual_gradient().data()); diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index d369ebf78f..566593e834 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -129,9 +129,9 @@ pdlp_solver_t::pdlp_solver_t(problem_t& op_problem, is_cupdlpx_(is_cupdlpx_restart(settings.hyper_params)), op_problem_scaled_( op_problem, false), // False to call the PDLP custom version of the problem copy constructor - unscaled_primal_avg_solution_{is_cupdlpx_ && !settings.get_pdlp_warm_start_data().is_populated() ? 0 : problem_size, + unscaled_primal_avg_solution_{is_cupdlpx_ && !settings.get_pdlp_warm_start_data().is_populated() ? 0 : static_cast(op_problem.n_variables), stream_view_}, - unscaled_dual_avg_solution_{is_cupdlpx_ && !settings.get_pdlp_warm_start_data().is_populated() ? 0 : problem_size, + unscaled_dual_avg_solution_{is_cupdlpx_ && !settings.get_pdlp_warm_start_data().is_populated() ? 0 : static_cast(op_problem.n_constraints), stream_view_}, primal_size_h_(op_problem.n_variables), dual_size_h_(op_problem.n_constraints), @@ -2466,9 +2466,9 @@ optimization_problem_solution_t pdlp_solver_t::run_solver(co raft::print_device_vector("Initial primal_step_size", primal_step_size_.data(), 1, std::cout); raft::print_device_vector("Initial dual_step_size", dual_step_size_.data(), 1, std::cout); } - if (!per_constraint_residual){ - op_problem_scaled_.combined_bounds.resize(0, handle_ptr->get_stream()); - problem_ptr->combined_bounds.resize(0, handle_ptr->get_stream()); + if (!settings_.per_constraint_residual){ + lighten_problem_for_pdlp(op_problem_scaled_, stream_view_); + lighten_problem_for_pdlp(*problem_ptr, stream_view_); } #ifdef CUPDLP_DEBUG_MODE raft::print_device_vector("Initial step_size", step_size_.data(), step_size_.size(), std::cout); From 84daee47fa7ed94ddf632f8123fc9c3452d92f70 Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Tue, 28 Apr 2026 14:31:03 +0200 Subject: [PATCH 16/30] linter --- cpp/src/pdlp/cusparse_view.cu | 5 +++-- cpp/src/pdlp/pdlp.cu | 10 +++++++--- 2 files changed, 10 insertions(+), 5 deletions(-) diff --git a/cpp/src/pdlp/cusparse_view.cu b/cpp/src/pdlp/cusparse_view.cu index 3c5447def9..bd912f541b 100644 --- a/cpp/src/pdlp/cusparse_view.cu +++ b/cpp/src/pdlp/cusparse_view.cu @@ -407,8 +407,9 @@ cusparse_view_t::cusparse_view_t( _tmp_primal.data(), CUSPARSE_ORDER_COL); - primal_gradient.create(current_saddle_point_state.get_primal_gradient().size(), // It is 0 in cupdlpx - current_saddle_point_state.get_primal_gradient().data()); + primal_gradient.create( + current_saddle_point_state.get_primal_gradient().size(), // It is 0 in cupdlpx + current_saddle_point_state.get_primal_gradient().data()); dual_gradient.create(op_problem_scaled.n_constraints, current_saddle_point_state.get_dual_gradient().data()); diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index 566593e834..b64d1e56e9 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -129,9 +129,13 @@ pdlp_solver_t::pdlp_solver_t(problem_t& op_problem, is_cupdlpx_(is_cupdlpx_restart(settings.hyper_params)), op_problem_scaled_( op_problem, false), // False to call the PDLP custom version of the problem copy constructor - unscaled_primal_avg_solution_{is_cupdlpx_ && !settings.get_pdlp_warm_start_data().is_populated() ? 0 : static_cast(op_problem.n_variables), + unscaled_primal_avg_solution_{is_cupdlpx_ && !settings.get_pdlp_warm_start_data().is_populated() + ? 0 + : static_cast(op_problem.n_variables), stream_view_}, - unscaled_dual_avg_solution_{is_cupdlpx_ && !settings.get_pdlp_warm_start_data().is_populated() ? 0 : static_cast(op_problem.n_constraints), + unscaled_dual_avg_solution_{is_cupdlpx_ && !settings.get_pdlp_warm_start_data().is_populated() + ? 0 + : static_cast(op_problem.n_constraints), stream_view_}, primal_size_h_(op_problem.n_variables), dual_size_h_(op_problem.n_constraints), @@ -2466,7 +2470,7 @@ optimization_problem_solution_t pdlp_solver_t::run_solver(co raft::print_device_vector("Initial primal_step_size", primal_step_size_.data(), 1, std::cout); raft::print_device_vector("Initial dual_step_size", dual_step_size_.data(), 1, std::cout); } - if (!settings_.per_constraint_residual){ + if (!settings_.per_constraint_residual) { lighten_problem_for_pdlp(op_problem_scaled_, stream_view_); lighten_problem_for_pdlp(*problem_ptr, stream_view_); } From edb13aea127d08cc4e87e5b43aa12b98422857cf Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Wed, 29 Apr 2026 16:37:28 +0200 Subject: [PATCH 17/30] nit coderabbit1 : more robust never_restart_to_average handling --- cpp/src/pdlp/pdlp.cu | 4 ++-- .../pdlp/restart_strategy/pdlp_restart_strategy.cu | 12 ++++++------ 2 files changed, 8 insertions(+), 8 deletions(-) diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index b64d1e56e9..985ff5a6b6 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -207,8 +207,8 @@ pdlp_solver_t::pdlp_solver_t(problem_t& op_problem, op_problem_scaled_, average_op_problem_evaluation_cusparse_view_, pdhg_solver_.get_cusparse_view(), - is_cupdlpx_ ? 0 : primal_size_h_, - is_cupdlpx_ ? 0 : dual_size_h_, + settings_.hyper_params.never_restart_to_average ? 0 : primal_size_h_, + settings_.hyper_params.never_restart_to_average ? 0 : dual_size_h_, initial_scaling_strategy_, settings_, climber_strategies_}, diff --git a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu index b09e87f36b..14561ff005 100644 --- a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu +++ b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu @@ -77,8 +77,8 @@ pdlp_restart_strategy_t::pdlp_restart_strategy_t( stream_view_(handle_ptr_->get_stream()), batch_mode_(climber_strategies.size() > 1), weighted_average_solution_{handle_ptr_, - is_cupdlpx_restart(hyper_params) ? 0 : primal_size, - is_cupdlpx_restart(hyper_params) ? 0 : dual_size, + hyper_params.never_restart_to_average ? 0 : primal_size, + hyper_params.never_restart_to_average ? 0 : dual_size, is_legacy_batch_mode}, primal_size_h_(primal_size), dual_size_h_(dual_size), @@ -92,13 +92,13 @@ pdlp_restart_strategy_t::pdlp_restart_strategy_t( restart_triggered_{0, stream_view_}, candidate_is_avg_{0, stream_view_}, avg_duality_gap_{handle_ptr_, - is_cupdlpx_restart(hyper_params) ? 0 : primal_size, - is_cupdlpx_restart(hyper_params) ? 0 : dual_size, + hyper_params.never_restart_to_average ? 0 : primal_size, + hyper_params.never_restart_to_average ? 0 : dual_size, climber_strategies, hyper_params}, current_duality_gap_{handle_ptr_, - is_cupdlpx_restart(hyper_params) ? 0 : primal_size, - is_cupdlpx_restart(hyper_params) ? 0 : dual_size, + hyper_params.never_restart_to_average ? 0 : primal_size, + hyper_params.never_restart_to_average ? 0 : dual_size, climber_strategies, hyper_params}, last_restart_duality_gap_{ From 92d6d76a5bf223e305342a7dbe28d461ce56d39a Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Wed, 29 Apr 2026 16:40:21 +0200 Subject: [PATCH 18/30] code rabbit 2-3 --- cpp/src/pdlp/pdlp.cu | 5 ----- 1 file changed, 5 deletions(-) diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index 985ff5a6b6..32f673456d 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -2287,7 +2287,6 @@ optimization_problem_solution_t pdlp_solver_t::run_solver(co // Remove unused device vectors in PDLP lighten_problem_for_pdlp(op_problem_scaled_, stream_view_); - lighten_problem_for_pdlp(*problem_ptr, stream_view_); if (!settings_.hyper_params.compute_initial_step_size_before_scaling && !settings_.get_initial_step_size().has_value()) @@ -2470,10 +2469,6 @@ optimization_problem_solution_t pdlp_solver_t::run_solver(co raft::print_device_vector("Initial primal_step_size", primal_step_size_.data(), 1, std::cout); raft::print_device_vector("Initial dual_step_size", dual_step_size_.data(), 1, std::cout); } - if (!settings_.per_constraint_residual) { - lighten_problem_for_pdlp(op_problem_scaled_, stream_view_); - lighten_problem_for_pdlp(*problem_ptr, stream_view_); - } #ifdef CUPDLP_DEBUG_MODE raft::print_device_vector("Initial step_size", step_size_.data(), step_size_.size(), std::cout); raft::print_device_vector( From cf541e9c18ce2dd3bad9efda30e9b2baad7a466c Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Wed, 29 Apr 2026 16:48:07 +0200 Subject: [PATCH 19/30] added comment fro is_cupdlpx --- cpp/src/pdlp/pdlp.cuh | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/cpp/src/pdlp/pdlp.cuh b/cpp/src/pdlp/pdlp.cuh index 9251b268e1..d59e63d259 100644 --- a/cpp/src/pdlp/pdlp.cuh +++ b/cpp/src/pdlp/pdlp.cuh @@ -143,11 +143,13 @@ class pdlp_solver_t { settings_.shared_sb_solved}; problem_t* problem_ptr; - bool is_cupdlpx_; // Combined bounds in op_problem_scaled_ will only be scaled if // compute_initial_primal_weight_before_scaling is false because of compute_initial_primal_weight problem_t op_problem_scaled_; + // Flag indicating whether the restart strategy is CUPDLPx (hyper_params.restart_strategy == static_cast(pdlp_restart_strategy_t::restart_strategy_t::CUPDLPX_RESTART)) + bool is_cupdlpx_; + rmm::device_uvector unscaled_primal_avg_solution_; rmm::device_uvector unscaled_dual_avg_solution_; From a426bf449c6c66b849f6890b897fe35fd7e9a673 Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Wed, 29 Apr 2026 17:27:23 +0200 Subject: [PATCH 20/30] removed default argument in saddle_point state and renamed it for more clarity --- cpp/src/pdlp/saddle_point.cu | 4 ++-- cpp/src/pdlp/saddle_point.hpp | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/cpp/src/pdlp/saddle_point.cu b/cpp/src/pdlp/saddle_point.cu index 30077d8512..e3685f7799 100644 --- a/cpp/src/pdlp/saddle_point.cu +++ b/cpp/src/pdlp/saddle_point.cu @@ -21,7 +21,7 @@ saddle_point_state_t::saddle_point_state_t(raft::handle_t const* handl const i_t primal_size, const i_t dual_size, const size_t batch_size, - const bool is_cupdlpx) + const bool need_primal_gradient) : primal_size_{primal_size}, dual_size_{dual_size}, primal_solution_{batch_size * primal_size, handle_ptr->get_stream()}, @@ -29,7 +29,7 @@ saddle_point_state_t::saddle_point_state_t(raft::handle_t const* handl delta_primal_{batch_size * primal_size, handle_ptr->get_stream()}, delta_dual_{batch_size * dual_size, handle_ptr->get_stream()}, // Primal gradient is only used in trust region restart mode which does not support batch mode - primal_gradient_{is_cupdlpx ? 0 : static_cast(primal_size), handle_ptr->get_stream()}, + primal_gradient_{need_primal_gradient ? static_cast(primal_size) : 0, handle_ptr->get_stream()}, dual_gradient_{batch_size * dual_size, handle_ptr->get_stream()}, current_AtY_{batch_size * primal_size, handle_ptr->get_stream()}, next_AtY_{batch_size * primal_size, handle_ptr->get_stream()} diff --git a/cpp/src/pdlp/saddle_point.hpp b/cpp/src/pdlp/saddle_point.hpp index bca633912b..5649703dfc 100644 --- a/cpp/src/pdlp/saddle_point.hpp +++ b/cpp/src/pdlp/saddle_point.hpp @@ -65,7 +65,7 @@ class saddle_point_state_t { i_t primal_size, i_t dual_size, size_t batch_size, - const bool is_cupdlpx = false); + const bool need_primal_gradient); /** * @brief Copies the values of the solutions in another saddle_point_state_t From 03fd3526bac7a013c25cdbebd1af351c5bd1f1f9 Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Wed, 29 Apr 2026 17:36:11 +0200 Subject: [PATCH 21/30] undid bug in current duality gap --- cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu index 14561ff005..fa53cc0da4 100644 --- a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu +++ b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu @@ -97,8 +97,8 @@ pdlp_restart_strategy_t::pdlp_restart_strategy_t( climber_strategies, hyper_params}, current_duality_gap_{handle_ptr_, - hyper_params.never_restart_to_average ? 0 : primal_size, - hyper_params.never_restart_to_average ? 0 : dual_size, + is_cupdlpx_restart(hyper_params) ? 0 : primal_size, + is_cupdlpx_restart(hyper_params) ? 0 : dual_size, climber_strategies, hyper_params}, last_restart_duality_gap_{ From 6fc6a24f98edffe5a484d075f2eb3f871840ac3f Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Thu, 30 Apr 2026 10:16:09 +0200 Subject: [PATCH 22/30] removed lighten_problem_for_pdlp --- cpp/src/mip_heuristics/problem/problem.cu | 10 +++++----- cpp/src/pdlp/pdlp.cu | 11 ----------- 2 files changed, 5 insertions(+), 16 deletions(-) diff --git a/cpp/src/mip_heuristics/problem/problem.cu b/cpp/src/mip_heuristics/problem/problem.cu index daaebe6c05..c9be19a431 100644 --- a/cpp/src/mip_heuristics/problem/problem.cu +++ b/cpp/src/mip_heuristics/problem/problem.cu @@ -278,7 +278,8 @@ problem_t::problem_t(const problem_t& problem_, bool no_deep deterministic(problem_.deterministic), handle_ptr(problem_.handle_ptr), integer_fixed_problem(problem_.integer_fixed_problem), - integer_fixed_variable_map(problem_.n_variables, handle_ptr->get_stream()), + integer_fixed_variable_map((!no_deep_copy) ? 0 : problem_.n_variables, + handle_ptr->get_stream()), n_variables(problem_.n_variables), n_constraints(problem_.n_constraints), n_binary_vars(problem_.n_binary_vars), @@ -343,9 +344,7 @@ problem_t::problem_t(const problem_t& problem_, bool no_deep ? rmm::device_uvector(problem_.combined_bounds, handle_ptr->get_stream()) : rmm::device_uvector(problem_.combined_bounds.size(), handle_ptr->get_stream())), variable_types( - (!no_deep_copy) - ? rmm::device_uvector(problem_.variable_types, handle_ptr->get_stream()) - : rmm::device_uvector(problem_.variable_types.size(), handle_ptr->get_stream())), + (!no_deep_copy) ? 0 : problem_.variable_types.size(), handle_ptr->get_stream()), integer_indices((!no_deep_copy) ? 0 : problem_.integer_indices.size(), handle_ptr->get_stream()), binary_indices((!no_deep_copy) ? 0 : problem_.binary_indices.size(), handle_ptr->get_stream()), @@ -354,7 +353,8 @@ problem_t::problem_t(const problem_t& problem_, bool no_deep is_binary_variable((!no_deep_copy) ? 0 : problem_.is_binary_variable.size(), handle_ptr->get_stream()), related_variables(problem_.related_variables, handle_ptr->get_stream()), - related_variables_offsets(problem_.related_variables_offsets, handle_ptr->get_stream()), + related_variables_offsets((!no_deep_copy) ? 0 : problem_.related_variables_offsets.size(), + handle_ptr->get_stream()), var_names(problem_.var_names), row_names(problem_.row_names), objective_name(problem_.objective_name), diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index 32f673456d..af6e839664 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -108,14 +108,6 @@ static size_t batch_size_handler(const problem_t& op_problem, return settings.new_bounds.size(); } -template -static void lighten_problem_for_pdlp(problem_t& problem, rmm::cuda_stream_view stream) -{ - problem.variable_types.resize(0, stream); - problem.related_variables_offsets.resize(0, stream); - problem.integer_fixed_variable_map.resize(0, stream); -} - template pdlp_solver_t::pdlp_solver_t(problem_t& op_problem, pdlp_solver_settings_t const& settings, @@ -2285,9 +2277,6 @@ optimization_problem_solution_t pdlp_solver_t::run_solver(co op_problem_scaled_.reverse_constraints.resize(0, stream_view_); op_problem_scaled_.reverse_offsets.resize(0, stream_view_); - // Remove unused device vectors in PDLP - lighten_problem_for_pdlp(op_problem_scaled_, stream_view_); - if (!settings_.hyper_params.compute_initial_step_size_before_scaling && !settings_.get_initial_step_size().has_value()) compute_initial_step_size(); From 453e6d9b1442c107198b3d19ba8d331a9e7ed723 Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Thu, 30 Apr 2026 16:07:18 +0200 Subject: [PATCH 23/30] renamed redirect_rows_and_cols to redirect_cusparse_csr_structure_pointers --- cpp/src/pdlp/cusparse_view.cu | 5 ++++- cpp/src/pdlp/cusparse_view.hpp | 6 +++--- cpp/src/pdlp/pdlp.cu | 2 +- 3 files changed, 8 insertions(+), 5 deletions(-) diff --git a/cpp/src/pdlp/cusparse_view.cu b/cpp/src/pdlp/cusparse_view.cu index bd912f541b..359bb7e928 100644 --- a/cpp/src/pdlp/cusparse_view.cu +++ b/cpp/src/pdlp/cusparse_view.cu @@ -1083,8 +1083,11 @@ void cusparse_view_t::update_mixed_precision_matrices() } } +// Redirects the cuSPARSE CSR structure pointers from op_problem_scaled_ to the original problem +// so the duplicated row/column buffers can be freed. template -void cusparse_view_t::redirect_rows_and_cols(const problem_t& original_problem) +void cusparse_view_t::redirect_cusparse_csr_structure_pointers( + const problem_t& original_problem) { RAFT_CUSPARSE_TRY(cusparseCsrSetPointers(A, const_cast(original_problem.offsets.data()), diff --git a/cpp/src/pdlp/cusparse_view.hpp b/cpp/src/pdlp/cusparse_view.hpp index eb919bbea9..c6d0ddea61 100644 --- a/cpp/src/pdlp/cusparse_view.hpp +++ b/cpp/src/pdlp/cusparse_view.hpp @@ -209,9 +209,9 @@ class cusparse_view_t { // Update FP32 matrix copies after scaling (must be called after scale_problem()) void update_mixed_precision_matrices(); - // Redirects the vectors of rows and columns from op_problem_scaled_ to the original problem so - // that they can be freed - void redirect_rows_and_cols(const problem_t& original_problem); + // Redirects the cuSPARSE CSR structure pointers from op_problem_scaled_ to the original problem + // so the duplicated row/column buffers can be freed. + void redirect_cusparse_csr_structure_pointers(const problem_t& original_problem); }; // Mixed precision SpMV: FP32 matrix with FP64 vectors and FP64 compute type diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index af6e839664..d9b00fa688 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -2271,7 +2271,7 @@ optimization_problem_solution_t pdlp_solver_t::run_solver(co // Redirect cuSPARSE descriptors to use the original problem's structural data (offsets, indices), // then free the duplicated structural vectors from the scaled copy to save device memory. - pdhg_solver_.get_cusparse_view().redirect_rows_and_cols(*problem_ptr); + pdhg_solver_.get_cusparse_view().redirect_cusparse_csr_structure_pointers(*problem_ptr); op_problem_scaled_.variables.resize(0, stream_view_); op_problem_scaled_.offsets.resize(0, stream_view_); op_problem_scaled_.reverse_constraints.resize(0, stream_view_); From 3dcccdcbf97a00c3a5abeef0255144c76433b5b7 Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Thu, 30 Apr 2026 16:10:44 +0200 Subject: [PATCH 24/30] changed const bool need_primal_gradient) in saddle_point_state constructor, to instead use hyper parameters --- cpp/src/pdlp/pdhg.cu | 2 +- cpp/src/pdlp/saddle_point.cu | 6 ++++-- cpp/src/pdlp/saddle_point.hpp | 4 +++- 3 files changed, 8 insertions(+), 4 deletions(-) diff --git a/cpp/src/pdlp/pdhg.cu b/cpp/src/pdlp/pdhg.cu index 0e3bf30edb..bc8819b99e 100644 --- a/cpp/src/pdlp/pdhg.cu +++ b/cpp/src/pdlp/pdhg.cu @@ -55,7 +55,7 @@ pdhg_solver_t::pdhg_solver_t( problem_ptr->n_variables, problem_ptr->n_constraints, climber_strategies.size(), - is_cupdlpx_restart(hyper_params)}, + hyper_params}, tmp_primal_{(climber_strategies.size() * problem_ptr->n_variables), stream_view_}, tmp_dual_{(climber_strategies.size() * problem_ptr->n_constraints), stream_view_}, potential_next_primal_solution_{(climber_strategies.size() * problem_ptr->n_variables), diff --git a/cpp/src/pdlp/saddle_point.cu b/cpp/src/pdlp/saddle_point.cu index e3685f7799..9301b2cc74 100644 --- a/cpp/src/pdlp/saddle_point.cu +++ b/cpp/src/pdlp/saddle_point.cu @@ -7,6 +7,7 @@ #include +#include #include #include @@ -21,7 +22,7 @@ saddle_point_state_t::saddle_point_state_t(raft::handle_t const* handl const i_t primal_size, const i_t dual_size, const size_t batch_size, - const bool need_primal_gradient) + const pdlp_hyper_params::pdlp_hyper_params_t& hyper_params) : primal_size_{primal_size}, dual_size_{dual_size}, primal_solution_{batch_size * primal_size, handle_ptr->get_stream()}, @@ -29,7 +30,8 @@ saddle_point_state_t::saddle_point_state_t(raft::handle_t const* handl delta_primal_{batch_size * primal_size, handle_ptr->get_stream()}, delta_dual_{batch_size * dual_size, handle_ptr->get_stream()}, // Primal gradient is only used in trust region restart mode which does not support batch mode - primal_gradient_{need_primal_gradient ? static_cast(primal_size) : 0, handle_ptr->get_stream()}, + primal_gradient_{!is_cupdlpx_restart(hyper_params) ? static_cast(primal_size) : 0, + handle_ptr->get_stream()}, dual_gradient_{batch_size * dual_size, handle_ptr->get_stream()}, current_AtY_{batch_size * primal_size, handle_ptr->get_stream()}, next_AtY_{batch_size * primal_size, handle_ptr->get_stream()} diff --git a/cpp/src/pdlp/saddle_point.hpp b/cpp/src/pdlp/saddle_point.hpp index 5649703dfc..eb6b8025cf 100644 --- a/cpp/src/pdlp/saddle_point.hpp +++ b/cpp/src/pdlp/saddle_point.hpp @@ -7,6 +7,8 @@ #pragma once +#include + #include #include @@ -65,7 +67,7 @@ class saddle_point_state_t { i_t primal_size, i_t dual_size, size_t batch_size, - const bool need_primal_gradient); + const pdlp_hyper_params::pdlp_hyper_params_t& hyper_params); /** * @brief Copies the values of the solutions in another saddle_point_state_t From 851be7c8f670ffc54bed634bab2b50e2ca0d90c0 Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Thu, 30 Apr 2026 16:15:40 +0200 Subject: [PATCH 25/30] removed is_cupdlpx bool --- cpp/src/pdlp/pdlp.cu | 7 ++++--- cpp/src/pdlp/pdlp.cuh | 3 --- 2 files changed, 4 insertions(+), 6 deletions(-) diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index d9b00fa688..c51ac9628d 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -118,14 +118,15 @@ pdlp_solver_t::pdlp_solver_t(problem_t& op_problem, stream_view_(handle_ptr_->get_stream()), settings_(settings), problem_ptr(&op_problem), - is_cupdlpx_(is_cupdlpx_restart(settings.hyper_params)), op_problem_scaled_( op_problem, false), // False to call the PDLP custom version of the problem copy constructor - unscaled_primal_avg_solution_{is_cupdlpx_ && !settings.get_pdlp_warm_start_data().is_populated() + unscaled_primal_avg_solution_{is_cupdlpx_restart(settings.hyper_params) && + !settings.get_pdlp_warm_start_data().is_populated() ? 0 : static_cast(op_problem.n_variables), stream_view_}, - unscaled_dual_avg_solution_{is_cupdlpx_ && !settings.get_pdlp_warm_start_data().is_populated() + unscaled_dual_avg_solution_{is_cupdlpx_restart(settings.hyper_params) && + !settings.get_pdlp_warm_start_data().is_populated() ? 0 : static_cast(op_problem.n_constraints), stream_view_}, diff --git a/cpp/src/pdlp/pdlp.cuh b/cpp/src/pdlp/pdlp.cuh index d59e63d259..68b4ac4590 100644 --- a/cpp/src/pdlp/pdlp.cuh +++ b/cpp/src/pdlp/pdlp.cuh @@ -146,9 +146,6 @@ class pdlp_solver_t { // Combined bounds in op_problem_scaled_ will only be scaled if // compute_initial_primal_weight_before_scaling is false because of compute_initial_primal_weight problem_t op_problem_scaled_; - - // Flag indicating whether the restart strategy is CUPDLPx (hyper_params.restart_strategy == static_cast(pdlp_restart_strategy_t::restart_strategy_t::CUPDLPX_RESTART)) - bool is_cupdlpx_; rmm::device_uvector unscaled_primal_avg_solution_; rmm::device_uvector unscaled_dual_avg_solution_; From 8f9a937a2e1ae6a9151cdff5a524956e3dd9bcbd Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Mon, 4 May 2026 11:55:09 +0200 Subject: [PATCH 26/30] pre-commit-run --- cpp/src/mip_heuristics/problem/problem.cu | 3 +-- cpp/src/pdlp/pdlp.cu | 21 ++++++++++--------- cpp/src/pdlp/pdlp.cuh | 2 +- .../restart_strategy/pdlp_restart_strategy.cu | 4 ++-- cpp/src/pdlp/saddle_point.cu | 16 +++++++------- 5 files changed, 24 insertions(+), 22 deletions(-) diff --git a/cpp/src/mip_heuristics/problem/problem.cu b/cpp/src/mip_heuristics/problem/problem.cu index c9be19a431..5de427f18a 100644 --- a/cpp/src/mip_heuristics/problem/problem.cu +++ b/cpp/src/mip_heuristics/problem/problem.cu @@ -343,8 +343,7 @@ problem_t::problem_t(const problem_t& problem_, bool no_deep (!no_deep_copy) ? rmm::device_uvector(problem_.combined_bounds, handle_ptr->get_stream()) : rmm::device_uvector(problem_.combined_bounds.size(), handle_ptr->get_stream())), - variable_types( - (!no_deep_copy) ? 0 : problem_.variable_types.size(), handle_ptr->get_stream()), + variable_types((!no_deep_copy) ? 0 : problem_.variable_types.size(), handle_ptr->get_stream()), integer_indices((!no_deep_copy) ? 0 : problem_.integer_indices.size(), handle_ptr->get_stream()), binary_indices((!no_deep_copy) ? 0 : problem_.binary_indices.size(), handle_ptr->get_stream()), diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index c51ac9628d..b4c161f0fd 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -195,16 +195,17 @@ pdlp_solver_t::pdlp_solver_t(problem_t& op_problem, is_legacy_batch_mode, climber_strategies_, settings_.hyper_params}, - average_termination_strategy_{handle_ptr_, - op_problem, - op_problem_scaled_, - average_op_problem_evaluation_cusparse_view_, - pdhg_solver_.get_cusparse_view(), - settings_.hyper_params.never_restart_to_average ? 0 : primal_size_h_, - settings_.hyper_params.never_restart_to_average ? 0 : dual_size_h_, - initial_scaling_strategy_, - settings_, - climber_strategies_}, + average_termination_strategy_{ + handle_ptr_, + op_problem, + op_problem_scaled_, + average_op_problem_evaluation_cusparse_view_, + pdhg_solver_.get_cusparse_view(), + settings_.hyper_params.never_restart_to_average ? 0 : primal_size_h_, + settings_.hyper_params.never_restart_to_average ? 0 : dual_size_h_, + initial_scaling_strategy_, + settings_, + climber_strategies_}, current_termination_strategy_{handle_ptr_, op_problem, op_problem_scaled_, diff --git a/cpp/src/pdlp/pdlp.cuh b/cpp/src/pdlp/pdlp.cuh index 68b4ac4590..d03430f150 100644 --- a/cpp/src/pdlp/pdlp.cuh +++ b/cpp/src/pdlp/pdlp.cuh @@ -146,7 +146,7 @@ class pdlp_solver_t { // Combined bounds in op_problem_scaled_ will only be scaled if // compute_initial_primal_weight_before_scaling is false because of compute_initial_primal_weight problem_t op_problem_scaled_; - + rmm::device_uvector unscaled_primal_avg_solution_; rmm::device_uvector unscaled_dual_avg_solution_; diff --git a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu index fa53cc0da4..f968b0f027 100644 --- a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu +++ b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu @@ -77,8 +77,8 @@ pdlp_restart_strategy_t::pdlp_restart_strategy_t( stream_view_(handle_ptr_->get_stream()), batch_mode_(climber_strategies.size() > 1), weighted_average_solution_{handle_ptr_, - hyper_params.never_restart_to_average ? 0 : primal_size, - hyper_params.never_restart_to_average ? 0 : dual_size, + hyper_params.never_restart_to_average ? 0 : primal_size, + hyper_params.never_restart_to_average ? 0 : dual_size, is_legacy_batch_mode}, primal_size_h_(primal_size), dual_size_h_(dual_size), diff --git a/cpp/src/pdlp/saddle_point.cu b/cpp/src/pdlp/saddle_point.cu index 9301b2cc74..f740176a3c 100644 --- a/cpp/src/pdlp/saddle_point.cu +++ b/cpp/src/pdlp/saddle_point.cu @@ -18,11 +18,12 @@ namespace cuopt::linear_programming::detail { template -saddle_point_state_t::saddle_point_state_t(raft::handle_t const* handle_ptr, - const i_t primal_size, - const i_t dual_size, - const size_t batch_size, - const pdlp_hyper_params::pdlp_hyper_params_t& hyper_params) +saddle_point_state_t::saddle_point_state_t( + raft::handle_t const* handle_ptr, + const i_t primal_size, + const i_t dual_size, + const size_t batch_size, + const pdlp_hyper_params::pdlp_hyper_params_t& hyper_params) : primal_size_{primal_size}, dual_size_{dual_size}, primal_solution_{batch_size * primal_size, handle_ptr->get_stream()}, @@ -30,8 +31,9 @@ saddle_point_state_t::saddle_point_state_t(raft::handle_t const* handl delta_primal_{batch_size * primal_size, handle_ptr->get_stream()}, delta_dual_{batch_size * dual_size, handle_ptr->get_stream()}, // Primal gradient is only used in trust region restart mode which does not support batch mode - primal_gradient_{!is_cupdlpx_restart(hyper_params) ? static_cast(primal_size) : 0, - handle_ptr->get_stream()}, + primal_gradient_{ + !is_cupdlpx_restart(hyper_params) ? static_cast(primal_size) : 0, + handle_ptr->get_stream()}, dual_gradient_{batch_size * dual_size, handle_ptr->get_stream()}, current_AtY_{batch_size * primal_size, handle_ptr->get_stream()}, next_AtY_{batch_size * primal_size, handle_ptr->get_stream()} From 29bd4ca218d3742b0bc86095f4018702a866e47a Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Mon, 4 May 2026 14:04:30 +0200 Subject: [PATCH 27/30] revrted back some vectors to full size because non-trivial to remove --- cpp/src/pdlp/pdlp.cu | 10 ++-------- cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu | 4 ++-- 2 files changed, 4 insertions(+), 10 deletions(-) diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index b4c161f0fd..0695eb6182 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -120,15 +120,9 @@ pdlp_solver_t::pdlp_solver_t(problem_t& op_problem, problem_ptr(&op_problem), op_problem_scaled_( op_problem, false), // False to call the PDLP custom version of the problem copy constructor - unscaled_primal_avg_solution_{is_cupdlpx_restart(settings.hyper_params) && - !settings.get_pdlp_warm_start_data().is_populated() - ? 0 - : static_cast(op_problem.n_variables), + unscaled_primal_avg_solution_{static_cast(op_problem.n_variables), stream_view_}, - unscaled_dual_avg_solution_{is_cupdlpx_restart(settings.hyper_params) && - !settings.get_pdlp_warm_start_data().is_populated() - ? 0 - : static_cast(op_problem.n_constraints), + unscaled_dual_avg_solution_{static_cast(op_problem.n_constraints), stream_view_}, primal_size_h_(op_problem.n_variables), dual_size_h_(op_problem.n_constraints), diff --git a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu index f968b0f027..28f5913ccf 100644 --- a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu +++ b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu @@ -77,8 +77,8 @@ pdlp_restart_strategy_t::pdlp_restart_strategy_t( stream_view_(handle_ptr_->get_stream()), batch_mode_(climber_strategies.size() > 1), weighted_average_solution_{handle_ptr_, - hyper_params.never_restart_to_average ? 0 : primal_size, - hyper_params.never_restart_to_average ? 0 : dual_size, + primal_size, + dual_size, is_legacy_batch_mode}, primal_size_h_(primal_size), dual_size_h_(dual_size), From 72afe0a00b3915e74837b073b9d236d7c053e157 Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Mon, 4 May 2026 14:12:45 +0200 Subject: [PATCH 28/30] style --- cpp/src/pdlp/pdlp.cu | 6 ++---- cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu | 5 +---- 2 files changed, 3 insertions(+), 8 deletions(-) diff --git a/cpp/src/pdlp/pdlp.cu b/cpp/src/pdlp/pdlp.cu index 0695eb6182..55441f76be 100644 --- a/cpp/src/pdlp/pdlp.cu +++ b/cpp/src/pdlp/pdlp.cu @@ -120,10 +120,8 @@ pdlp_solver_t::pdlp_solver_t(problem_t& op_problem, problem_ptr(&op_problem), op_problem_scaled_( op_problem, false), // False to call the PDLP custom version of the problem copy constructor - unscaled_primal_avg_solution_{static_cast(op_problem.n_variables), - stream_view_}, - unscaled_dual_avg_solution_{static_cast(op_problem.n_constraints), - stream_view_}, + unscaled_primal_avg_solution_{static_cast(op_problem.n_variables), stream_view_}, + unscaled_dual_avg_solution_{static_cast(op_problem.n_constraints), stream_view_}, primal_size_h_(op_problem.n_variables), dual_size_h_(op_problem.n_constraints), primal_step_size_{climber_strategies_.size(), stream_view_}, diff --git a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu index 28f5913ccf..682a544911 100644 --- a/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu +++ b/cpp/src/pdlp/restart_strategy/pdlp_restart_strategy.cu @@ -76,10 +76,7 @@ pdlp_restart_strategy_t::pdlp_restart_strategy_t( : handle_ptr_(handle_ptr), stream_view_(handle_ptr_->get_stream()), batch_mode_(climber_strategies.size() > 1), - weighted_average_solution_{handle_ptr_, - primal_size, - dual_size, - is_legacy_batch_mode}, + weighted_average_solution_{handle_ptr_, primal_size, dual_size, is_legacy_batch_mode}, primal_size_h_(primal_size), dual_size_h_(dual_size), problem_ptr(&op_problem), From cea749f3ae3828cc6468a1446f5cdd4be836a304 Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Tue, 5 May 2026 11:57:15 +0200 Subject: [PATCH 29/30] bump max-allowed-size-compressed --- ci/validate_wheel.sh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ci/validate_wheel.sh b/ci/validate_wheel.sh index 4d6180637f..c4731b7aaf 100755 --- a/ci/validate_wheel.sh +++ b/ci/validate_wheel.sh @@ -26,7 +26,7 @@ if [[ "${package_dir}" == "python/libcuopt" ]]; then ) else PYDISTCHECK_ARGS+=( - --max-allowed-size-compressed '505Mi' + --max-allowed-size-compressed '550Mi' ) fi elif [[ "${package_dir}" != "python/cuopt" ]] && \ From 1e2b582ac8d70f94b103c33019c0f064bb334eee Mon Sep 17 00:00:00 2001 From: Bulle Mostovoi Date: Sun, 10 May 2026 18:40:39 +0200 Subject: [PATCH 30/30] empty commit