From cf701be621426c007425a990b6f913c374a52ae5 Mon Sep 17 00:00:00 2001 From: Yuwen Chen Date: Fri, 4 Sep 2026 07:39:54 -0700 Subject: [PATCH 1/2] Use deterministic cuSPARSE SpMV algorithm when cudss_deterministic is set The barrier solver produced different results run to run even with --cudss-deterministic 1: two runs were bit-identical for the first ten IPM iterations, then differed by 1 ULP and amplified to the fifth significant digit by iteration 100. The source was the SpMV algorithm selection. Per the cuSPARSE documentation CUSPARSE_SPMV_CSR_ALG1 "may produce slightly different results during different runs with the same input parameters", while ALG2 "provides deterministic (bit-wise) results for each run" - the reverse of what the previous comment here claimed. Measured on cuSPARSE 12.5.10 / sm_100a with the barrier's own operands, ALG1 differed on 199 of 199 repeats (5 of 7990 entries, 1 ULP) and gave 5 distinct results across 8 processes; ALG2 was bit-identical over 200 repeats and across processes. A single perturbed entry is enough to change the trajectory: the SpMV computing r = b - A*x feeds d_h_, which becomes the constraint block of the augmented system RHS, so cuDSS is handed a different right-hand side and returns a different search direction. cuDSS itself was verified deterministic for fixed input on both 0.7.1 and 0.9.0. Select ALG2 when determinism is requested, and thread the flag through cusparse_view_t so both spmv() and transpose_spmv() use it. The default path is unchanged, so runs that do not ask for determinism keep the faster ALG1. The pre-existing note about ALG2 and beta=1 accumulate mode on cuSPARSE < 13.0 is retained for the default path; that behaviour did not reproduce on 12.5.10, where ALG2 with beta=1 matched a CPU reference to 3.6e-15. Signed-off-by: Yuwen Chen Signed-off-by: yuwenchen95 --- cpp/src/barrier/barrier.cu | 4 ++-- cpp/src/barrier/cusparse_view.cu | 20 +++++++++++--------- cpp/src/barrier/cusparse_view.hpp | 5 ++++- 3 files changed, 17 insertions(+), 12 deletions(-) diff --git a/cpp/src/barrier/barrier.cu b/cpp/src/barrier/barrier.cu index e411a0d5dc..ad8284aeb7 100644 --- a/cpp/src/barrier/barrier.cu +++ b/cpp/src/barrier/barrier.cu @@ -270,8 +270,8 @@ class iteration_data_t { Hchol(0, 0), A(lp.A), Q(Qin), - cusparse_Q_view_(lp.handle_ptr, Q), - cusparse_view_(lp.handle_ptr, lp.A), + cusparse_Q_view_(lp.handle_ptr, Q, settings.cudss_deterministic), + cusparse_view_(lp.handle_ptr, lp.A, settings.cudss_deterministic), cusparse_info(lp.handle_ptr), device_AD(lp.num_cols, lp.num_rows, 0, lp.handle_ptr->get_stream()), device_A(lp.num_cols, lp.num_rows, 0, lp.handle_ptr->get_stream()), diff --git a/cpp/src/barrier/cusparse_view.cu b/cpp/src/barrier/cusparse_view.cu index 477200c5e9..0ad85f3434 100644 --- a/cpp/src/barrier/cusparse_view.cu +++ b/cpp/src/barrier/cusparse_view.cu @@ -115,13 +115,13 @@ void my_cusparsespmv_preprocess(cusparseHandle_t handle, } #endif -static cusparseSpMVAlg_t get_spmv_alg([[maybe_unused]] int num_rows) +static cusparseSpMVAlg_t get_spmv_alg([[maybe_unused]] int num_rows, bool deterministic) { - // ALG2 has a bug in cuSPARSE < 13.0 where beta=1 accumulate mode ignores existing y values. - // ALG1 uses a deterministic row-split algorithm, while ALG2 uses a merge-based - // algorithm that may be faster but can use atomics. ALG1 is safe for reproducibility. + // ALG2 is reported to have a bug in cuSPARSE < 13.0 where beta=1 accumulate mode ignores the + // existing y values. constexpr int cusparse_version = CUSPARSE_VER_MAJOR * 1000 + CUSPARSE_VER_MINOR * 100 + CUSPARSE_VER_PATCH; + if (deterministic) { return CUSPARSE_SPMV_CSR_ALG2; } if (cusparse_version < 13000) { return CUSPARSE_SPMV_CSR_ALG1; } return CUSPARSE_SPMV_CSR_ALG2; } @@ -133,7 +133,7 @@ void cusparse_view_t::init_spmv_buffer_and_preprocess(cusparseSpMatDes rmm::device_buffer& buffer, i_t rows) { - const auto spmv_alg = get_spmv_alg(rows); + const auto spmv_alg = get_spmv_alg(rows, deterministic_); size_t buffer_size_spmv = 0; RAFT_CUSPARSE_TRY( raft::sparse::detail::cusparsespmv_buffersize(handle_ptr_->get_cusparse_handle(), @@ -162,7 +162,8 @@ void cusparse_view_t::init_spmv_buffer_and_preprocess(cusparseSpMatDes template cusparse_view_t::cusparse_view_t(raft::handle_t const* handle_ptr, - const csc_matrix_t& A) + const csc_matrix_t& A, + bool deterministic) : handle_ptr_(handle_ptr), A_offsets_(0, handle_ptr->get_stream()), A_indices_(0, handle_ptr->get_stream()), @@ -174,7 +175,8 @@ cusparse_view_t::cusparse_view_t(raft::handle_t const* handle_ptr, spmv_buffer_transpose_(0, handle_ptr->get_stream()), d_one_(one_v, handle_ptr->get_stream()), d_minus_one_(neg_one_v, handle_ptr->get_stream()), - d_zero_(zero_v, handle_ptr->get_stream()) + d_zero_(zero_v, handle_ptr->get_stream()), + deterministic_(deterministic) { RAFT_CUBLAS_TRY(raft::linalg::detail::cublassetpointermode( handle_ptr->get_cublas_handle(), CUBLAS_POINTER_MODE_DEVICE, handle_ptr->get_stream())); @@ -302,7 +304,7 @@ void cusparse_view_t::spmv(f_t alpha, x, d_beta->data(), y, - get_spmv_alg(rows_), + get_spmv_alg(rows_, deterministic_), (f_t*)spmv_buffer_.data(), handle_ptr_->get_stream()); } @@ -357,7 +359,7 @@ void cusparse_view_t::transpose_spmv(f_t alpha, x, d_beta->data(), y, - get_spmv_alg(A_T_offsets_.size() - 1), + get_spmv_alg(A_T_offsets_.size() - 1, deterministic_), (f_t*)spmv_buffer_transpose_.data(), handle_ptr_->get_stream()); } diff --git a/cpp/src/barrier/cusparse_view.hpp b/cpp/src/barrier/cusparse_view.hpp index ea6bf363b9..db91df3f45 100644 --- a/cpp/src/barrier/cusparse_view.hpp +++ b/cpp/src/barrier/cusparse_view.hpp @@ -28,7 +28,9 @@ class cusparse_view_t { public: // Copy CSC -> owned CSR + CSC-transpose, with preprocess. Supports forward and transpose SpMV. // TMP matrix data should already be on the GPU and in CSR not CSC - cusparse_view_t(raft::handle_t const* handle_ptr, const csc_matrix_t& A); + cusparse_view_t(raft::handle_t const* handle_ptr, + const csc_matrix_t& A, + bool deterministic = false); ~cusparse_view_t(); pdlp::cusparse_dn_vec_descr_wrapper_t create_vector(rmm::device_uvector const& vec); @@ -80,5 +82,6 @@ class cusparse_view_t { rmm::device_scalar d_minus_one_; rmm::device_scalar d_zero_; i_t rows_{0}; + bool deterministic_{false}; }; } // namespace cuopt::mathematical_optimization::barrier From 173d37cc078034496e8812d547be92b000a8dc60 Mon Sep 17 00:00:00 2001 From: yuwenchen95 Date: Mon, 7 Sep 2026 02:34:12 -0700 Subject: [PATCH 2/2] Only apply ALG1 when ALG2 is problematic in certain conditions of CUDA 12 Signed-off-by: yuwenchen95 --- cpp/src/barrier/barrier.cu | 4 ++-- cpp/src/barrier/cusparse_view.cu | 38 +++++++++++++++---------------- cpp/src/barrier/cusparse_view.hpp | 9 ++++---- 3 files changed, 25 insertions(+), 26 deletions(-) diff --git a/cpp/src/barrier/barrier.cu b/cpp/src/barrier/barrier.cu index ad8284aeb7..e411a0d5dc 100644 --- a/cpp/src/barrier/barrier.cu +++ b/cpp/src/barrier/barrier.cu @@ -270,8 +270,8 @@ class iteration_data_t { Hchol(0, 0), A(lp.A), Q(Qin), - cusparse_Q_view_(lp.handle_ptr, Q, settings.cudss_deterministic), - cusparse_view_(lp.handle_ptr, lp.A, settings.cudss_deterministic), + cusparse_Q_view_(lp.handle_ptr, Q), + cusparse_view_(lp.handle_ptr, lp.A), cusparse_info(lp.handle_ptr), device_AD(lp.num_cols, lp.num_rows, 0, lp.handle_ptr->get_stream()), device_A(lp.num_cols, lp.num_rows, 0, lp.handle_ptr->get_stream()), diff --git a/cpp/src/barrier/cusparse_view.cu b/cpp/src/barrier/cusparse_view.cu index 0ad85f3434..a455bbd90e 100644 --- a/cpp/src/barrier/cusparse_view.cu +++ b/cpp/src/barrier/cusparse_view.cu @@ -115,14 +115,15 @@ void my_cusparsespmv_preprocess(cusparseHandle_t handle, } #endif -static cusparseSpMVAlg_t get_spmv_alg([[maybe_unused]] int num_rows, bool deterministic) +static cusparseSpMVAlg_t get_spmv_alg(int num_rows, int num_cols) { - // ALG2 is reported to have a bug in cuSPARSE < 13.0 where beta=1 accumulate mode ignores the - // existing y values. + // ALG2 provides deterministic (bit-wise) results but older ALG2 has a bug with matrices that have + // a single row or column, so fall back to ALG1 for those on cuSPARSE < 13.0.0. constexpr int cusparse_version = CUSPARSE_VER_MAJOR * 1000 + CUSPARSE_VER_MINOR * 100 + CUSPARSE_VER_PATCH; - if (deterministic) { return CUSPARSE_SPMV_CSR_ALG2; } - if (cusparse_version < 13000) { return CUSPARSE_SPMV_CSR_ALG1; } + if ((num_rows <= 1 || num_cols <= 1) && cusparse_version < 13000) { + return CUSPARSE_SPMV_CSR_ALG1; + } return CUSPARSE_SPMV_CSR_ALG2; } @@ -131,9 +132,10 @@ void cusparse_view_t::init_spmv_buffer_and_preprocess(cusparseSpMatDes cusparseDnVecDescr_t x, cusparseDnVecDescr_t y, rmm::device_buffer& buffer, - i_t rows) + i_t rows, + i_t cols) { - const auto spmv_alg = get_spmv_alg(rows, deterministic_); + const auto spmv_alg = get_spmv_alg(rows, cols); size_t buffer_size_spmv = 0; RAFT_CUSPARSE_TRY( raft::sparse::detail::cusparsespmv_buffersize(handle_ptr_->get_cusparse_handle(), @@ -162,8 +164,7 @@ void cusparse_view_t::init_spmv_buffer_and_preprocess(cusparseSpMatDes template cusparse_view_t::cusparse_view_t(raft::handle_t const* handle_ptr, - const csc_matrix_t& A, - bool deterministic) + const csc_matrix_t& A) : handle_ptr_(handle_ptr), A_offsets_(0, handle_ptr->get_stream()), A_indices_(0, handle_ptr->get_stream()), @@ -175,8 +176,7 @@ cusparse_view_t::cusparse_view_t(raft::handle_t const* handle_ptr, spmv_buffer_transpose_(0, handle_ptr->get_stream()), d_one_(one_v, handle_ptr->get_stream()), d_minus_one_(neg_one_v, handle_ptr->get_stream()), - d_zero_(zero_v, handle_ptr->get_stream()), - deterministic_(deterministic) + d_zero_(zero_v, handle_ptr->get_stream()) { RAFT_CUBLAS_TRY(raft::linalg::detail::cublassetpointermode( handle_ptr->get_cublas_handle(), CUBLAS_POINTER_MODE_DEVICE, handle_ptr->get_stream())); @@ -188,7 +188,7 @@ cusparse_view_t::cusparse_view_t(raft::handle_t const* handle_ptr, csr_matrix_t A_csr(A.m, A.n, 1); A.to_compressed_row(A_csr); rows_ = A_csr.m; - i_t cols = A_csr.n; + cols_ = A_csr.n; i_t nnz = A_csr.x.size(); const std::vector& offsets = A_csr.row_start; const std::vector& indices = A_csr.j; @@ -204,7 +204,7 @@ cusparse_view_t::cusparse_view_t(raft::handle_t const* handle_ptr, cusparseCreateCsr(&A_, rows_, - cols, + cols_, nnz, A_offsets_.data(), A_indices_.data(), @@ -215,7 +215,7 @@ cusparse_view_t::cusparse_view_t(raft::handle_t const* handle_ptr, CUDA_R_64F); cusparseCreateCsr(&A_T_, - cols, + cols_, rows_, nnz, A_T_offsets_.data(), @@ -229,13 +229,13 @@ cusparse_view_t::cusparse_view_t(raft::handle_t const* handle_ptr, // Tmp just to init the buffer size and preprocess cusparseDnVecDescr_t x; cusparseDnVecDescr_t y; - rmm::device_uvector d_x(cols, handle_ptr_->get_stream()); + rmm::device_uvector d_x(cols_, handle_ptr_->get_stream()); rmm::device_uvector d_y(rows_, handle_ptr_->get_stream()); RAFT_CUSPARSE_TRY(raft::sparse::detail::cusparsecreatednvec(&x, d_x.size(), d_x.data())); RAFT_CUSPARSE_TRY(raft::sparse::detail::cusparsecreatednvec(&y, d_y.size(), d_y.data())); - init_spmv_buffer_and_preprocess(A_, x, y, spmv_buffer_, rows_); - init_spmv_buffer_and_preprocess(A_T_, y, x, spmv_buffer_transpose_, A_T_offsets_.size() - 1); + init_spmv_buffer_and_preprocess(A_, x, y, spmv_buffer_, rows_, cols_); + init_spmv_buffer_and_preprocess(A_T_, y, x, spmv_buffer_transpose_, cols_, rows_); RAFT_CUSPARSE_TRY(cusparseDestroyDnVec(x)); RAFT_CUSPARSE_TRY(cusparseDestroyDnVec(y)); @@ -304,7 +304,7 @@ void cusparse_view_t::spmv(f_t alpha, x, d_beta->data(), y, - get_spmv_alg(rows_, deterministic_), + get_spmv_alg(rows_, cols_), (f_t*)spmv_buffer_.data(), handle_ptr_->get_stream()); } @@ -359,7 +359,7 @@ void cusparse_view_t::transpose_spmv(f_t alpha, x, d_beta->data(), y, - get_spmv_alg(A_T_offsets_.size() - 1, deterministic_), + get_spmv_alg(cols_, rows_), (f_t*)spmv_buffer_transpose_.data(), handle_ptr_->get_stream()); } diff --git a/cpp/src/barrier/cusparse_view.hpp b/cpp/src/barrier/cusparse_view.hpp index db91df3f45..63e6706609 100644 --- a/cpp/src/barrier/cusparse_view.hpp +++ b/cpp/src/barrier/cusparse_view.hpp @@ -28,9 +28,7 @@ class cusparse_view_t { public: // Copy CSC -> owned CSR + CSC-transpose, with preprocess. Supports forward and transpose SpMV. // TMP matrix data should already be on the GPU and in CSR not CSC - cusparse_view_t(raft::handle_t const* handle_ptr, - const csc_matrix_t& A, - bool deterministic = false); + cusparse_view_t(raft::handle_t const* handle_ptr, const csc_matrix_t& A); ~cusparse_view_t(); pdlp::cusparse_dn_vec_descr_wrapper_t create_vector(rmm::device_uvector const& vec); @@ -66,7 +64,8 @@ class cusparse_view_t { cusparseDnVecDescr_t x, cusparseDnVecDescr_t y, rmm::device_buffer& buffer, - i_t rows); + i_t rows, + i_t cols); rmm::device_uvector A_offsets_; rmm::device_uvector A_indices_; @@ -82,6 +81,6 @@ class cusparse_view_t { rmm::device_scalar d_minus_one_; rmm::device_scalar d_zero_; i_t rows_{0}; - bool deterministic_{false}; + i_t cols_{0}; }; } // namespace cuopt::mathematical_optimization::barrier