Skip to content

Load balanced based multi probe #390

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 25 commits into
base: branch-25.10
Choose a base branch
from
Open

Load balanced based multi probe #390

Show file tree
Hide file tree
Changes from 24 commits
Commits
Show all changes
25 commits
Select commit Hold shift + click to select a range
8b5cdfd
in progress. removing variable_lower/upper bounds
kaatish Sep 2, 2025
2035c9a
in progress
kaatish Sep 3, 2025
8df30a6
in progress
kaatish Sep 3, 2025
4987b67
debug
kaatish Sep 3, 2025
d6e48ec
fixed
kaatish Sep 4, 2025
96eaf24
remove dead code
kaatish Sep 4, 2025
77c94d8
timer fix
kaatish Sep 4, 2025
12e131d
:Merge branch 'branch-25.10' of github.com:NVIDIA/cuopt into combined…
kaatish Sep 4, 2025
cacd567
pr review fixes
kaatish Sep 5, 2025
e062f51
add back sync
kaatish Sep 5, 2025
f403cbb
multi probe files
kaatish Sep 6, 2025
4970cca
comp
kaatish Sep 7, 2025
439ebf5
debug
kaatish Sep 7, 2025
d490307
h100 debug
kaatish Sep 8, 2025
5272f72
fixed
kaatish Sep 8, 2025
af959b1
constraint prop
kaatish Sep 9, 2025
d77bfb2
Merge branch 'branch-25.10' of github.com:NVIDIA/cuopt into lb_multi_…
kaatish Sep 9, 2025
c4ba103
add unique_ptr to problem
kaatish Sep 9, 2025
02b7c58
integration wip
kaatish Sep 10, 2025
f965358
fix multi_probe
kaatish Sep 13, 2025
360d11c
integrate lb probe and bounds with problem
kaatish Sep 14, 2025
bafeb6c
lb integrated
kaatish Sep 15, 2025
4b7aa8c
Merge branch 'branch-25.10' of github.com:NVIDIA/cuopt into lb_multi_…
kaatish Sep 15, 2025
40b4dbe
cleanup test
kaatish Sep 15, 2025
2a84cc7
fix build
kaatish Sep 15, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 6 additions & 0 deletions cpp/src/mip/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,9 +42,15 @@ set(MIP_NON_LP_FILES
${CMAKE_CURRENT_SOURCE_DIR}/presolve/bounds_presolve.cu
${CMAKE_CURRENT_SOURCE_DIR}/presolve/bounds_update_data.cu
${CMAKE_CURRENT_SOURCE_DIR}/presolve/conditional_bound_strengthening.cu
${CMAKE_CURRENT_SOURCE_DIR}/presolve/lb_bounds_presolve.cu
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Trying to understand the structure: we are using lb versions throughout the code and want to phase-out the non-lb version, right? If that's the case, I would put the non-lb versions under legacy dir to be removed completely.

${CMAKE_CURRENT_SOURCE_DIR}/presolve/lb_bounds_update_data.cu
${CMAKE_CURRENT_SOURCE_DIR}/presolve/lb_problem.cu
${CMAKE_CURRENT_SOURCE_DIR}/presolve/lb_multi_probe.cu
${CMAKE_CURRENT_SOURCE_DIR}/presolve/load_balanced_bounds_presolve.cu
${CMAKE_CURRENT_SOURCE_DIR}/presolve/multi_probe.cu
${CMAKE_CURRENT_SOURCE_DIR}/presolve/probing_cache.cu
${CMAKE_CURRENT_SOURCE_DIR}/presolve/trivial_presolve.cu
${CMAKE_CURRENT_SOURCE_DIR}/problem/load_balanced_problem.cu
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It is a bit confusing to have lb_problem.cu and load_balanced_problem.cu

${CMAKE_CURRENT_SOURCE_DIR}/feasibility_jump/feasibility_jump.cu
${CMAKE_CURRENT_SOURCE_DIR}/feasibility_jump/feasibility_jump_kernels.cu
)
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/mip/diversity/diversity_manager.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -411,7 +411,7 @@ solution_t<i_t, f_t> diversity_manager_t<i_t, f_t>::run_solver()

