Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
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
51 changes: 36 additions & 15 deletions cpp/src/mip_heuristics/diversity/diversity_manager.cu
Original file line number Diff line number Diff line change
Expand Up @@ -493,15 +493,22 @@ solution_t<i_t, f_t> diversity_manager_t<i_t, f_t>::run_solver()
timer_t lp_timer(lp_time_limit);
auto lp_result = solve_lp_with_method<i_t, f_t>(*problem_ptr, pdlp_settings, lp_timer);

// The concurrent root LP can fail to produce a usable solution -- e.g. the barrier
// hits a numerical error on an infeasible problem and PDLP returns NumericalError
// with empty primal/dual. In that case we must not copy or hand off the empty
// result (copying n elements from an empty buffer throws), and we must still
// release B&B's root-relaxation wait so it proceeds with its own dual-simplex root
// instead of spinning forever.
const bool root_lp_usable =
lp_result.get_termination_status() != pdlp_termination_status_t::NumericalError &&
lp_result.get_primal_solution().size() == lp_optimal_solution.size() &&
lp_result.get_dual_solution().size() == lp_dual_optimal_solution.size();

bool use_staged_simplex_solution = false;
{
std::lock_guard<std::mutex> guard(relaxed_solution_mutex);
use_staged_simplex_solution = simplex_solution_exists.load();
if (!use_staged_simplex_solution) {
cuopt_assert(lp_result.get_primal_solution().size() == lp_optimal_solution.size(),
"LP optimal solution size mismatch");
cuopt_assert(lp_result.get_dual_solution().size() == lp_dual_optimal_solution.size(),
"LP dual optimal solution size mismatch");
if (!use_staged_simplex_solution && root_lp_usable) {
raft::copy(lp_optimal_solution.data(),
lp_result.get_primal_solution().data(),
lp_optimal_solution.size(),
Expand All @@ -513,14 +520,26 @@ solution_t<i_t, f_t> diversity_manager_t<i_t, f_t>::run_solver()
}
}
if (use_staged_simplex_solution) { consume_staged_simplex_solution(lp_state); }
cuopt_assert(thrust::all_of(problem_ptr->handle_ptr->get_thrust_policy(),
lp_optimal_solution.begin(),
lp_optimal_solution.end(),
[] __host__ __device__(f_t val) { return std::isfinite(val); }),
"LP optimal solution contains non-finite values");
if (use_staged_simplex_solution || root_lp_usable) {
cuopt_assert(thrust::all_of(problem_ptr->handle_ptr->get_thrust_policy(),
lp_optimal_solution.begin(),
lp_optimal_solution.end(),
[] __host__ __device__(f_t val) { return std::isfinite(val); }),
"LP optimal solution contains non-finite values");
}
ls.lp_optimal_exists = true;
if (!use_staged_simplex_solution) {
if (lp_result.get_termination_status() == pdlp_termination_status_t::Optimal) {
if (!root_lp_usable) {
// The concurrent root LP produced no usable solution. Do not hand an empty
// solution to B&B; instead release its root-relaxation wait loop so it falls
// back to its own dual-simplex root rather than deadlocking.
CUOPT_LOG_DEBUG("Root LP produced no usable solution (status %d); releasing B&B root solve",
(int)lp_result.get_termination_status());
ls.lp_optimal_exists = false;
if (context.branch_and_bound_ptr != nullptr) {
context.branch_and_bound_ptr->set_root_concurrent_halt(1);
}
} else if (lp_result.get_termination_status() == pdlp_termination_status_t::Optimal) {
solution_t<i_t, f_t> lp_sol(*problem_ptr);
lp_sol.copy_new_assignment(lp_optimal_solution);
const bool consider_integrality = false;
Expand All @@ -541,9 +560,11 @@ solution_t<i_t, f_t> diversity_manager_t<i_t, f_t>::run_solver()
}
}

// Send relaxed solution to branch and bound only if PDLP found it (not dual simplex via
// set_simplex_solution)
if (!use_staged_simplex_solution &&
// Hand the root relaxation off to branch and bound when we have a usable solution
// (sets root_crossover_solution_set_, releasing B&B's wait). When the root LP failed
// the wait is instead released above via set_root_concurrent_halt, and a staged
// dual-simplex solution is owned by B&B already, so neither needs this hand-off.
if (!use_staged_simplex_solution && root_lp_usable &&
problem_ptr->set_root_relaxation_solution_callback != nullptr) {
auto& d_primal_solution = lp_result.get_primal_solution();
auto& d_dual_solution = lp_result.get_dual_solution();
Expand Down Expand Up @@ -576,7 +597,7 @@ solution_t<i_t, f_t> diversity_manager_t<i_t, f_t>::run_solver()
host_primal, host_dual, host_reduced_costs, solver_obj, user_obj, iterations, method);
}

if (!use_staged_simplex_solution) {
if (!use_staged_simplex_solution && root_lp_usable) {
// in case the pdlp returned var boudns that are out of bounds
clamp_within_var_bounds(lp_optimal_solution, problem_ptr, problem_ptr->handle_ptr);
}
Expand Down
43 changes: 34 additions & 9 deletions cpp/src/utilities/manual_cuda_graph.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -24,14 +24,21 @@ namespace cuopt {
// cuSPARSE calls inside the captured region are preserved.
//
// Invalidation recovery:
// If cudaStreamEndCapture returns cudaErrorStreamCaptureInvalidated
// (typically because another thread issued a synchronous CUDA call --
// A concurrent thread that issues a capture-hostile CUDA call --
// cudaDeviceSynchronize, cudaMalloc, cudaFree, or a library first-use that
// internally syncs the device -- concurrently with this capture window),
// the captured work has NOT been issued to the device. The wrapper drains
// the sticky error, re-executes `work` eagerly so the current iteration
// still produces correct results, and leaves itself uninitialized so the
// next `run` call retries capture.
// internally syncs the device (e.g. the cuDSS barrier's handle init) -- during
// this capture window invalidates the capture. That shows up in one of two ways,
// both handled here:
// 1. cudaStreamEndCapture returns cudaErrorStreamCaptureInvalidated, or
// 2. a CUDA / cuBLAS / cuSPARSE call inside `work` observes the invalidated
// capture and throws. Because cuBLAS/cuSPARSE cannot return a CUDA error
// code, this surfaces as e.g. CUBLAS_STATUS_INTERNAL_ERROR rather than the
// clean cudaErrorStreamCaptureInvalidated.
// In both cases the captured work has NOT been issued to the device. The wrapper
// drains the sticky error, re-executes `work` eagerly (no capture, so the
// concurrent op cannot break it) so the current iteration still produces correct
// results, and leaves itself uninitialized so the next `run` retries capture.
// A throw whose capture is NOT invalidated is a genuine error and is rethrown.
// IMPORTANT: because `work` is invoked a second time on recovery, any
// host-side mutations inside the callable will run twice -- keep `work`
// host-idempotent or move host bookkeeping (counters, flags, hash updates,
Expand Down Expand Up @@ -75,9 +82,27 @@ class manual_cuda_graph_t {
RAFT_CUDA_TRY(cudaStreamBeginCapture(stream.value(), cudaStreamCaptureModeThreadLocal));
guard.capture_active = true;

work();

cudaGraph_t captured = nullptr;
try {
work();
} catch (...) {
// A CUDA / cuBLAS / cuSPARSE call inside `work` threw mid-capture (commonly
// CUBLAS_STATUS_INTERNAL_ERROR when a concurrent capture-hostile op
// invalidated this capture and the failure was observed inside the library
// call). End the capture and let its status disambiguate: if the capture was
// invalidated the recorded work was never issued, so recover by re-running
// `work` eagerly; otherwise the error is genuine and is rethrown.
cudaError_t catch_end_err = cudaStreamEndCapture(stream.value(), &captured);
guard.capture_active = false;
if (catch_end_err == cudaErrorStreamCaptureInvalidated) {
cudaGetLastError();
work();
return;
}
if (captured != nullptr) { RAFT_CUDA_TRY_NO_THROW(cudaGraphDestroy(captured)); }
throw;
}

cudaError_t end_err = cudaStreamEndCapture(stream.value(), &captured);
guard.capture_active = false;

Expand Down
Loading