Incorporate determinism into GPU heuristics + parallel determinsitic B&B+GPU exploration - #986
Incorporate determinism into GPU heuristics + parallel determinsitic B&B+GPU exploration#986aliceb-nv wants to merge 307 commits into
Conversation
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 5
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (1)
cpp/src/branch_and_bound/pseudo_costs.cpp (1)
1184-1202:⚠️ Potential issue | 🟠 MajorForward
work_unit_contextintostrong_branch_helper(...)call.At Line 1184,
strong_branch_helper(...)is invoked without the newwork_unit_contextparameter, so Line 307 falls back tonullptrand deterministic work accounting is dropped in this DS path.Suggested fix
strong_branch_helper(start, end, start_time, original_lp, settings, var_types, fractional, root_solution.x, root_vstatus, edge_norms, root_obj, upper_bound, simplex_iteration_limit, pc, dual_simplex_obj_down, dual_simplex_obj_up, dual_simplex_status_down, dual_simplex_status_up, - sb_view); + sb_view, + work_unit_context);As per coding guidelines: "Validate algorithm correctness in optimization logic: simplex pivots, branch-and-bound decisions, routing heuristics, and constraint/objective handling must produce correct results."
🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/src/branch_and_bound/pseudo_costs.cpp` around lines 1184 - 1202, The call to strong_branch_helper(...) at the shown site omits the new work_unit_context parameter so the helper receives nullptr (losing deterministic work accounting); update the invocation to pass the existing work_unit_context variable into strong_branch_helper, and ensure the helper's signature and all other call sites accept the extra parameter (verify function strong_branch_helper and its declaration are updated accordingly) so deterministic work accounting is preserved on this DS path.
♻️ Duplicate comments (7)
cpp/src/branch_and_bound/pseudo_costs.cpp (1)
2086-2094:⚠️ Potential issue | 🟡 MinorGuard pseudo-cost updates against near-zero fractional distances.
Line 2088 and Line 2093 divide by
fracwithout a minimum threshold; near-integer values can amplify noise or spike pseudo-costs.Suggested guard
for (i_t k = 0; k < fractional.size(); k++) { const i_t j = fractional[k]; + constexpr f_t frac_eps = 1e-8; if (!std::isnan(strong_branch_down[k])) { const f_t frac = root_soln[j] - std::floor(root_soln[j]); - pseudo_cost_sum_down[j] += strong_branch_down[k] / frac; - pseudo_cost_num_down[j]++; + if (frac > frac_eps) { + pseudo_cost_sum_down[j] += strong_branch_down[k] / frac; + pseudo_cost_num_down[j]++; + } } if (!std::isnan(strong_branch_up[k])) { const f_t frac = std::ceil(root_soln[j]) - root_soln[j]; - pseudo_cost_sum_up[j] += strong_branch_up[k] / frac; - pseudo_cost_num_up[j]++; + if (frac > frac_eps) { + pseudo_cost_sum_up[j] += strong_branch_up[k] / frac; + pseudo_cost_num_up[j]++; + } } }As per coding guidelines: "Check numerical stability: prevent overflow/underflow, precision loss, division by zero/near-zero, and use epsilon comparisons for floating-point equality checks."
🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/src/branch_and_bound/pseudo_costs.cpp` around lines 2086 - 2094, The pseudo-cost updates in the block that updates pseudo_cost_sum_down[j]/pseudo_cost_sum_up[j] (using strong_branch_down[k], strong_branch_up[k] and root_soln[j]) divide by frac computed as root_soln[j]-floor(root_soln[j]) and ceil(root_soln[j])-root_soln[j] without guarding tiny denominators; change the logic in that section to protect against near-zero frac by either clamping frac = std::max(frac, eps) or skipping the update when frac < eps (choose a small constant eps, e.g. 1e-12 or a configurable constant), so that pseudo_cost_num_down/up are only incremented when a safe division occurred and no large spikes are introduced.cpp/src/mip_heuristics/local_search/rounding/lb_constraint_prop.cuh (1)
103-103:⚠️ Potential issue | 🟡 MinorLatent issue:
termination_checker_tmember lacks initializer.The
termination_checker_t max_timer;declaration has no initializer, buttermination_checker_thas no default constructor (per context snippet 1 showing only explicit constructors). This would cause a compilation error.However, based on learnings, this
lb_constraint_proppath is not currently compiled, so this is a latent issue that will only surface if the path is re-enabled. When that happens, the member will need explicit initialization similar toconstraint_prop_t:termination_checker_t max_timer{0.0, cuopt::termination_checker_t::root_tag_t{}};🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/src/mip_heuristics/local_search/rounding/lb_constraint_prop.cuh` at line 103, The member declaration termination_checker_t max_timer in lb_constraint_prop lacks an initializer and will fail to compile because termination_checker_t has no default constructor; initialize max_timer explicitly (same pattern used for constraint_prop_t) by constructing it with a zero timeout and the root_tag_t sentinel (use termination_checker_t{0.0, cuopt::termination_checker_t::root_tag_t{}} or equivalent) inside the lb_constraint_prop class so the member is properly initialized when that path is enabled.cpp/src/mip_heuristics/local_search/local_search.cu (1)
233-249:⚠️ Potential issue | 🟠 MajorAccount the helper CPUFJ threads against one deterministic budget.
In deterministic mode, each
ls_cpu_fjworker gets an independentwork_budget = time_limit(line 247). This means with multiple helpers, the aggregate work consumed can exceed the caller's budget, and timestamps won't align with the single work-total that B&B synchronizes against.Unlike
start_cpufj_deterministic()(lines 148-190) which registers withproducer_sync, these helpers operate with isolated budgets.🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/src/mip_heuristics/local_search/local_search.cu` around lines 233 - 249, The helper CPUFJ threads in ls_cpu_fj are being given independent deterministic budgets (setting cpu_fj.fj_cpu->work_units_elapsed = 0.0 and cpu_fj.fj_cpu->work_budget = time_limit), which allows aggregate work to exceed the single deterministic budget; instead, mirror start_cpufj_deterministic() and register each helper with the shared producer_sync so they are charged against the common deterministic budget (or otherwise attach the same shared budget handle) after create_cpu_climber(...) rather than assigning an isolated work_budget, ensuring helpers use the synchronized producer_sync accounting.cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cu (2)
1037-1041:⚠️ Potential issue | 🟠 MajorGraph batch can overshoot deterministic budget.
When
deterministic_work_estimateis true,use_graphremainstrue, sorun_step_device()executes a full 50-iteration batch. Thedeterministic_batch_workis only charged to the timer after the batch completes (lines 1107-1112). This allows a tiny positive budget to overshoot by an entire batch, shifting work timestamps used for deterministic publication/sync.Consider disabling graph mode when the remaining deterministic budget is less than a full batch's estimated work.
🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cu` around lines 1037 - 1041, The current logic keeps use_graph=true even when deterministic_work_estimate is set, allowing run_step_device() to execute a full iterations_per_graph batch and overshoot the deterministic budget; modify the decision so that when deterministic_work_estimate is true and the remaining deterministic budget (compare against deterministic_batch_work) is less than the estimated work for a full graph batch, disable graph mode or force iterations_per_batch=1 (i.e., set use_graph=false or set iterations_per_batch = 1) before calling run_step_device(), referencing the variables deterministic_work_estimate, use_graph, iterations_per_batch, run_step_device, and deterministic_batch_work so the small remaining budget is respected and work is charged before any multi-iteration graph run.
995-1017:⚠️ Potential issue | 🟠 MajorAdd CUDA error checking after the kernel launch.
The
init_lhs_and_violated_constraintskernel launch at line 996 lacks immediate error checking. Any failure will only surface later from the subsequentthrust::transform_reducecalls, making the fault harder to localize.💡 Suggested fix
data.violated_constraints.clear(stream); init_lhs_and_violated_constraints<i_t, f_t><<<4096, 256, 0, stream>>>(v); + RAFT_CHECK_CUDA(stream); // both transformreduce could be fused; but oh well hardly a bottleneckAs per coding guidelines, "Every CUDA kernel launch and memory operation must have error checking with CUDA_CHECK or equivalent verification."
🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cu` around lines 995 - 1017, The kernel launch init_lhs_and_violated_constraints<i_t, f_t><<<4096, 256, 0, stream>>>(v) lacks immediate CUDA error checking; add a post-launch check using the project's CUDA_CHECK (or equivalent) macro/function right after that launch (and before any thrust::transform_reduce calls) to capture launch failures on stream, e.g., call CUDA_CHECK(cudaGetLastError()) and/or CUDA_CHECK(cudaPeekAtLastError()/cudaStreamSynchronize(stream)) as appropriate so failures are reported at the kernel boundary; update surrounding code in feasibility_jump.cu near init_lhs_and_violated_constraints, data.violated_constraints.clear, and before the transform_reduce invocations to perform this check.cpp/tests/mip/diversity_test.cu (2)
147-153:⚠️ Potential issue | 🟠 MajorUse
solver.context.gpu_heur_loopas the deterministic work context.This helper creates an ad-hoc
work_limit_context_tinstead of using the shared context fromsolver.context.gpu_heur_loop. The nested timers inside local search/FJ readsolver.context.gpu_heur_loop, not this ad-hoc context, so they can bypass the work-accounting path the test is supposed to validate.💡 Suggested fix
detail::diversity_manager_t<int, double> diversity_manager(solver.context); solver.context.diversity_manager_ptr = &diversity_manager; - work_limit_context_t work_limit_context("DiversityManager"); - work_limit_context.deterministic = true; - diversity_manager.timer = termination_checker_t(work_limit_context, 60000, timer); + solver.context.gpu_heur_loop.deterministic = true; + diversity_manager.timer = + termination_checker_t(solver.context.gpu_heur_loop, 60000, timer);🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/tests/mip/diversity_test.cu` around lines 147 - 153, The test creates an ad-hoc work_limit_context_t and assigns diversity_manager.timer from it, which bypasses the shared GPU heuristic work accounting; replace that ad-hoc context with the shared solver.context.gpu_heur_loop: remove the local work_limit_context_t creation and set diversity_manager.timer = termination_checker_t(solver.context.gpu_heur_loop, 60000, timer) so nested timers inside local search/FJ use solver.context.gpu_heur_loop; keep the rest of the setup (diversity_manager, diversity_manager.diversity_config.initial_solution_only, solver.context.diversity_manager_ptr and diversity_manager.run_solver()) unchanged.
196-202:⚠️ Potential issue | 🟠 MajorUse
solver.context.gpu_heur_loopas the deterministic work context.Same issue as
test_initial_solution_determinism- the ad-hocwork_limit_context_tis not connected to the nested timers in FJ/local search.💡 Suggested fix
detail::diversity_manager_t<int, double> diversity_manager(solver.context); solver.context.diversity_manager_ptr = &diversity_manager; - work_limit_context_t work_limit_context("DiversityManager"); - work_limit_context.deterministic = true; - diversity_manager.timer = termination_checker_t(work_limit_context, 60000, timer); + solver.context.gpu_heur_loop.deterministic = true; + diversity_manager.timer = + termination_checker_t(solver.context.gpu_heur_loop, 60000, timer);🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/tests/mip/diversity_test.cu` around lines 196 - 202, The test creates an ad-hoc work_limit_context_t and uses it to build the termination_checker_t for diversity_manager (detail::diversity_manager_t<int,double> diversity_manager(...)), but this isn’t hooked into nested timers; replace the local work_limit_context_t with the existing solver.context.gpu_heur_loop: set solver.context.gpu_heur_loop.deterministic = true (or ensure it is true), construct diversity_manager.timer = termination_checker_t(solver.context.gpu_heur_loop, 60000, timer), remove the standalone work_limit_context_t, and keep diversity_manager.diversity_config.dry_run = true before calling diversity_manager.run_solver() so the nested FJ/local search timers use the correct deterministic context.
🧹 Nitpick comments (5)
cpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cu (2)
74-90: Redundant condition and dead code in determinism mode block.The check
!std::isinf(work_limit)on line 79 is redundant becausedeterminism_modeon line 74 is only true whenstd::isfinite(work_limit)(line 68), which already excludes infinity. Consequently, theelsebranch (lines 88-90) settingestim_iters = std::numeric_limits<int>::max()is dead code.♻️ Proposed simplification
if (determinism_mode) { // try to estimate the iteration count based on the requested work limit // TODO: replace with an actual model. this is a rather ugly hack to avoid having // to touch the PDLP code for this initial PR estim_iters = 100; - if (!std::isinf(work_limit)) { - do { - // TODO: use an actual predictor model here - double estim_ms = 313 + 200 * op_problem.n_variables - 400 * op_problem.n_constraints + - 600 * op_problem.coefficients.size() + 7100 * estim_iters; - estim_ms = std::max(0.0, estim_ms); - if (estim_ms > work_limit * 1000) { break; } - estim_iters += 100; - } while (true); - } else { - estim_iters = std::numeric_limits<int>::max(); - } + do { + // TODO: use an actual predictor model here + double estim_ms = 313 + 200 * op_problem.n_variables - 400 * op_problem.n_constraints + + 600 * op_problem.coefficients.size() + 7100 * estim_iters; + estim_ms = std::max(0.0, estim_ms); + if (estim_ms > work_limit * 1000) { break; } + estim_iters += 100; + } while (true);🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cu` around lines 74 - 90, The determinism_mode block contains a redundant check for !std::isinf(work_limit) and a dead else branch; remove the inner if/else and always run the estimators loop because determinism_mode is only set when std::isfinite(work_limit) earlier, so delete the else that sets estim_iters = std::numeric_limits<int>::max() and collapse the loop to run unconditionally inside the determinism_mode branch, keeping the estim_iters initialization and the loop that computes estim_ms using op_problem.n_variables, op_problem.n_constraints, op_problem.coefficients.size(), and updating estim_iters.
23-23: Unused include:<atomic>is not used in this file.The
#include <atomic>header was added but nostd::atomictypes are used. The static counter on line 53 uses plainuint64_t. Consider removing the unused include or convertinglp_call_countertostd::atomic<uint64_t>if thread-safety is intended.🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cu` at line 23, The include <atomic> is unused or lp_call_counter should be atomic for thread-safety; either remove the unused `#include` <atomic> or change the static counter lp_call_counter (currently a uint64_t) to std::atomic<uint64_t> and update its increment/usesites accordingly (use fetch_add or ++ semantics) so thread-safe increments are correct; ensure to include <atomic> only when converting lp_call_counter to std::atomic<uint64_t>.cpp/tests/mip/local_search_test.cu (2)
196-197: Unusedspin_stream_raii_tvariables.These RAII objects are declared but appear unused in the test. If they're intended to keep CUDA streams alive during the test, consider adding a comment explaining their purpose. Otherwise, they can be removed.
🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/tests/mip/local_search_test.cu` around lines 196 - 197, The two RAII stream objects spin_stream_raii_t spin_stream_1 and spin_stream_raii_t spin_stream_2 are declared but never used; either remove these unused declarations or explicitly document their purpose if they are intended to keep CUDA streams alive (add a comment above spin_stream_1/spin_stream_2 stating they intentionally hold/flush streams for the test), and if they must be used, ensure they're constructed with the appropriate stream handles where the test creates CUDA streams so the RAII actually manages those resources.
231-236: Consider enablingFJ_ON_ZEROmode in the test suite.The
FJ_ON_ZEROmode is defined in the enum (line 74) and has a corresponding execution path (lines 162-164), but it's commented out in the test instantiation. If this mode is intended to be deterministic, it should be tested.🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/tests/mip/local_search_test.cu` around lines 231 - 236, Enable the commented-out FJ_ON_ZERO test case in the test instantiation: update the INSTANTIATE_TEST_SUITE_P call for LocalSearchTests/LocalSearchTestParams to include std::make_tuple(local_search_mode_t::FJ_ON_ZERO) alongside FP, FJ_LINE_SEGMENT and FJ_ANNEALING so the FJ_ON_ZERO execution path is covered by the test suite; ensure you remove the comment markers around the std::make_tuple(local_search_mode_t::FJ_ON_ZERO) entry so the test framework will instantiate it.cpp/tests/mip/diversity_test.cu (1)
240-241: Structured binding discards the third return value.The
recombinefunction returnsstd::tuple<solution, bool, double>but only two elements are captured. While this works in C++, explicitly ignoring the third element makes intent clearer.💡 Suggested fix
- auto [offspring, success] = + auto [offspring, success, /* work_done */] = diversity_manager.recombine(pop_vector[i], pop_vector[j], recombiner);🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/tests/mip/diversity_test.cu` around lines 240 - 241, The structured binding is discarding the third value from diversity_manager.recombine which returns std::tuple<solution, bool, double>; update the capture to explicitly bind the third element (for example capture as a named variable like "metric" or use std::ignore if you intend to discard it) so the intent is clear when calling recombine and avoid silent discarding of the double return; update the line that currently binds "offspring" and "success" to also bind the third value from recombine.
🤖 Prompt for all review comments with AI agents
Verify each finding against the current code and only fix it if needed.
Inline comments:
In `@cpp/src/mip_heuristics/diversity/recombiners/fp_recombiner.cuh`:
- Around line 68-69: The current TODO placeholder sets work to a small static
value and then returns 0.0 on LP-based early exits, losing the work already
spent by get_relaxed_lp_solution; update the code in fp_recombiner.cuh so the
work variable (currently computed from n_vars_from_other) accumulates the LP
pass cost and is returned on all early-exit paths. Concretely: in the scope
around get_relaxed_lp_solution(...) ensure you add the LP work contribution to
the existing work variable (reference n_vars_from_other and work) and return
that accumulated work instead of 0.0; apply the same change to the other
early-exit sites mentioned (around the lines corresponding to 100-105) so
recombine_stats.set_recombiner_work(work) receives the correct nonzero value.
In `@cpp/src/mip_heuristics/problem/presolve_data.cu`:
- Around line 236-241: The code synchronizes the wrong handle's stream: replace
the call to problem.handle_ptr->sync_stream() with a sync on the actual handle
used for subsequent GPU work (the local pointer h), e.g. call h->sync_stream()
or otherwise synchronize h->get_stream() before invoking cuopt::host_copy and
other GPU operations; update any related comments to reflect that
synchronization targets the chosen handle (h) to avoid mismatched streams when
handle_override is provided.
In `@cpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cu`:
- Around line 53-54: The static non-atomic lp_call_counter causes a data race
when multiple threads call get_relaxed_lp_solution; change lp_call_counter to an
atomic (e.g., std::atomic<uint64_t>) and obtain lp_call_id using an atomic
fetch-add (or equivalent) so increments are thread-safe, updating references to
lp_call_counter and lp_call_id in get_relaxed_lp_solution accordingly.
In `@cpp/src/mip_heuristics/solver.cu`:
- Around line 491-496: The async lambda that calls
RAFT_CUDA_TRY(cudaSetDevice(bb_device_id)) must verify the result and fail fast
on error: replace the bare cudaSetDevice call with the project’s CUDA
error-checking helper (e.g., RAFT_CUDA_TRY or CUDA_CHECK) or explicitly check
cudaError_t and handle failures before calling branch_and_bound->solve; update
the lambda that captures bb_device_id and calls
branch_and_bound->solve(branch_and_bound_solution) so it logs/propagates the
cudaSetDevice error (using the same error handling pattern used elsewhere) and
returns/throws an error instead of proceeding to solve when device
initialization fails.
In `@cpp/tests/mip/diversity_test.cu`:
- Around line 267-279: The block that collects hashes and computes final_hash
(using diversity_manager.get_population_pointer(), pop->population_to_vector(),
detail::compute_hash(hashes), and the printf that references path and
final_hash) is dead/unreachable because an earlier return happens; remove this
unreachable code or move its logic before the earlier return so only reachable
code remains—specifically delete the loop that pushes sol.get_hash(), the
compute_hash call, the printf, and the trailing return final_hash; ensure no
references to final_hash or hashes remain after the earlier return.
---
Outside diff comments:
In `@cpp/src/branch_and_bound/pseudo_costs.cpp`:
- Around line 1184-1202: The call to strong_branch_helper(...) at the shown site
omits the new work_unit_context parameter so the helper receives nullptr (losing
deterministic work accounting); update the invocation to pass the existing
work_unit_context variable into strong_branch_helper, and ensure the helper's
signature and all other call sites accept the extra parameter (verify function
strong_branch_helper and its declaration are updated accordingly) so
deterministic work accounting is preserved on this DS path.
---
Duplicate comments:
In `@cpp/src/branch_and_bound/pseudo_costs.cpp`:
- Around line 2086-2094: The pseudo-cost updates in the block that updates
pseudo_cost_sum_down[j]/pseudo_cost_sum_up[j] (using strong_branch_down[k],
strong_branch_up[k] and root_soln[j]) divide by frac computed as
root_soln[j]-floor(root_soln[j]) and ceil(root_soln[j])-root_soln[j] without
guarding tiny denominators; change the logic in that section to protect against
near-zero frac by either clamping frac = std::max(frac, eps) or skipping the
update when frac < eps (choose a small constant eps, e.g. 1e-12 or a
configurable constant), so that pseudo_cost_num_down/up are only incremented
when a safe division occurred and no large spikes are introduced.
In `@cpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cu`:
- Around line 1037-1041: The current logic keeps use_graph=true even when
deterministic_work_estimate is set, allowing run_step_device() to execute a full
iterations_per_graph batch and overshoot the deterministic budget; modify the
decision so that when deterministic_work_estimate is true and the remaining
deterministic budget (compare against deterministic_batch_work) is less than the
estimated work for a full graph batch, disable graph mode or force
iterations_per_batch=1 (i.e., set use_graph=false or set iterations_per_batch =
1) before calling run_step_device(), referencing the variables
deterministic_work_estimate, use_graph, iterations_per_batch, run_step_device,
and deterministic_batch_work so the small remaining budget is respected and work
is charged before any multi-iteration graph run.
- Around line 995-1017: The kernel launch init_lhs_and_violated_constraints<i_t,
f_t><<<4096, 256, 0, stream>>>(v) lacks immediate CUDA error checking; add a
post-launch check using the project's CUDA_CHECK (or equivalent) macro/function
right after that launch (and before any thrust::transform_reduce calls) to
capture launch failures on stream, e.g., call CUDA_CHECK(cudaGetLastError())
and/or CUDA_CHECK(cudaPeekAtLastError()/cudaStreamSynchronize(stream)) as
appropriate so failures are reported at the kernel boundary; update surrounding
code in feasibility_jump.cu near init_lhs_and_violated_constraints,
data.violated_constraints.clear, and before the transform_reduce invocations to
perform this check.
In `@cpp/src/mip_heuristics/local_search/local_search.cu`:
- Around line 233-249: The helper CPUFJ threads in ls_cpu_fj are being given
independent deterministic budgets (setting cpu_fj.fj_cpu->work_units_elapsed =
0.0 and cpu_fj.fj_cpu->work_budget = time_limit), which allows aggregate work to
exceed the single deterministic budget; instead, mirror
start_cpufj_deterministic() and register each helper with the shared
producer_sync so they are charged against the common deterministic budget (or
otherwise attach the same shared budget handle) after create_cpu_climber(...)
rather than assigning an isolated work_budget, ensuring helpers use the
synchronized producer_sync accounting.
In `@cpp/src/mip_heuristics/local_search/rounding/lb_constraint_prop.cuh`:
- Line 103: The member declaration termination_checker_t max_timer in
lb_constraint_prop lacks an initializer and will fail to compile because
termination_checker_t has no default constructor; initialize max_timer
explicitly (same pattern used for constraint_prop_t) by constructing it with a
zero timeout and the root_tag_t sentinel (use termination_checker_t{0.0,
cuopt::termination_checker_t::root_tag_t{}} or equivalent) inside the
lb_constraint_prop class so the member is properly initialized when that path is
enabled.
In `@cpp/tests/mip/diversity_test.cu`:
- Around line 147-153: The test creates an ad-hoc work_limit_context_t and
assigns diversity_manager.timer from it, which bypasses the shared GPU heuristic
work accounting; replace that ad-hoc context with the shared
solver.context.gpu_heur_loop: remove the local work_limit_context_t creation and
set diversity_manager.timer =
termination_checker_t(solver.context.gpu_heur_loop, 60000, timer) so nested
timers inside local search/FJ use solver.context.gpu_heur_loop; keep the rest of
the setup (diversity_manager,
diversity_manager.diversity_config.initial_solution_only,
solver.context.diversity_manager_ptr and diversity_manager.run_solver())
unchanged.
- Around line 196-202: The test creates an ad-hoc work_limit_context_t and uses
it to build the termination_checker_t for diversity_manager
(detail::diversity_manager_t<int,double> diversity_manager(...)), but this isn’t
hooked into nested timers; replace the local work_limit_context_t with the
existing solver.context.gpu_heur_loop: set
solver.context.gpu_heur_loop.deterministic = true (or ensure it is true),
construct diversity_manager.timer =
termination_checker_t(solver.context.gpu_heur_loop, 60000, timer), remove the
standalone work_limit_context_t, and keep
diversity_manager.diversity_config.dry_run = true before calling
diversity_manager.run_solver() so the nested FJ/local search timers use the
correct deterministic context.
---
Nitpick comments:
In `@cpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cu`:
- Around line 74-90: The determinism_mode block contains a redundant check for
!std::isinf(work_limit) and a dead else branch; remove the inner if/else and
always run the estimators loop because determinism_mode is only set when
std::isfinite(work_limit) earlier, so delete the else that sets estim_iters =
std::numeric_limits<int>::max() and collapse the loop to run unconditionally
inside the determinism_mode branch, keeping the estim_iters initialization and
the loop that computes estim_ms using op_problem.n_variables,
op_problem.n_constraints, op_problem.coefficients.size(), and updating
estim_iters.
- Line 23: The include <atomic> is unused or lp_call_counter should be atomic
for thread-safety; either remove the unused `#include` <atomic> or change the
static counter lp_call_counter (currently a uint64_t) to std::atomic<uint64_t>
and update its increment/usesites accordingly (use fetch_add or ++ semantics) so
thread-safe increments are correct; ensure to include <atomic> only when
converting lp_call_counter to std::atomic<uint64_t>.
In `@cpp/tests/mip/diversity_test.cu`:
- Around line 240-241: The structured binding is discarding the third value from
diversity_manager.recombine which returns std::tuple<solution, bool, double>;
update the capture to explicitly bind the third element (for example capture as
a named variable like "metric" or use std::ignore if you intend to discard it)
so the intent is clear when calling recombine and avoid silent discarding of the
double return; update the line that currently binds "offspring" and "success" to
also bind the third value from recombine.
In `@cpp/tests/mip/local_search_test.cu`:
- Around line 196-197: The two RAII stream objects spin_stream_raii_t
spin_stream_1 and spin_stream_raii_t spin_stream_2 are declared but never used;
either remove these unused declarations or explicitly document their purpose if
they are intended to keep CUDA streams alive (add a comment above
spin_stream_1/spin_stream_2 stating they intentionally hold/flush streams for
the test), and if they must be used, ensure they're constructed with the
appropriate stream handles where the test creates CUDA streams so the RAII
actually manages those resources.
- Around line 231-236: Enable the commented-out FJ_ON_ZERO test case in the test
instantiation: update the INSTANTIATE_TEST_SUITE_P call for
LocalSearchTests/LocalSearchTestParams to include
std::make_tuple(local_search_mode_t::FJ_ON_ZERO) alongside FP, FJ_LINE_SEGMENT
and FJ_ANNEALING so the FJ_ON_ZERO execution path is covered by the test suite;
ensure you remove the comment markers around the
std::make_tuple(local_search_mode_t::FJ_ON_ZERO) entry so the test framework
will instantiate it.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Pro
Run ID: 85b5b753-f0c9-41b8-be6a-0d56784ce4d5
📒 Files selected for processing (55)
cpp/CMakeLists.txtcpp/include/cuopt/linear_programming/constants.hcpp/include/cuopt/linear_programming/utilities/internals.hppcpp/src/branch_and_bound/pseudo_costs.cppcpp/src/mip_heuristics/diversity/diversity_manager.cucpp/src/mip_heuristics/diversity/diversity_manager.cuhcpp/src/mip_heuristics/diversity/population.cucpp/src/mip_heuristics/diversity/population.cuhcpp/src/mip_heuristics/diversity/recombiners/bound_prop_recombiner.cuhcpp/src/mip_heuristics/diversity/recombiners/fp_recombiner.cuhcpp/src/mip_heuristics/diversity/recombiners/line_segment_recombiner.cuhcpp/src/mip_heuristics/early_heuristic.cuhcpp/src/mip_heuristics/feasibility_jump/feasibility_jump.cucpp/src/mip_heuristics/local_search/feasibility_pump/feasibility_pump.cucpp/src/mip_heuristics/local_search/feasibility_pump/feasibility_pump.cuhcpp/src/mip_heuristics/local_search/line_segment_search/line_segment_search.cucpp/src/mip_heuristics/local_search/line_segment_search/line_segment_search.cuhcpp/src/mip_heuristics/local_search/local_search.cucpp/src/mip_heuristics/local_search/local_search.cuhcpp/src/mip_heuristics/local_search/rounding/bounds_repair.cucpp/src/mip_heuristics/local_search/rounding/bounds_repair.cuhcpp/src/mip_heuristics/local_search/rounding/constraint_prop.cucpp/src/mip_heuristics/local_search/rounding/constraint_prop.cuhcpp/src/mip_heuristics/local_search/rounding/lb_bounds_repair.cucpp/src/mip_heuristics/local_search/rounding/lb_bounds_repair.cuhcpp/src/mip_heuristics/local_search/rounding/lb_constraint_prop.cucpp/src/mip_heuristics/local_search/rounding/lb_constraint_prop.cuhcpp/src/mip_heuristics/presolve/bounds_presolve.cucpp/src/mip_heuristics/presolve/bounds_presolve.cuhcpp/src/mip_heuristics/presolve/bounds_update_data.cucpp/src/mip_heuristics/presolve/lb_probing_cache.cucpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve.cucpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve.cuhcpp/src/mip_heuristics/presolve/multi_probe.cucpp/src/mip_heuristics/presolve/multi_probe.cuhcpp/src/mip_heuristics/presolve/probing_cache.cucpp/src/mip_heuristics/presolve/probing_cache.cuhcpp/src/mip_heuristics/problem/presolve_data.cucpp/src/mip_heuristics/problem/presolve_data.cuhcpp/src/mip_heuristics/problem/problem.cucpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cucpp/src/mip_heuristics/solve.cucpp/src/mip_heuristics/solver.cucpp/src/mip_heuristics/solver_context.cuhcpp/src/utilities/copy_helpers.hppcpp/src/utilities/cuda_helpers.cuhcpp/src/utilities/termination_checker.hppcpp/tests/CMakeLists.txtcpp/tests/mip/CMakeLists.txtcpp/tests/mip/determinism_test.cucpp/tests/mip/diversity_test.cucpp/tests/mip/incumbent_callback_test.cucpp/tests/mip/local_search_test.cucpp/tests/mip/presolve_test.cuskills/cuopt-developer/SKILL.md
✅ Files skipped from review due to trivial changes (3)
- skills/cuopt-developer/SKILL.md
- cpp/tests/CMakeLists.txt
- cpp/src/mip_heuristics/diversity/diversity_manager.cu
🚧 Files skipped from review as they are similar to previous changes (19)
- cpp/src/mip_heuristics/presolve/bounds_update_data.cu
- cpp/src/mip_heuristics/presolve/multi_probe.cuh
- cpp/src/mip_heuristics/presolve/load_balanced_bounds_presolve.cu
- cpp/src/mip_heuristics/local_search/feasibility_pump/feasibility_pump.cuh
- cpp/src/mip_heuristics/local_search/rounding/constraint_prop.cuh
- cpp/src/mip_heuristics/presolve/lb_probing_cache.cu
- cpp/src/mip_heuristics/local_search/rounding/bounds_repair.cuh
- cpp/src/mip_heuristics/presolve/bounds_presolve.cu
- cpp/src/mip_heuristics/local_search/rounding/lb_bounds_repair.cu
- cpp/src/mip_heuristics/problem/presolve_data.cuh
- cpp/src/mip_heuristics/local_search/rounding/lb_constraint_prop.cu
- cpp/src/utilities/cuda_helpers.cuh
- cpp/src/mip_heuristics/solver_context.cuh
- cpp/src/mip_heuristics/local_search/line_segment_search/line_segment_search.cu
- cpp/src/mip_heuristics/problem/problem.cu
- cpp/src/mip_heuristics/presolve/probing_cache.cu
- cpp/tests/mip/determinism_test.cu
- cpp/src/mip_heuristics/local_search/rounding/bounds_repair.cu
- cpp/tests/mip/presolve_test.cu
| // TODO: CHANGE | ||
| double work = static_cast<double>(n_vars_from_other) / 1e8; |
There was a problem hiding this comment.
Don't drop recombiner work on LP-based early exits.
After get_relaxed_lp_solution(...) runs, these branches return 0.0 even though the infeasibility-detection path already spent work. Since cpp/src/mip_heuristics/diversity/diversity_manager.cu:1146-1150 forwards this value into recombine_stats.set_recombiner_work(work), the current code undercounts exactly the expensive failure cases. At minimum, return the accumulated work here; ideally replace the // TODO: CHANGE placeholder with a cost that also covers the LP pass.
Suggested minimal fix
- // TODO: CHANGE
double work = static_cast<double>(n_vars_from_other) / 1e8;
@@
if (lp_response.get_termination_status() == pdlp_termination_status_t::PrimalInfeasible ||
lp_response.get_termination_status() == pdlp_termination_status_t::DualInfeasible ||
lp_response.get_termination_status() == pdlp_termination_status_t::TimeLimit) {
CUOPT_LOG_DEBUG("FP recombiner failed because LP found infeasible!");
- return std::make_tuple(offspring, false, 0.0);
+ return std::make_tuple(offspring, false, work);
}Also applies to: 100-105
🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed.
In `@cpp/src/mip_heuristics/diversity/recombiners/fp_recombiner.cuh` around lines
68 - 69, The current TODO placeholder sets work to a small static value and then
returns 0.0 on LP-based early exits, losing the work already spent by
get_relaxed_lp_solution; update the code in fp_recombiner.cuh so the work
variable (currently computed from n_vars_from_other) accumulates the LP pass
cost and is returned on all early-exit paths. Concretely: in the scope around
get_relaxed_lp_solution(...) ensure you add the LP work contribution to the
existing work variable (reference n_vars_from_other and work) and return that
accumulated work instead of 0.0; apply the same change to the other early-exit
sites mentioned (around the lines corresponding to 100-105) so
recombine_stats.set_recombiner_work(work) receives the correct nonzero value.
| const auto* h = handle_override ? handle_override : problem.handle_ptr; | ||
| cuopt_assert(assignment.size() == papilo_reduced_to_original_map.size(), | ||
| "Papilo uncrush assignment size mismatch"); | ||
| auto h_assignment = cuopt::host_copy(assignment, problem.handle_ptr->get_stream()); | ||
| problem.handle_ptr->sync_stream(); | ||
| auto h_assignment = cuopt::host_copy(assignment, h->get_stream()); | ||
| std::vector<f_t> full_assignment; |
There was a problem hiding this comment.
🧩 Analysis chain
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# 1) Inspect all papilo_uncrush_assignment call sites and passed handle argument.
rg -nP -C3 --type=cu --type=cpp --type=cuh '\bpapilo_uncrush_assignment\s*\('
# 2) Inspect the exact sync/copy sequence in the implementation under review.
rg -nP -C8 'handle_override|sync_stream|host_copy\(assignment|raft::copy\(assignment\.data' cpp/src/mip_heuristics/problem/presolve_data.cu
# 3) Inspect host_copy helper implementation to confirm synchronization behavior.
rg -nP -C8 --type=hpp --type=h --type=cu --type=cpp '\bhost_copy\s*\('Repository: NVIDIA/cuopt
Length of output: 82
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Search for papilo_uncrush_assignment call sites
echo "=== papilo_uncrush_assignment call sites ==="
rg -nP -C3 'papilo_uncrush_assignment\s*\(' --glob='*.cu' --glob='*.cpp' --glob='*.cuh'
echo ""
echo "=== Sync/copy sequence in presolve_data.cu ==="
rg -nP -C8 'handle_override|sync_stream|host_copy\(assignment|raft::copy\(assignment\.data' cpp/src/mip_heuristics/problem/presolve_data.cu
echo ""
echo "=== host_copy helper signature ==="
rg -nP -C8 'host_copy\s*\(' --glob='*.hpp' --glob='*.h' --glob='*.cu' --glob='*.cpp'Repository: NVIDIA/cuopt
Length of output: 50369
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Read the exact lines from the file to understand context
echo "=== Content around lines 230-250 of presolve_data.cu ==="
sed -n '230,250p' cpp/src/mip_heuristics/problem/presolve_data.cu | cat -n
echo ""
echo "=== Full function context (broader view) ==="
# Find the function containing line 236 and show more context
awk 'NR==200,NR==260' cpp/src/mip_heuristics/problem/presolve_data.cu | cat -nRepository: NVIDIA/cuopt
Length of output: 3700
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Find where handle_override is defined and used in this function
echo "=== Function containing the code at line 236 ==="
sed -n '180,280p' cpp/src/mip_heuristics/problem/presolve_data.cu | cat -nRepository: NVIDIA/cuopt
Length of output: 3522
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Find host_copy implementation
echo "=== Searching for host_copy definition/implementation ==="
find . -type f \( -name "*.hpp" -o -name "*.h" -o -name "*.cu" -o -name "*.cpp" \) -exec grep -l "host_copy" {} \; | head -20
# Look for the template definition
echo ""
echo "=== host_copy template definition ==="
rg -A 15 'template.*host_copy\s*\(' --glob='*.hpp' --glob='*.h' | head -80Repository: NVIDIA/cuopt
Length of output: 2270
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Check what host_copy is - look in header files for implementation
echo "=== Looking for host_copy in utilities ==="
find . -path ./build -prune -o -type f -name "*.hpp" -o -name "*.h" | xargs grep -l 'host_copy' | head -10
# Find it with context
echo ""
echo "=== host_copy usage pattern analysis ==="
rg 'host_copy\(' --glob='*.cu' -B2 -A2 | head -100Repository: NVIDIA/cuopt
Length of output: 7907
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Find and display host_copy implementation
echo "=== host_copy implementation in copy_helpers.hpp ==="
cat cpp/src/utilities/copy_helpers.hpp | head -200Repository: NVIDIA/cuopt
Length of output: 4676
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# See the rest of host_copy implementation
echo "=== Continuation of copy_helpers.hpp ==="
sed -n '160,220p' cpp/src/utilities/copy_helpers.hppRepository: NVIDIA/cuopt
Length of output: 1859
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Check where assignment data comes from in the call sites
echo "=== Call site 1: solution_callbacks.cuh:102 context ==="
sed -n '95,110p' cpp/src/mip_heuristics/solution_callbacks.cuh
echo ""
echo "=== Call site 2: solve.cu:113 context ==="
sed -n '108,120p' cpp/src/mip_heuristics/solve.cu
# Also check if temp_sol is using a different stream
echo ""
echo "=== solution_t structure and stream handling ==="
rg -A 5 'struct solution_t|class solution_t' --glob='*.cuh' --glob='*.hpp' cpp/src/mip_heuristics/solution/ | head -60Repository: NVIDIA/cuopt
Length of output: 2489
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Now let's understand the rmm::device_uvector and stream associations
echo "=== Check how assignment is created/initialized ==="
rg -B5 -A5 'rmm::device_uvector<f_t>& assignment' cpp/src/mip_heuristics/problem/presolve_data.cu | head -40
# Check what stream assignment is associated with
echo ""
echo "=== Looking for assignment initialization in solution_t ==="
rg -B3 -A3 'assignment.*resize|assignment.*rmm' cpp/src/mip_heuristics/solution/solution.cu | head -50
# Check if there are any operations on assignment before papilo_uncrush_assignment is called
echo ""
echo "=== In solution_callbacks.cuh, what happens to assignment before papilo_uncrush_assignment ==="
sed -n '85,105p' cpp/src/mip_heuristics/solution_callbacks.cuhRepository: NVIDIA/cuopt
Length of output: 4351
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Let's understand stream lifetime in solution_t copy constructor
echo "=== solution_t copy constructor ==="
rg -A 20 'solution_t<i_t, f_t>::solution_t\(const solution_t' cpp/src/mip_heuristics/solution/solution.cu
echo ""
echo "=== solution_t constructor from problem ==="
rg -A 15 'solution_t<i_t, f_t>::solution_t\(problem_t' cpp/src/mip_heuristics/solution/solution.cu | head -40
# Let me check if assignment retains its stream affinity across copies
echo ""
echo "=== rmm::device_uvector stream handling in copy ==="
rg 'assignment = other.assignment|assignment\(' cpp/src/mip_heuristics/solution/solution.cu | head -20Repository: NVIDIA/cuopt
Length of output: 3059
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Let me trace through the actual execution in solution_callbacks.cuh more carefully
echo "=== Check solution_t constructor with handle_override ==="
sed -n '90,110p' cpp/src/mip_heuristics/solution_callbacks.cuh
echo ""
echo "=== Now let's look at the actual fix - does syncing problem.handle_ptr help or hurt? ==="
# The key question: when handle_override differs from problem.handle_ptr, what stream is assignment on?
echo "When papilo_uncrush_assignment is called with handle_override:"
echo "- h = handle_override (if not nullptr)"
echo "- h = problem.handle_ptr (if nullptr)"
echo ""
echo "The assignment device_uvector - which handle does it use?"
echo "Looking at line 240-246 operations:"
rg -B2 -A10 'const auto\* h = handle_override' cpp/src/mip_heuristics/problem/presolve_data.cu
echo ""
echo "=== Check if there are similar mixed-handle sync patterns elsewhere ==="
rg 'problem\.handle_ptr->sync|h->sync|handle_override' cpp/src/mip_heuristics/problem/presolve_data.cuRepository: NVIDIA/cuopt
Length of output: 2706
🏁 Script executed:
#!/bin/bash
set -euo pipefail
# Search for similar patterns: sync on one handle, use another handle's stream
echo "=== Looking for mixed-handle sync/stream patterns ==="
rg -B3 'sync_stream\(\)' cpp/src/mip_heuristics/problem/ --glob='*.cu' | grep -A3 -B3 'handle_override\|h->get_stream\|h->sync'
echo ""
echo "=== Confirm the issue: in solution_callbacks.cuh, sol.handle_ptr vs problem_ptr->handle_ptr ==="
sed -n '88,105p' cpp/src/mip_heuristics/solution_callbacks.cuh | cat -nRepository: NVIDIA/cuopt
Length of output: 2727
Synchronize the correct stream before GPU operations
At line 240, the code syncs problem.handle_ptr->sync_stream(), but the subsequent GPU operations on lines 241 and 244 use h->get_stream(), which may differ when handle_override is provided (as in solution_callbacks.cuh:102). This mixed-handle synchronization is incorrect; the initial sync should target the stream actually used for GPU work.
Suggested fix
- problem.handle_ptr->sync_stream();
+ h->sync_stream();
auto h_assignment = cuopt::host_copy(assignment, h->get_stream());📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
| const auto* h = handle_override ? handle_override : problem.handle_ptr; | |
| cuopt_assert(assignment.size() == papilo_reduced_to_original_map.size(), | |
| "Papilo uncrush assignment size mismatch"); | |
| auto h_assignment = cuopt::host_copy(assignment, problem.handle_ptr->get_stream()); | |
| problem.handle_ptr->sync_stream(); | |
| auto h_assignment = cuopt::host_copy(assignment, h->get_stream()); | |
| std::vector<f_t> full_assignment; | |
| const auto* h = handle_override ? handle_override : problem.handle_ptr; | |
| cuopt_assert(assignment.size() == papilo_reduced_to_original_map.size(), | |
| "Papilo uncrush assignment size mismatch"); | |
| h->sync_stream(); | |
| auto h_assignment = cuopt::host_copy(assignment, h->get_stream()); | |
| std::vector<f_t> full_assignment; |
🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed.
In `@cpp/src/mip_heuristics/problem/presolve_data.cu` around lines 236 - 241, The
code synchronizes the wrong handle's stream: replace the call to
problem.handle_ptr->sync_stream() with a sync on the actual handle used for
subsequent GPU work (the local pointer h), e.g. call h->sync_stream() or
otherwise synchronize h->get_stream() before invoking cuopt::host_copy and other
GPU operations; update any related comments to reflect that synchronization
targets the chosen handle (h) to avoid mismatched streams when handle_override
is provided.
| int bb_device_id = context.handle_ptr->get_device(); | ||
| branch_and_bound_status_future = | ||
| std::async(std::launch::async, [&branch_and_bound, &branch_and_bound_solution, bb_device_id] { | ||
| RAFT_CUDA_TRY(cudaSetDevice(bb_device_id)); | ||
| return branch_and_bound->solve(branch_and_bound_solution); | ||
| }); |
There was a problem hiding this comment.
Add CUDA error checking after cudaSetDevice.
The cudaSetDevice call inside the async B&B thread should have error checking to catch device initialization failures early.
💡 Suggested fix
branch_and_bound_status_future =
std::async(std::launch::async, [&branch_and_bound, &branch_and_bound_solution, bb_device_id] {
- RAFT_CUDA_TRY(cudaSetDevice(bb_device_id));
+ cudaError_t err = cudaSetDevice(bb_device_id);
+ if (err != cudaSuccess) {
+ CUOPT_LOG_ERROR("Failed to set CUDA device %d: %s", bb_device_id, cudaGetErrorString(err));
+ return dual_simplex::mip_status_t::NUMERICAL_ISSUES;
+ }
return branch_and_bound->solve(branch_and_bound_solution);
});As per coding guidelines, "Every CUDA kernel launch and memory operation must have error checking with CUDA_CHECK or equivalent verification."
📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
| int bb_device_id = context.handle_ptr->get_device(); | |
| branch_and_bound_status_future = | |
| std::async(std::launch::async, [&branch_and_bound, &branch_and_bound_solution, bb_device_id] { | |
| RAFT_CUDA_TRY(cudaSetDevice(bb_device_id)); | |
| return branch_and_bound->solve(branch_and_bound_solution); | |
| }); | |
| int bb_device_id = context.handle_ptr->get_device(); | |
| branch_and_bound_status_future = | |
| std::async(std::launch::async, [&branch_and_bound, &branch_and_bound_solution, bb_device_id] { | |
| cudaError_t err = cudaSetDevice(bb_device_id); | |
| if (err != cudaSuccess) { | |
| CUOPT_LOG_ERROR("Failed to set CUDA device %d: %s", bb_device_id, cudaGetErrorString(err)); | |
| return dual_simplex::mip_status_t::NUMERICAL_ISSUES; | |
| } | |
| return branch_and_bound->solve(branch_and_bound_solution); | |
| }); |
🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed.
In `@cpp/src/mip_heuristics/solver.cu` around lines 491 - 496, The async lambda
that calls RAFT_CUDA_TRY(cudaSetDevice(bb_device_id)) must verify the result and
fail fast on error: replace the bare cudaSetDevice call with the project’s CUDA
error-checking helper (e.g., RAFT_CUDA_TRY or CUDA_CHECK) or explicitly check
cudaError_t and handle failures before calling branch_and_bound->solve; update
the lambda that captures bb_device_id and calls
branch_and_bound->solve(branch_and_bound_solution) so it logs/propagates the
cudaSetDevice error (using the same error handling pattern used elsewhere) and
returns/throws an error instead of proceeding to solve when device
initialization fails.
|
|
||
| auto pop = diversity_manager.get_population_pointer(); | ||
| for (const auto& sol : pop->population_to_vector()) { | ||
| hashes.push_back(sol.get_hash()); | ||
| } | ||
|
|
||
| uint32_t final_hash = detail::compute_hash(hashes); | ||
| printf("%s: final hash: 0x%x, pop size %d\n", | ||
| path.c_str(), | ||
| final_hash, | ||
| (int)pop->population_to_vector().size()); | ||
| return final_hash; | ||
| } |
There was a problem hiding this comment.
Remove unreachable dead code.
Lines 267-279 are unreachable because line 266 returns early. This appears to be leftover code from an earlier iteration.
💡 Suggested fix
return detail::compute_hash(hashes);
-
- auto pop = diversity_manager.get_population_pointer();
- for (const auto& sol : pop->population_to_vector()) {
- hashes.push_back(sol.get_hash());
- }
-
- uint32_t final_hash = detail::compute_hash(hashes);
- printf("%s: final hash: 0x%x, pop size %d\n",
- path.c_str(),
- final_hash,
- (int)pop->population_to_vector().size());
- return final_hash;
}📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
| auto pop = diversity_manager.get_population_pointer(); | |
| for (const auto& sol : pop->population_to_vector()) { | |
| hashes.push_back(sol.get_hash()); | |
| } | |
| uint32_t final_hash = detail::compute_hash(hashes); | |
| printf("%s: final hash: 0x%x, pop size %d\n", | |
| path.c_str(), | |
| final_hash, | |
| (int)pop->population_to_vector().size()); | |
| return final_hash; | |
| } | |
| return detail::compute_hash(hashes); | |
| } |
🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed.
In `@cpp/tests/mip/diversity_test.cu` around lines 267 - 279, The block that
collects hashes and computes final_hash (using
diversity_manager.get_population_pointer(), pop->population_to_vector(),
detail::compute_hash(hashes), and the printf that references path and
final_hash) is dead/unreachable because an earlier return happens; remove this
unreachable code or move its logic before the earlier return so only reachable
code remains—specifically delete the loop that pushes sol.get_hash(), the
compute_hash call, the printf, and the trailing return final_hash; ensure no
references to final_hash or hashes remain after the earlier return.
There was a problem hiding this comment.
Actionable comments posted: 1
🧹 Nitpick comments (6)
cpp/tests/mip/diversity_test.cu (4)
69-69: Empty function should be removed or documented.
setup_device_symbolsis a no-op function. If it's a placeholder for future implementation, add a TODO comment; otherwise, remove it to avoid confusion.💡 Suggested fix
Either remove the function and its call sites, or document the intent:
-static void setup_device_symbols(rmm::cuda_stream_view stream_view) { (void)stream_view; } +// TODO: Implement device symbol initialization if needed for determinism tests +// static void setup_device_symbols(rmm::cuda_stream_view stream_view) { (void)stream_view; }🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/tests/mip/diversity_test.cu` at line 69, The function setup_device_symbols(rmm::cuda_stream_view stream_view) is a no-op; either remove this unused function (and clean up any call sites that only exist to invoke it) or make its purpose explicit by adding a TODO/comment indicating it's a deliberate placeholder for future device-symbol setup; update references to setup_device_symbols in tests (e.g., any calls in diversity_test.cu) to reflect the removal or to clarify that the call is intentionally a no-op.
71-73: Unusedseedparameter in test runner functions.The
seedparameter intest_full_run_determinism(line 72) andtest_initial_solution_determinism(line 121) has a default value but is never used within the function body. The global seed is set externally viacuopt::seed_generator::set_seed(seed)in the calling test. Either remove the parameter or use it to seed something locally.💡 Suggested fix
If the seed is intended to be set locally within these functions:
static uint32_t test_full_run_determinism(std::string path, - unsigned long seed = std::random_device{}(), + unsigned long seed, float work_limit = 10.0f) { + cuopt::seed_generator::set_seed(seed); const raft::handle_t handle_{};Or remove the parameter if external seeding is intentional.
Also applies to: 120-121
🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/tests/mip/diversity_test.cu` around lines 71 - 73, The two test runner functions test_full_run_determinism and test_initial_solution_determinism declare an unused seed parameter; either remove the seed parameter from both function signatures (and update any callers) if external seeding is intended, or make the functions actually use the value by calling cuopt::seed_generator::set_seed(seed) (or equivalent local seeding) at the start of each function to apply the provided seed; pick one approach and apply it consistently to both functions.
382-394: Test instantiation includes duplicate entries.Lines 387-388 instantiate
mip/neos5.mpsandmip/gen-ip054.mps, which appear both in active entries and as commented-out entries (lines 385-386). The commented-out duplicates should be removed to avoid confusion.💡 Suggested cleanup
INSTANTIATE_TEST_SUITE_P(DiversityTest, DiversityTestParams, testing::Values( - // std::make_tuple("mip/gen-ip054.mps", 5.0f), - // std::make_tuple("mip/pk1.mps", 5.0f), std::make_tuple("mip/neos5.mps", 5.0f), std::make_tuple("mip/gen-ip054.mps", 5.0f), std::make_tuple("mip/pk1.mps", 5.0f), - // std::make_tuple("uccase9.mps"), - // std::make_tuple("mip/neos5.mps", 5.0f), std::make_tuple("mip/50v-10.mps", 5.0f) - // std::make_tuple("mip/rmatr200-p5.mps", 5.0f) ));🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/tests/mip/diversity_test.cu` around lines 382 - 394, The INSTANTIATE_TEST_SUITE_P block for DiversityTest (using DiversityTestParams) contains commented-out entries that duplicate active tuples (e.g., "mip/neos5.mps" and "mip/gen-ip054.mps"); to clean up, remove the commented duplicate lines so only one declaration remains per test case in the INSTANTIATE_TEST_SUITE_P(DiversityTest, DiversityTestParams, testing::Values(...)) invocation and leave the active std::make_tuple entries (such as "mip/neos5.mps", "mip/gen-ip054.mps", "mip/pk1.mps", "mip/50v-10.mps") intact.
283-314: Test structure is sound but consider removing debug logging in CI.The test correctly validates recombiner determinism by comparing hashes across runs. However, the debug-level logging setup (lines 286-288) may produce excessive output in CI. Consider either:
- Removing the logger configuration, or
- Guarding it behind an environment variable check
🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/tests/mip/diversity_test.cu` around lines 283 - 314, The test recombiners_deterministic currently forces debug logging by calling cuopt::default_logger().set_pattern, set_level and flush_on at the start of the test; remove these three logger configuration calls or wrap them in a runtime guard (e.g., check an env var like CUOPT_DEBUG_LOG) so CI runs don't emit excessive debug output; locate the calls in the TEST_P(DiversityTestParams, recombiners_deterministic) body and either delete the cuopt::default_logger().set_pattern(...), cuopt::default_logger().set_level(...), and cuopt::default_logger().flush_on(...) lines or condition them on std::getenv("CUOPT_DEBUG_LOG") before applying them.cpp/tests/mip/determinism_test.cu (1)
103-136: Consider removing unused helper function.
count_callbacks_with_origin(lines 115-121) is defined but not used anywhere in this file. If it's intended for future use, consider adding a comment; otherwise, it could be removed to avoid dead code.🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/tests/mip/determinism_test.cu` around lines 103 - 136, The helper function count_callbacks_with_origin is currently unused in this file; either remove its definition to eliminate dead code or, if you plan to use it later, add a brief explanatory comment above count_callbacks_with_origin indicating its intended future use (or mark it with a clear TODO), so reviewers know it’s intentional; locate the function by name and delete it or add the comment accordingly while leaving the other helpers (is_gpu_callback_origin, count_gpu_callbacks, count_branch_and_bound_callbacks) unchanged.cpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cu (1)
82-84: Model produces unreliable estimates for constraint-heavy problems.The negative coefficient for
n_constraints(-400) causes the model to underestimate work for problems with many constraints relative to variables. For such problems,estim_msclamps to 0, and the loop may overshoot the iteration estimate significantly.Since this is marked as a temporary placeholder (per the TODO), consider adding a tracking issue or ensuring the predictor replacement is prioritized.
Do you want me to open a tracking issue for replacing this heuristic model?
🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed. In `@cpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cu` around lines 82 - 84, The current heuristic for estim_ms uses a negative weight on op_problem.n_constraints which causes underestimation and clamping to 0 for constraint-heavy problems; update the formula in relaxed_lp.cu (the estim_ms computation using op_problem.n_variables, op_problem.n_constraints, op_problem.coefficients.size(), and estim_iters) to remove or make the n_constraints coefficient non-negative (e.g., replace -400 with a small positive or zero weight) and ensure a reasonable lower bound (e.g., min_ms > 0) instead of hard clamping to 0; also add a TODO comment and open a tracking issue to prioritize replacing this temporary predictor as referenced by the existing TODO so the heuristic is not relied on long-term.
🤖 Prompt for all review comments with AI agents
Verify each finding against the current code and only fix it if needed.
Inline comments:
In `@cpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cu`:
- Around line 79-90: The loop that increases estim_iters can overflow because
estim_iters is an int and may wrap if work_limit is huge; update the logic in
the block that computes estim_ms and increments estim_iters so it cannot run
unbounded: either change estim_iters to a wider integer type (e.g., int64_t)
and/or add an explicit upper bound check (compare against a safe MAX_ITERS
constant or std::numeric_limits<int>::max()/a guarded cap) and break when
reached, and ensure the estim_ms > work_limit*1000 check is preserved; locate
and modify the variables/loop referencing estim_iters, estim_ms, and work_limit
inside the do/while to add the guard and avoid signed overflow.
---
Nitpick comments:
In `@cpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cu`:
- Around line 82-84: The current heuristic for estim_ms uses a negative weight
on op_problem.n_constraints which causes underestimation and clamping to 0 for
constraint-heavy problems; update the formula in relaxed_lp.cu (the estim_ms
computation using op_problem.n_variables, op_problem.n_constraints,
op_problem.coefficients.size(), and estim_iters) to remove or make the
n_constraints coefficient non-negative (e.g., replace -400 with a small positive
or zero weight) and ensure a reasonable lower bound (e.g., min_ms > 0) instead
of hard clamping to 0; also add a TODO comment and open a tracking issue to
prioritize replacing this temporary predictor as referenced by the existing TODO
so the heuristic is not relied on long-term.
In `@cpp/tests/mip/determinism_test.cu`:
- Around line 103-136: The helper function count_callbacks_with_origin is
currently unused in this file; either remove its definition to eliminate dead
code or, if you plan to use it later, add a brief explanatory comment above
count_callbacks_with_origin indicating its intended future use (or mark it with
a clear TODO), so reviewers know it’s intentional; locate the function by name
and delete it or add the comment accordingly while leaving the other helpers
(is_gpu_callback_origin, count_gpu_callbacks, count_branch_and_bound_callbacks)
unchanged.
In `@cpp/tests/mip/diversity_test.cu`:
- Line 69: The function setup_device_symbols(rmm::cuda_stream_view stream_view)
is a no-op; either remove this unused function (and clean up any call sites that
only exist to invoke it) or make its purpose explicit by adding a TODO/comment
indicating it's a deliberate placeholder for future device-symbol setup; update
references to setup_device_symbols in tests (e.g., any calls in
diversity_test.cu) to reflect the removal or to clarify that the call is
intentionally a no-op.
- Around line 71-73: The two test runner functions test_full_run_determinism and
test_initial_solution_determinism declare an unused seed parameter; either
remove the seed parameter from both function signatures (and update any callers)
if external seeding is intended, or make the functions actually use the value by
calling cuopt::seed_generator::set_seed(seed) (or equivalent local seeding) at
the start of each function to apply the provided seed; pick one approach and
apply it consistently to both functions.
- Around line 382-394: The INSTANTIATE_TEST_SUITE_P block for DiversityTest
(using DiversityTestParams) contains commented-out entries that duplicate active
tuples (e.g., "mip/neos5.mps" and "mip/gen-ip054.mps"); to clean up, remove the
commented duplicate lines so only one declaration remains per test case in the
INSTANTIATE_TEST_SUITE_P(DiversityTest, DiversityTestParams,
testing::Values(...)) invocation and leave the active std::make_tuple entries
(such as "mip/neos5.mps", "mip/gen-ip054.mps", "mip/pk1.mps", "mip/50v-10.mps")
intact.
- Around line 283-314: The test recombiners_deterministic currently forces debug
logging by calling cuopt::default_logger().set_pattern, set_level and flush_on
at the start of the test; remove these three logger configuration calls or wrap
them in a runtime guard (e.g., check an env var like CUOPT_DEBUG_LOG) so CI runs
don't emit excessive debug output; locate the calls in the
TEST_P(DiversityTestParams, recombiners_deterministic) body and either delete
the cuopt::default_logger().set_pattern(...),
cuopt::default_logger().set_level(...), and
cuopt::default_logger().flush_on(...) lines or condition them on
std::getenv("CUOPT_DEBUG_LOG") before applying them.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Pro
Run ID: 6d186eaf-cff0-4117-9f13-2cb2e79896ba
📒 Files selected for processing (5)
cpp/src/branch_and_bound/pseudo_costs.cppcpp/src/mip_heuristics/local_search/rounding/constraint_prop.cucpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cucpp/tests/mip/determinism_test.cucpp/tests/mip/diversity_test.cu
🚧 Files skipped from review as they are similar to previous changes (2)
- cpp/src/branch_and_bound/pseudo_costs.cpp
- cpp/src/mip_heuristics/local_search/rounding/constraint_prop.cu
| if (!std::isinf(work_limit)) { | ||
| do { | ||
| // TODO: use an actual predictor model here | ||
| double estim_ms = 313 + 200 * op_problem.n_variables - 400 * op_problem.n_constraints + | ||
| 600 * op_problem.coefficients.size() + 7100 * estim_iters; | ||
| estim_ms = std::max(0.0, estim_ms); | ||
| if (estim_ms > work_limit * 1000) { break; } | ||
| estim_iters += 100; | ||
| } while (true); | ||
| } else { | ||
| estim_iters = std::numeric_limits<int>::max(); | ||
| } |
There was a problem hiding this comment.
Potential unbounded loop and signed integer overflow.
When work_limit is finite but very large (e.g., work_limit = 1e6 seconds), this loop may iterate until estim_iters overflows. Since estim_iters is a signed int, overflow is undefined behavior.
Consider adding bounds:
🛡️ Proposed fix to bound iterations
+ constexpr int max_estim_iters = std::numeric_limits<int>::max() - 100;
if (!std::isinf(work_limit)) {
do {
// TODO: use an actual predictor model here
double estim_ms = 313 + 200 * op_problem.n_variables - 400 * op_problem.n_constraints +
600 * op_problem.coefficients.size() + 7100 * estim_iters;
estim_ms = std::max(0.0, estim_ms);
if (estim_ms > work_limit * 1000) { break; }
+ if (estim_iters >= max_estim_iters) { break; }
estim_iters += 100;
} while (true);🤖 Prompt for AI Agents
Verify each finding against the current code and only fix it if needed.
In `@cpp/src/mip_heuristics/relaxed_lp/relaxed_lp.cu` around lines 79 - 90, The
loop that increases estim_iters can overflow because estim_iters is an int and
may wrap if work_limit is huge; update the logic in the block that computes
estim_ms and increments estim_iters so it cannot run unbounded: either change
estim_iters to a wider integer type (e.g., int64_t) and/or add an explicit upper
bound check (compare against a safe MAX_ITERS constant or
std::numeric_limits<int>::max()/a guarded cap) and break when reached, and
ensure the estim_ms > work_limit*1000 check is preserved; locate and modify the
variables/loop referencing estim_iters, estim_ms, and work_limit inside the
do/while to add the guard and avoid signed overflow.
|
🔔 Hi @anandhkb, this pull request has had no activity for 7 days. Please update or let us know if it can be closed. Thank you! If this is an "epic" issue, then please add the "epic" label to this issue. |
4 similar comments
|
🔔 Hi @anandhkb, this pull request has had no activity for 7 days. Please update or let us know if it can be closed. Thank you! If this is an "epic" issue, then please add the "epic" label to this issue. |
|
🔔 Hi @anandhkb, this pull request has had no activity for 7 days. Please update or let us know if it can be closed. Thank you! If this is an "epic" issue, then please add the "epic" label to this issue. |
|
🔔 Hi @anandhkb, this pull request has had no activity for 7 days. Please update or let us know if it can be closed. Thank you! If this is an "epic" issue, then please add the "epic" label to this issue. |
|
🔔 Hi @anandhkb, this pull request has had no activity for 7 days. Please update or let us know if it can be closed. Thank you! If this is an "epic" issue, then please add the "epic" label to this issue. |
|
🔔 Hi @anandhkb, this pull request has had no activity for 7 days. Please update or let us know if it can be closed. Thank you! If this is an "epic" issue, then please add the "epic" label to this issue. |
6 similar comments
|
🔔 Hi @anandhkb, this pull request has had no activity for 7 days. Please update or let us know if it can be closed. Thank you! If this is an "epic" issue, then please add the "epic" label to this issue. |
|
🔔 Hi @anandhkb, this pull request has had no activity for 7 days. Please update or let us know if it can be closed. Thank you! If this is an "epic" issue, then please add the "epic" label to this issue. |
|
🔔 Hi @anandhkb, this pull request has had no activity for 7 days. Please update or let us know if it can be closed. Thank you! If this is an "epic" issue, then please add the "epic" label to this issue. |
|
🔔 Hi @anandhkb, this pull request has had no activity for 7 days. Please update or let us know if it can be closed. Thank you! If this is an "epic" issue, then please add the "epic" label to this issue. |
|
🔔 Hi @anandhkb, this pull request has had no activity for 7 days. Please update or let us know if it can be closed. Thank you! If this is an "epic" issue, then please add the "epic" label to this issue. |
|
🔔 Hi @anandhkb, this pull request has had no activity for 7 days. Please update or let us know if it can be closed. Thank you! If this is an "epic" issue, then please add the "epic" label to this issue. |
3 participants
This PR implements the integration of GPU heuristics into the deterministic codepath of the solver. This enables running full end-to-end solves with presolve, parallel B&B, and GPU heuristics, with full run-to-run determinism and heterogeneous (CPU+GPU) execution.
For this purpose, the GPU heuristics were modified to eliminate sources of nondeterminism such as non-ordered floating point reductions, atomic operations, scheduling-dependent operations... Extensive testing has been added to the CI runs to ensure individual components (such as probing, FJ, recombiners, FP...) are bitwise deterministic, and their overall orchestration maintains this invariant.
Furthermore, to reduce changes to the codebase and splits in the codepath, wall-clock timers on the heuristics side were replaced with unified termination checking objects, that check for wall time (opportunistic mode), or work total (deterministic mode).
These timers are hierarchical, and thus allow for checking against multiple terminations sources, for example a local algo's work budget, the global wall clock time limit, or in future PRs user-controlled termination signals like Ctrl-C or callback handlers.
B&B acts as the authority for incumbent publishing to the user. At every horizon sync, the queue of GPU heuristic solutions is checked, and if solutions whose timestamp is <= the current B&B work total, they are drained, replayed, and incorporated into the tree exploration.
A new get_solution_callback_ext callback has been introduced to prevent ABI breaks and pass more metadata with each published solution. For now, only the solution origin and its work unit timestamp are exposed, but additional fields may be later added without breaking the ABI.
The benchmark runner now relies on such callbacks to record the incumbents and their timestamps - it is no longer necessary to parse the logs to compute the primal integral.
A few other bugfixes have been included to the existing B&B deterministic exploration regarding node queues. Other fixes include uninitialized memory accesses, and improvements to allow running compute-sanitizer initcheck without false positives. Determinism logs have also been added thorought the code, disabled by default.
Some features remain unsupported in deterministic mode, such as Papilo's Probing presolver, concurrent root solve, RINS, or the sub-MIP recombiner. These will be incorporated in later PRs.
Furthermore, the PR has been modified to take the ML estimators out for the sake of reducing code review workload. They will be incorporated again in later PRs, along with improved tuning.
No intentional API or ABI breaking changes introduced.
Opportunistic mode performance remains unchanged:
Primal gap 12.5%, 226 feasible, primal integral 19.1%.
Deterministic mode benchmarks pending.
Closes #144
Closes #882