// test problem is not ii
cuopt_func_call(
ls.constraint_prop.bounds_update.calculate_activity_on_problem_bounds(*problem_ptr));
ls.constraint_prop.bounds_update.calculate_constraint_slack_on_problem_bounds(*problem_ptr));
cuopt_assert(
ls.constraint_prop.bounds_update.calculate_infeasible_redundant_constraints(*problem_ptr),
"The problem must not be ii");
Expand Down
45 changes: 26 additions & 19 deletions cpp/src/mip/local_search/rounding/bounds_repair.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,7 @@ namespace cuopt::linear_programming::detail {

template <typename i_t, typename f_t>
bounds_repair_t<i_t, f_t>::bounds_repair_t(const problem_t<i_t, f_t>& pb,
bound_presolve_t<i_t, f_t>& bound_presolve_)
lb_bound_presolve_t<i_t, f_t>& bound_presolve_)
: bound_presolve(bound_presolve_),
candidates(pb.handle_ptr),
best_bounds(pb.handle_ptr),
Expand Down Expand Up @@ -67,25 +67,25 @@ void bounds_repair_t<i_t, f_t>::reset()
template <typename i_t, typename f_t>
f_t bounds_repair_t<i_t, f_t>::get_ii_violation(problem_t<i_t, f_t>& problem)
{
bound_presolve.calculate_activity_on_problem_bounds(problem);
bound_presolve.calculate_constraint_slack_on_problem_bounds(problem);
// calculate the violation and mark of violated constraints
thrust::for_each(
handle_ptr->get_thrust_policy(),
thrust::make_counting_iterator(0),
thrust::make_counting_iterator(0) + problem.n_constraints,
[pb_v = problem.view(),
violated_cstr_map = violated_cstr_map.data(),
min_act = bound_presolve.upd.min_activity.data(),
max_act = bound_presolve.upd.max_activity.data(),
cnst_slack = make_span(bound_presolve.upd.cnst_slack),
cstr_violations_up = cstr_violations_up.data(),
cstr_violations_down = cstr_violations_down.data(),
total_vio = total_vio.data()] __device__(i_t cstr_idx) {
f_t cnst_lb = pb_v.constraint_lower_bounds[cstr_idx];
f_t cnst_ub = pb_v.constraint_upper_bounds[cstr_idx];
f_t eps = get_cstr_tolerance<i_t, f_t>(
cnst_lb, cnst_ub, pb_v.tolerances.absolute_tolerance, pb_v.tolerances.relative_tolerance);
f_t curr_cstr_violation_up = max(0., min_act[cstr_idx] - (cnst_ub + eps));
f_t curr_cstr_violation_down = max(0., cnst_lb - eps - max_act[cstr_idx]);
auto slack = cnst_slack[cstr_idx];
f_t curr_cstr_violation_up = max(0., -get_lower(slack) - eps);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Does the slack computation go negative or is there a min(0, slack) check? If it doesn't go negative this will not work. If it can go negative, we need to check other places if there is an implicit assumption that the slack is non-negative.

f_t curr_cstr_violation_down = max(0., get_upper(slack) - eps);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Isn't max_act - cnst_lb == upper_slack ?

f_t violation = max(curr_cstr_violation_up, curr_cstr_violation_down);
if (violation >= ROUNDOFF_TOLERANCE) {
violated_cstr_map[cstr_idx] = 1;
Expand Down Expand Up @@ -208,8 +208,7 @@ __global__ void compute_damages_kernel(typename problem_t<i_t, f_t>::view_t prob
typename candidates_t<i_t, f_t>::view_t candidates,
raft::device_span<f_t> cstr_violations_up,
raft::device_span<f_t> cstr_violations_down,
raft::device_span<f_t> minimum_activity,
raft::device_span<f_t> maximum_activity)
raft::device_span<typename type_2<f_t>::type> cnst_slack)
{
i_t var_idx = candidates.variable_index[blockIdx.x];
f_t shift_amount = candidates.bound_shift[blockIdx.x];
Expand All @@ -233,14 +232,23 @@ __global__ void compute_damages_kernel(typename problem_t<i_t, f_t>::view_t prob
f_t cnst_lb = problem.constraint_lower_bounds[c];
f_t cnst_ub = problem.constraint_upper_bounds[c];
f_t shift_in_activities = shift_amount * coeff;
f_t new_min_act = minimum_activity[c] + shift_in_activities;
f_t new_max_act = maximum_activity[c] + shift_in_activities;
auto slack = cnst_slack[c];
f_t eps = get_cstr_tolerance<i_t, f_t>(cnst_lb,
cnst_ub,
problem.tolerances.absolute_tolerance,
problem.tolerances.relative_tolerance);
f_t new_violations_up = max(0., new_min_act - (cnst_ub + eps));
f_t new_violations_down = max(0., cnst_lb - eps - new_max_act);
// Given
// f_t new_min_act = minimum_activity[c] + shift_in_activities;
// f_t new_max_act = maximum_activity[c] + shift_in_activities;
//
// f_t new_violations_up = max(0., new_min_act - (cnst_ub + eps));
// f_t new_violations_down = max(0., cnst_lb - eps - new_max_act);
// becomes
// f_t new_violations_up = max(0., shift_in_activities - (cnst_ub - min_act) - eps);
// f_t new_violations_down = max(0., (cnst_lb - max_act) - shift_in_activities - eps);
// becomes
f_t new_violations_up = max(0., shift_in_activities - get_lower(slack) - eps);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

note: This is valid given that slack calculation allows negative.

f_t new_violations_down = max(0., get_upper(slack) - shift_in_activities - eps);
f_t new_vio = max(new_violations_up, new_violations_down);
i_t curr_cstr_delta = i_t(curr_vio < ROUNDOFF_TOLERANCE) - i_t(new_vio < ROUNDOFF_TOLERANCE);
n_infeasible_cstr_delta += curr_cstr_delta;
Expand All @@ -262,13 +270,12 @@ void bounds_repair_t<i_t, f_t>::compute_damages(problem_t<i_t, f_t>& problem, i_
CUOPT_LOG_TRACE("Bounds repair: Computing damanges!");
// TODO check performance, we can apply load balancing here
const i_t TPB = 256;
compute_damages_kernel<i_t, f_t><<<n_candidates, TPB, 0, handle_ptr->get_stream()>>>(
problem.view(),
candidates.view(),
make_span(cstr_violations_up),
make_span(cstr_violations_down),
make_span(bound_presolve.upd.min_activity),
make_span(bound_presolve.upd.max_activity));
compute_damages_kernel<i_t, f_t>
<<<n_candidates, TPB, 0, handle_ptr->get_stream()>>>(problem.view(),
candidates.view(),
make_span(cstr_violations_up),
make_span(cstr_violations_down),
make_span(bound_presolve.upd.cnst_slack));
RAFT_CHECK_CUDA(handle_ptr->get_stream());
auto sort_iterator = thrust::make_zip_iterator(
thrust::make_tuple(candidates.cstr_delta.data(), candidates.damage.data()));
Expand Down
40 changes: 13 additions & 27 deletions cpp/src/mip/local_search/rounding/bounds_repair.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,7 @@

#pragma once

#include <mip/presolve/bounds_presolve.cuh>
#include <mip/presolve/lb_bounds_presolve.cuh>
#include <mip/problem/problem.cuh>
#include <mip/solution/solution.cuh>
#include <utilities/copy_helpers.hpp>
Expand All @@ -32,39 +32,25 @@ constexpr int MAX_CYCLE_SEQUENCE = 5;

template <typename i_t, typename f_t>
struct bounds_t {
bounds_t(const raft::handle_t* handle_ptr)
: lb(0, handle_ptr->get_stream()), ub(0, handle_ptr->get_stream())
{
}
bounds_t(const raft::handle_t* handle_ptr) : bounds(0, handle_ptr->get_stream()) {}
void resize(i_t var_size, const raft::handle_t* handle_ptr)
{
lb.resize(var_size, handle_ptr->get_stream());
ub.resize(var_size, handle_ptr->get_stream());
bounds.resize(var_size, handle_ptr->get_stream());
}
void update_from(const problem_t<i_t, f_t>& pb, const raft::handle_t* handle_ptr)
{
cuopt_assert(lb.size() == pb.variable_bounds.size(), "");
cuopt_assert(ub.size() == pb.variable_bounds.size(), "");
thrust::transform(
handle_ptr->get_thrust_policy(),
pb.variable_bounds.begin(),
pb.variable_bounds.end(),
thrust::make_zip_iterator(thrust::make_tuple(lb.begin(), ub.begin())),
[] __device__(auto i) { return thrust::make_tuple(get_lower(i), get_upper(i)); });
cuopt_assert(bounds.size() == pb.variable_bounds.size(), "");
raft::copy(bounds.data(),
pb.variable_bounds.data(),
pb.variable_bounds.size(),
handle_ptr->get_stream());
};
void update_to(problem_t<i_t, f_t>& pb, const raft::handle_t* handle_ptr)
{
cuopt_assert(lb.size() == pb.variable_bounds.size(), "");
cuopt_assert(ub.size() == pb.variable_bounds.size(), "");
using f_t2 = typename type_2<f_t>::type;
thrust::transform(handle_ptr->get_thrust_policy(),
thrust::make_zip_iterator(thrust::make_tuple(lb.begin(), ub.begin())),
thrust::make_zip_iterator(thrust::make_tuple(lb.end(), ub.end())),
pb.variable_bounds.begin(),
[] __device__(auto i) { return f_t2{thrust::get<0>(i), thrust::get<1>(i)}; });
cuopt_assert(bounds.size() == pb.variable_bounds.size(), "");
raft::copy(pb.variable_bounds.data(), bounds.data(), bounds.size(), handle_ptr->get_stream());
};
rmm::device_uvector<f_t> lb;
rmm::device_uvector<f_t> ub;
rmm::device_uvector<typename type_2<f_t>::type> bounds;
};

template <typename i_t, typename f_t>
Expand Down Expand Up @@ -118,7 +104,7 @@ struct candidates_t {
template <typename i_t, typename f_t>
class bounds_repair_t {
public:
bounds_repair_t(const problem_t<i_t, f_t>& p, bound_presolve_t<i_t, f_t>& bound_presolve);
bounds_repair_t(const problem_t<i_t, f_t>& p, lb_bound_presolve_t<i_t, f_t>& bound_presolve);
void resize(const problem_t<i_t, f_t>& problem);
void reset();
f_t get_ii_violation(problem_t<i_t, f_t>& problem);
Expand All @@ -141,7 +127,7 @@ class bounds_repair_t {
f_t best_damage,
i_t n_candidates);

bound_presolve_t<i_t, f_t>& bound_presolve;
lb_bound_presolve_t<i_t, f_t>& bound_presolve;
candidates_t<i_t, f_t> candidates;
bounds_t<i_t, f_t> best_bounds;
rmm::device_uvector<f_t> cstr_violations_up;
Expand Down
Loading
Loading