From 933778ac2d9225bbf8ee207906591af674c2ff1c Mon Sep 17 00:00:00 2001 From: akif Date: Wed, 10 Jun 2026 15:02:11 +0200 Subject: [PATCH 1/2] fix hang and cublas bug --- .../diversity/diversity_manager.cu | 51 +++++++++++++------ cpp/src/utilities/manual_cuda_graph.cuh | 43 ++++++++++++---- det_one.py | 40 +++++++++++++++ determinism_milp_test.py | 43 ++++++++++++++++ 4 files changed, 153 insertions(+), 24 deletions(-) create mode 100644 det_one.py create mode 100644 determinism_milp_test.py 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; diff --git a/det_one.py b/det_one.py new file mode 100644 index 0000000000..31b57edb0c --- /dev/null +++ b/det_one.py @@ -0,0 +1,40 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +import sys +import numpy as np +from cuopt.linear_programming.problem import Problem, INTEGER, MINIMIZE +from cuopt.linear_programming import SolverSettings +from cuopt.linear_programming.solver.solver_parameters import ( + CUOPT_MIP_DETERMINISM_MODE, + CUOPT_TIME_LIMIT, +) + +mode = sys.argv[1] +rng = np.random.default_rng(0) +n = 120 +a = rng.uniform(1.0, 10.0, n) +b = a + rng.uniform(-0.5, 0.5, n) +need_a = 0.80 * a.sum() +cap_b = 0.30 * b.sum() + +p = Problem("infeasible_milp") +x = [ + p.addVariable(lb=0.0, ub=1.0, vtype=INTEGER, name=f"x{i}") + for i in range(n) +] +p.setObjective(sum(x), sense=MINIMIZE) +p.addConstraint( + sum(float(a[i]) * x[i] for i in range(n)) >= float(need_a), name="need_a" +) +p.addConstraint( + sum(float(b[i]) * x[i] for i in range(n)) <= float(cap_b), name="cap_b" +) + +s = SolverSettings() +s.set_parameter(CUOPT_TIME_LIMIT, 15.0) +if mode == "deterministic": + s.set_parameter(CUOPT_MIP_DETERMINISM_MODE, 1) +print(f"[det_one] solving mode={mode}", flush=True) +p.solve(s) +print("STATUS=" + str(getattr(p.Status, "name", p.Status)), flush=True) diff --git a/determinism_milp_test.py b/determinism_milp_test.py new file mode 100644 index 0000000000..24b445c4a2 --- /dev/null +++ b/determinism_milp_test.py @@ -0,0 +1,43 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +import subprocess +import sys +import textwrap + +WORKER = textwrap.dedent(""" + import sys + import numpy as np + from cuopt.linear_programming.problem import Problem, INTEGER, MINIMIZE + from cuopt.linear_programming import SolverSettings + from cuopt.linear_programming.solver.solver_parameters import ( + CUOPT_MIP_DETERMINISM_MODE, CUOPT_TIME_LIMIT) + + mode = sys.argv[1] + rng = np.random.default_rng(0) + n = 120 + a = rng.uniform(1.0, 10.0, n) + b = a + rng.uniform(-0.5, 0.5, n) # correlated second weight vector + need_a = 0.80 * a.sum() # capture >=80% of a-value ... + cap_b = 0.30 * b.sum() # ... using <=30% of b-value -> infeasible (a~b) + + p = Problem("infeasible_milp") + x = [p.addVariable(lb=0.0, ub=1.0, vtype=INTEGER, name=f"x{i}") for i in range(n)] + p.setObjective(sum(x), sense=MINIMIZE) + p.addConstraint(sum(float(a[i]) * x[i] for i in range(n)) >= float(need_a), name="need_a") + p.addConstraint(sum(float(b[i]) * x[i] for i in range(n)) <= float(cap_b), name="cap_b") + + s = SolverSettings() + s.set_parameter(CUOPT_TIME_LIMIT, 15.0) + if mode == "deterministic": + s.set_parameter(CUOPT_MIP_DETERMINISM_MODE, 1) + p.solve(s) + print("STATUS=" + str(getattr(p.Status, "name", p.Status))) +""") + +for mode in ["deterministic", "opportunistic"]: + r = subprocess.run( + [sys.executable, "-c", WORKER, mode], capture_output=True, text=True + ) + out = r.stdout.strip() or (r.stderr.strip().splitlines() or [""])[-1] + print(f"{mode:>13} -> exit {r.returncode:<4} | {out}") From e246d1f7b64296796b591eb9dad2c7cf8f2b38fd Mon Sep 17 00:00:00 2001 From: akif Date: Wed, 10 Jun 2026 15:02:40 +0200 Subject: [PATCH 2/2] remove test files --- det_one.py | 40 ------------------------------------- determinism_milp_test.py | 43 ---------------------------------------- 2 files changed, 83 deletions(-) delete mode 100644 det_one.py delete mode 100644 determinism_milp_test.py diff --git a/det_one.py b/det_one.py deleted file mode 100644 index 31b57edb0c..0000000000 --- a/det_one.py +++ /dev/null @@ -1,40 +0,0 @@ -# SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. -# SPDX-License-Identifier: Apache-2.0 - -import sys -import numpy as np -from cuopt.linear_programming.problem import Problem, INTEGER, MINIMIZE -from cuopt.linear_programming import SolverSettings -from cuopt.linear_programming.solver.solver_parameters import ( - CUOPT_MIP_DETERMINISM_MODE, - CUOPT_TIME_LIMIT, -) - -mode = sys.argv[1] -rng = np.random.default_rng(0) -n = 120 -a = rng.uniform(1.0, 10.0, n) -b = a + rng.uniform(-0.5, 0.5, n) -need_a = 0.80 * a.sum() -cap_b = 0.30 * b.sum() - -p = Problem("infeasible_milp") -x = [ - p.addVariable(lb=0.0, ub=1.0, vtype=INTEGER, name=f"x{i}") - for i in range(n) -] -p.setObjective(sum(x), sense=MINIMIZE) -p.addConstraint( - sum(float(a[i]) * x[i] for i in range(n)) >= float(need_a), name="need_a" -) -p.addConstraint( - sum(float(b[i]) * x[i] for i in range(n)) <= float(cap_b), name="cap_b" -) - -s = SolverSettings() -s.set_parameter(CUOPT_TIME_LIMIT, 15.0) -if mode == "deterministic": - s.set_parameter(CUOPT_MIP_DETERMINISM_MODE, 1) -print(f"[det_one] solving mode={mode}", flush=True) -p.solve(s) -print("STATUS=" + str(getattr(p.Status, "name", p.Status)), flush=True) diff --git a/determinism_milp_test.py b/determinism_milp_test.py deleted file mode 100644 index 24b445c4a2..0000000000 --- a/determinism_milp_test.py +++ /dev/null @@ -1,43 +0,0 @@ -# SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. -# SPDX-License-Identifier: Apache-2.0 - -import subprocess -import sys -import textwrap - -WORKER = textwrap.dedent(""" - import sys - import numpy as np - from cuopt.linear_programming.problem import Problem, INTEGER, MINIMIZE - from cuopt.linear_programming import SolverSettings - from cuopt.linear_programming.solver.solver_parameters import ( - CUOPT_MIP_DETERMINISM_MODE, CUOPT_TIME_LIMIT) - - mode = sys.argv[1] - rng = np.random.default_rng(0) - n = 120 - a = rng.uniform(1.0, 10.0, n) - b = a + rng.uniform(-0.5, 0.5, n) # correlated second weight vector - need_a = 0.80 * a.sum() # capture >=80% of a-value ... - cap_b = 0.30 * b.sum() # ... using <=30% of b-value -> infeasible (a~b) - - p = Problem("infeasible_milp") - x = [p.addVariable(lb=0.0, ub=1.0, vtype=INTEGER, name=f"x{i}") for i in range(n)] - p.setObjective(sum(x), sense=MINIMIZE) - p.addConstraint(sum(float(a[i]) * x[i] for i in range(n)) >= float(need_a), name="need_a") - p.addConstraint(sum(float(b[i]) * x[i] for i in range(n)) <= float(cap_b), name="cap_b") - - s = SolverSettings() - s.set_parameter(CUOPT_TIME_LIMIT, 15.0) - if mode == "deterministic": - s.set_parameter(CUOPT_MIP_DETERMINISM_MODE, 1) - p.solve(s) - print("STATUS=" + str(getattr(p.Status, "name", p.Status))) -""") - -for mode in ["deterministic", "opportunistic"]: - r = subprocess.run( - [sys.executable, "-c", WORKER, mode], capture_output=True, text=True - ) - out = r.stdout.strip() or (r.stderr.strip().splitlines() or [""])[-1] - print(f"{mode:>13} -> exit {r.returncode:<4} | {out}")