From e004040582a9c836ea750b7ce897907f3ce25204 Mon Sep 17 00:00:00 2001 From: akif Date: Thu, 20 Aug 2026 19:40:22 +0200 Subject: [PATCH 1/3] revert presolve tolerance --- cpp/src/mip_heuristics/presolve/third_party_presolve.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/mip_heuristics/presolve/third_party_presolve.cpp b/cpp/src/mip_heuristics/presolve/third_party_presolve.cpp index fdb5a0c88b..3d324d2007 100644 --- a/cpp/src/mip_heuristics/presolve/third_party_presolve.cpp +++ b/cpp/src/mip_heuristics/presolve/third_party_presolve.cpp @@ -724,7 +724,7 @@ void set_presolve_options(papilo::Presolve& presolver, { presolver.getPresolveOptions().tlim = time_limit; presolver.getPresolveOptions().threads = num_cpu_threads; // user setting or 0 (automatic) - presolver.getPresolveOptions().feastol = absolute_tolerance; + presolver.getPresolveOptions().feastol = 1e-5; if (max_rounds > 0) { presolver.getPresolveOptions().maxrounds = max_rounds; } if (dual_postsolve) { presolver.getPresolveOptions().componentsmaxint = -1; From 65b1a7cead10bacd7c3e793dd9ef4764d97562e4 Mon Sep 17 00:00:00 2001 From: akif Date: Thu, 27 Aug 2026 11:04:01 +0200 Subject: [PATCH 2/3] fix: synchronize cuDSS before concurrent halt Wait for interrupted cuDSS work before unwinding so asynchronous kernels cannot outlive solver state. Signed-off-by: akif --- cpp/src/barrier/sparse_cholesky.cuh | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/cpp/src/barrier/sparse_cholesky.cuh b/cpp/src/barrier/sparse_cholesky.cuh index 01045847d1..bc0637fe13 100644 --- a/cpp/src/barrier/sparse_cholesky.cuh +++ b/cpp/src/barrier/sparse_cholesky.cuh @@ -483,6 +483,8 @@ class sparse_cholesky_cudss_t : public sparse_cholesky_base_t { status = cudssExecute(handle, CUDSS_PHASE_REORDERING, solverConfig, solverData, A, cudss_x, cudss_b); if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { + RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); + handle_ptr_->get_stream().synchronize(); return CONCURRENT_HALT_RETURN; } if (status != CUDSS_STATUS_SUCCESS) { @@ -499,6 +501,8 @@ class sparse_cholesky_cudss_t : public sparse_cholesky_base_t { status = cudssExecute( handle, CUDSS_PHASE_SYMBOLIC_FACTORIZATION, solverConfig, solverData, A, cudss_x, cudss_b); if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { + RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); + handle_ptr_->get_stream().synchronize(); return CONCURRENT_HALT_RETURN; } if (status != CUDSS_STATUS_SUCCESS) { @@ -556,6 +560,8 @@ class sparse_cholesky_cudss_t : public sparse_cholesky_base_t { status = cudssExecute( handle, CUDSS_PHASE_FACTORIZATION, solverConfig, solverData, A, cudss_x, cudss_b); if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { + RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); + handle_ptr_->get_stream().synchronize(); return CONCURRENT_HALT_RETURN; } if (status != CUDSS_STATUS_SUCCESS) { @@ -572,6 +578,8 @@ class sparse_cholesky_cudss_t : public sparse_cholesky_base_t { f_t numeric_time = toc(start_numeric); if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { + RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); + handle_ptr_->get_stream().synchronize(); return CONCURRENT_HALT_RETURN; } @@ -701,6 +709,8 @@ class sparse_cholesky_cudss_t : public sparse_cholesky_base_t { f_t reorder_time = toc(start_analysis); if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { + RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); + handle_ptr_->get_stream().synchronize(); return CONCURRENT_HALT_RETURN; } @@ -767,6 +777,8 @@ class sparse_cholesky_cudss_t : public sparse_cholesky_base_t { f_t numeric_time = toc(start_numeric); if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { + RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); + handle_ptr_->get_stream().synchronize(); return CONCURRENT_HALT_RETURN; } @@ -832,6 +844,8 @@ class sparse_cholesky_cudss_t : public sparse_cholesky_base_t { status = cudssExecute(handle, CUDSS_PHASE_SOLVE, solverConfig, solverData, A, cudss_x, cudss_b); if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { + RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); + handle_ptr_->get_stream().synchronize(); return CONCURRENT_HALT_RETURN; } if (status != CUDSS_STATUS_SUCCESS) { From 8455505fc48d25f693a4047c3a3dc8812b271d6a Mon Sep 17 00:00:00 2001 From: akif Date: Thu, 27 Aug 2026 21:08:13 +0200 Subject: [PATCH 3/3] fix: preserve cuDSS status during concurrent halt --- cpp/src/barrier/sparse_cholesky.cuh | 17 +++++++++++------ 1 file changed, 11 insertions(+), 6 deletions(-) diff --git a/cpp/src/barrier/sparse_cholesky.cuh b/cpp/src/barrier/sparse_cholesky.cuh index bc0637fe13..4b6bb0adf2 100644 --- a/cpp/src/barrier/sparse_cholesky.cuh +++ b/cpp/src/barrier/sparse_cholesky.cuh @@ -702,17 +702,22 @@ class sparse_cholesky_cudss_t : public sparse_cholesky_base_t { return CONCURRENT_HALT_RETURN; } f_t start_analysis = tic(); - CUDSS_CALL_AND_CHECK( - cudssExecute(handle, CUDSS_PHASE_REORDERING, solverConfig, solverData, A, cudss_x, cudss_b), - status, - "cudssExecute for reordering"); - - f_t reorder_time = toc(start_analysis); + status = + cudssExecute(handle, CUDSS_PHASE_REORDERING, solverConfig, solverData, A, cudss_x, cudss_b); if (settings_.concurrent_halt != nullptr && *settings_.concurrent_halt == 1) { RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); handle_ptr_->get_stream().synchronize(); return CONCURRENT_HALT_RETURN; } + if (status != CUDSS_STATUS_SUCCESS) { + settings_.log.printf( + "FAILED: CUDSS call ended unsuccessfully with status = %d, details: cuDSSExecute for " + "reordering\n", + status); + return -1; + } + + f_t reorder_time = toc(start_analysis); f_t start_symbolic = tic();