diff --git a/cpp/src/mip_heuristics/diversity/diversity_manager.cu b/cpp/src/mip_heuristics/diversity/diversity_manager.cu index 7b038d6fa6..b88099d8ae 100644 --- a/cpp/src/mip_heuristics/diversity/diversity_manager.cu +++ b/cpp/src/mip_heuristics/diversity/diversity_manager.cu @@ -493,15 +493,22 @@ solution_t diversity_manager_t::run_solver() timer_t lp_timer(lp_time_limit); auto lp_result = solve_lp_with_method(*problem_ptr, pdlp_settings, lp_timer); + // The concurrent root LP can fail to produce a usable solution -- e.g. the barrier + // 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 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(), @@ -513,14 +520,26 @@ solution_t diversity_manager_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 lp_sol(*problem_ptr); lp_sol.copy_new_assignment(lp_optimal_solution); const bool consider_integrality = false; @@ -541,9 +560,11 @@ solution_t diversity_manager_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(); @@ -576,7 +597,7 @@ solution_t diversity_manager_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); } diff --git a/cpp/src/utilities/manual_cuda_graph.cuh b/cpp/src/utilities/manual_cuda_graph.cuh index 68b37b7c71..d61cf04af8 100644 --- a/cpp/src/utilities/manual_cuda_graph.cuh +++ b/cpp/src/utilities/manual_cuda_graph.cuh @@ -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, @@ -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;