diff --git a/.github/workflows/linux.yml b/.github/workflows/linux.yml
index c7e2cfa..eb1463b 100644
--- a/.github/workflows/linux.yml
+++ b/.github/workflows/linux.yml
@@ -82,6 +82,13 @@
           apt-get install -y nvidia-cuda-toolkit
           echo "CUDACXX=/usr/lib/nvidia-cuda-toolkit/bin/nvcc" >> $GITHUB_ENV
 
+      - name: Setup cuDSS
+        if: matrix.gpu == 'cuda'
+        run: |
+          wget https://developer.download.nvidia.com/compute/cudss/redist/libcudss/linux-x86_64/libcudss-linux-x86_64-0.3.0.9_cuda12-archive.tar.xz
+          tar -xf libcudss-linux-x86_64-0.3.0.9_cuda12-archive.tar.xz
+          echo "cudss_DIR=${{github.workspace}}/libcudss-linux-x86_64-0.3.0.9_cuda12-archive/lib/cmake/cudss" >> $GITHUB_ENV
+
       - name: Cache Build
         id: cache-build
         uses: actions/cache@v4
diff --git a/.github/workflows/windows.yml b/.github/workflows/windows.yml
index eaffa79..182ed30 100644
--- a/.github/workflows/windows.yml
+++ b/.github/workflows/windows.yml
@@ -133,6 +133,17 @@
           echo "CUDA_PATH_V12_5=$CUDA_PATH" | Out-File -FilePath $env:GITHUB_ENV -Encoding utf8 -Append
           echo "$CUDA_PATH/bin" | Out-File -FilePath $env:GITHUB_PATH -Encoding utf8 -Append
 
+      - name: Setup cuDSS
+        if: matrix.gpu == 'cuda'
+        run: |
+          Invoke-WebRequest https://developer.download.nvidia.com/compute/cudss/redist/libcudss/windows-x86_64/libcudss-windows-x86_64-0.3.0.9_cuda12-archive.zip -OutFile libcudss-windows-x86_64-0.3.0.9_cuda12-archive.zip
+          Expand-Archive -Path .\libcudss-windows-x86_64-0.3.0.9_cuda12-archive.zip -DestinationPath "C:/"
+          Remove-Item .\libcudss-windows-x86_64-0.3.0.9_cuda12-archive.zip
+          $CUDSS_PATH = "C:/libcudss-windows-x86_64-0.3.0.9_cuda12-archive/"
+          $CUDSS_CMAKE_CONFIG = "$CUDSS_PATH/lib/cmake/cudss/cudss-config.cmake"
+          echo "cudss_DIR=$CUDSS_PATH/lib/cmake/cudss" | Out-File -FilePath $env:GITHUB_ENV -Encoding utf8 -Append
+          echo "$CUDSS_PATH/bin" | Out-File -FilePath $env:GITHUB_PATH -Encoding utf8 -Append
+
       - name: Cache gflags
         id: cache-gflags
         uses: actions/cache@v4
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 6cbc942..b92437c 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -278,10 +278,21 @@
     CUDA::cudart${CERES_CUDA_TARGET_SUFFIX}
     CUDA::cusolver${CERES_CUDA_TARGET_SUFFIX}
     CUDA::cusparse${CERES_CUDA_TARGET_SUFFIX})
+
+  find_package(cudss CONFIG)
+  if (cudss_FOUND)
+    set(cudss_DEPENDENCY "find_dependency(cudss ${cudss_VERSION})")
+    list(APPEND CERES_CUDA_LIBRARIES cudss${CERES_CUDA_TARGET_SUFFIX})
+  else (cudss_FOUND)
+    list(APPEND CERES_COMPILE_OPTIONS CERES_NO_CUDSS)
+    message("-- Did not find cuDSS, SPARSE_SCHUR and SPARSE_NORMAL_CHOLESKY with CUDA_SPARSE will not be available.")
+  endif (cudss_FOUND)
+
   unset (CERES_CUDA_TARGET_SUFFIX)
   set(CMAKE_CUDA_RUNTIME_LIBRARY NONE)
 else (USE_CUDA)
   message("-- Building without CUDA.")
+  list(APPEND CERES_COMPILE_OPTIONS CERES_NO_CUDSS)
   list(APPEND CERES_COMPILE_OPTIONS CERES_NO_CUDA)
 endif (USE_CUDA)
 
@@ -380,10 +391,10 @@
 endif()
 
 # Ensure that the user understands they have disabled all sparse libraries.
-if (NOT SUITESPARSE AND NOT EIGENSPARSE AND NOT ACCELERATESPARSE)
+if (NOT SUITESPARSE AND NOT EIGENSPARSE AND NOT ACCELERATESPARSE AND NOT cudss_FOUND)
   message("   ===============================================================")
-  message("   Compiling without any sparse library: SuiteSparse, ")
-  message("   EigenSparse & Apple's Accelerate are all disabled or unavailable.  ")
+  message("   Compiling without any sparse library: SuiteSparse, EigenSparse")
+  message("   Apple's Accelerate & cuDSS are all disabled or unavailable.  ")
   message("   No sparse linear solvers (SPARSE_NORMAL_CHOLESKY & SPARSE_SCHUR)")
   message("   will be available when Ceres is used.")
   message("   ===============================================================")
diff --git a/cmake/CeresCompileOptionsToComponents.cmake b/cmake/CeresCompileOptionsToComponents.cmake
index 64634d5..f4ed9ee 100644
--- a/cmake/CeresCompileOptionsToComponents.cmake
+++ b/cmake/CeresCompileOptionsToComponents.cmake
@@ -79,6 +79,8 @@
   add_to_output_if_not_found(CURRENT_CERES_COMPILE_OPTIONS ${CERES_COMPONENTS_VAR}
     CERES_NO_ACCELERATE_SPARSE "AccelerateSparse;SparseLinearAlgebraLibrary")
   add_to_output_if_not_found(CURRENT_CERES_COMPILE_OPTIONS ${CERES_COMPONENTS_VAR}
+    CERES_NO_CUDSS "cuDSS;SparseLinearAlgebraLibrary")
+  add_to_output_if_not_found(CURRENT_CERES_COMPILE_OPTIONS ${CERES_COMPONENTS_VAR}
     CERES_RESTRICT_SCHUR_SPECIALIZATION "SchurSpecializations")
   # Remove duplicates of SparseLinearAlgebraLibrary if multiple sparse backends
   # are present.
diff --git a/cmake/CeresConfig.cmake.in b/cmake/CeresConfig.cmake.in
index ceb7e26..b929c82 100644
--- a/cmake/CeresConfig.cmake.in
+++ b/cmake/CeresConfig.cmake.in
@@ -180,6 +180,7 @@
 @METIS_DEPENDENCY@
 @SuiteSparse_DEPENDENCY@
 @CUDAToolkit_DEPENDENCY@
+@cudss_DEPENDENCY@
 @Threads_DEPENDENCY@
 
 # As imported CMake targets are not re-exported when a dependent target is
diff --git a/cmake/config.h.in b/cmake/config.h.in
index 1566795..51db4a0 100644
--- a/cmake/config.h.in
+++ b/cmake/config.h.in
@@ -53,11 +53,15 @@
 // If defined, Ceres was compiled without CUDA.
 @CERES_NO_CUDA@
 
+// If defined, Ceres was compiled without cuDSS.
+@CERES_NO_CUDSS@
+
 // If defined, Ceres was compiled without Apple's Accelerate framework solvers.
 @CERES_NO_ACCELERATE_SPARSE@
 
 #if defined(CERES_NO_SUITESPARSE) &&              \
     defined(CERES_NO_ACCELERATE_SPARSE) &&        \
+    defined(CERES_NO_CUDSS) &&                    \
     !defined(CERES_USE_EIGEN_SPARSE)  // NOLINT
 // If defined Ceres was compiled without any sparse linear algebra support.
 #define CERES_NO_SPARSE
@@ -90,6 +94,9 @@
 #if defined(CERES_USE_EIGEN_SPARSE)
 #error CERES_NO_SPARSE requires !CERES_USE_EIGEN_SPARSE
 #endif
+#if !defined(CERES_NO_CUDSS)
+#error CERES_NO_SPARSE required CERES_NO_CUDSS
+#endif
 #endif
 
 #endif  // CERES_PUBLIC_INTERNAL_CONFIG_H_
diff --git a/examples/bundle_adjuster.cc b/examples/bundle_adjuster.cc
index 582ae2e..8df0d8a 100644
--- a/examples/bundle_adjuster.cc
+++ b/examples/bundle_adjuster.cc
@@ -96,8 +96,7 @@
             "Use power series expansion to initialize the solution in ITERATIVE_SCHUR linear solver.");
 
 DEFINE_string(sparse_linear_algebra_library, "suite_sparse",
-              "Options are: suite_sparse, accelerate_sparse, eigen_sparse, and "
-              "cuda_sparse.");
+              "Options are: suite_sparse, accelerate_sparse, eigen_sparse and cuda_sparse");
 DEFINE_string(dense_linear_algebra_library, "eigen",
               "Options are: eigen, lapack, and cuda");
 DEFINE_string(ordering_type, "amd", "Options are: amd, nesdis");
@@ -381,6 +380,7 @@
 int main(int argc, char** argv) {
   GFLAGS_NAMESPACE::ParseCommandLineFlags(&argc, &argv, true);
   google::InitGoogleLogging(argv[0]);
+
   if (CERES_GET_FLAG(FLAGS_input).empty()) {
     LOG(ERROR) << "Usage: bundle_adjuster --input=bal_problem";
     return 1;
diff --git a/include/ceres/types.h b/include/ceres/types.h
index 6e19c51..f5b66a9 100644
--- a/include/ceres/types.h
+++ b/include/ceres/types.h
@@ -175,7 +175,7 @@
   // Apple's Accelerate framework sparse linear algebra routines.
   ACCELERATE_SPARSE,
 
-  // Nvidia's cuSPARSE library.
+  // Nvidia's cuDSS and cuSPARSE libraries.
   CUDA_SPARSE,
 
   // No sparse linear solver should be used.  This does not necessarily
diff --git a/internal/ceres/CMakeLists.txt b/internal/ceres/CMakeLists.txt
index f4e40fb..6c1fe93 100644
--- a/internal/ceres/CMakeLists.txt
+++ b/internal/ceres/CMakeLists.txt
@@ -170,6 +170,7 @@
     cuda_partitioned_block_sparse_crs_view.cc
     cuda_block_structure.cc
     cuda_sparse_matrix.cc
+    cuda_sparse_cholesky.cc
     cuda_vector.cc
     dense_cholesky.cc
     dense_normal_cholesky_solver.cc
diff --git a/internal/ceres/context_impl.cc b/internal/ceres/context_impl.cc
index 2b9d9cc..073dc7a 100644
--- a/internal/ceres/context_impl.cc
+++ b/internal/ceres/context_impl.cc
@@ -60,6 +60,12 @@
     cusparseDestroy(cusparse_handle_);
     cusparse_handle_ = nullptr;
   }
+#ifndef CERES_NO_CUDSS
+  if (cudss_handle_ != nullptr) {
+    cudssDestroy(cudss_handle_);
+    cudss_handle_ = nullptr;
+  }
+#endif  // CERES_NO_CUDSS
   for (auto& s : streams_) {
     if (s != nullptr) {
       cudaStreamDestroy(s);
@@ -157,6 +163,13 @@
     return false;
   }
   event_logger.AddEvent("cusparseCreate");
+#ifndef CERES_NO_CUDSS
+  if (cudssCreate(&cudss_handle_) != CUDSS_STATUS_SUCCESS) {
+    *message = "CUDA initialization failed because cudssCreate() failed.";
+    TearDown();
+    return false;
+  }
+#endif  // CERES_NO_CUDSS
   for (auto& s : streams_) {
     if (cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking) != cudaSuccess) {
       *message =
@@ -172,7 +185,11 @@
       cublasSetStream(cublas_handle_, DefaultStream()) !=
           CUBLAS_STATUS_SUCCESS ||
       cusparseSetStream(cusparse_handle_, DefaultStream()) !=
-          CUSPARSE_STATUS_SUCCESS) {
+          CUSPARSE_STATUS_SUCCESS
+#ifndef CERES_NO_CUDSS
+      || cudssSetStream(cudss_handle_, DefaultStream()) != CUDSS_STATUS_SUCCESS
+#endif  // CERES_NO_CUDSS
+  ) {
     *message = "CUDA initialization failed because SetStream failed.";
     TearDown();
     return false;
diff --git a/internal/ceres/context_impl.h b/internal/ceres/context_impl.h
index 46692e6..cc301ee 100644
--- a/internal/ceres/context_impl.h
+++ b/internal/ceres/context_impl.h
@@ -47,6 +47,9 @@
 #include "cuda_runtime.h"
 #include "cusolverDn.h"
 #include "cusparse.h"
+#ifndef CERES_NO_CUDSS
+#include "cudss.h"
+#endif  // CERES_NO_CUDSS
 #endif  // CERES_NO_CUDA
 
 #include "ceres/thread_pool.h"
@@ -122,6 +125,9 @@
 
   cusolverDnHandle_t cusolver_handle_ = nullptr;
   cublasHandle_t cublas_handle_ = nullptr;
+#ifndef CERES_NO_CUDSS
+  cudssHandle_t cudss_handle_ = nullptr;
+#endif  // CERES_NO_CUDSS
 
   // Default stream.
   // Kernel invocations and memory copies on this stream can be left without
diff --git a/internal/ceres/cuda_buffer.h b/internal/ceres/cuda_buffer.h
index 40048fd..95504ca 100644
--- a/internal/ceres/cuda_buffer.h
+++ b/internal/ceres/cuda_buffer.h
@@ -36,6 +36,8 @@
 
 #ifndef CERES_NO_CUDA
 
+#include <cstddef>
+#include <utility>
 #include <vector>
 
 #include "cuda_runtime.h"
@@ -165,6 +167,55 @@
   size_t size_ = 0;
   ContextImpl* context_ = nullptr;
 };
+
+// This class wraps host memory region allocated via cudaMallocHost. Such memory
+// region is page-locked, hence enabling direct transfer to/from device,
+// avoiding implicit buffering under the hood of CUDA API.
+template <typename T>
+class CudaPinnedHostBuffer {
+ public:
+  CudaPinnedHostBuffer() noexcept = default;
+  CudaPinnedHostBuffer(int size) { Reserve(size); }
+  CudaPinnedHostBuffer(CudaPinnedHostBuffer&& other) noexcept
+      : data_(std::exchange(other.data_, nullptr)),
+        size_(std::exchange(other.size_, 0)) {}
+  CudaPinnedHostBuffer(const CudaPinnedHostBuffer&) = delete;
+  CudaPinnedHostBuffer& operator=(const CudaPinnedHostBuffer&) = delete;
+  CudaPinnedHostBuffer& operator=(CudaPinnedHostBuffer&& other) noexcept {
+    Free();
+    data_ = std::exchange(other.data_, nullptr);
+    size_ = std::exchange(other.size_, 0);
+    return *this;
+  }
+  ~CudaPinnedHostBuffer() { Free(); }
+
+  void Reserve(const std::size_t size) {
+    if (size > size_) {
+      Free();
+      CHECK_EQ(cudaMallocHost(&data_, size * sizeof(T)), cudaSuccess)
+          << "Failed to allocate " << size * sizeof(T)
+          << " bytes of pinned host memory";
+      size_ = size;
+    }
+  }
+
+  T* data() noexcept { return data_; }
+  const T* data() const noexcept { return data_; }
+  std::size_t size() const noexcept { return size_; }
+
+ private:
+  void Free() {
+    if (data_ != nullptr) {
+      CHECK_EQ(cudaFreeHost(data_), cudaSuccess);
+      data_ = nullptr;
+      size_ = 0;
+    }
+  }
+
+  T* data_ = nullptr;
+  std::size_t size_ = 0;
+};
+
 }  // namespace ceres::internal
 
 #endif  // CERES_NO_CUDA
diff --git a/internal/ceres/cuda_sparse_cholesky.cc b/internal/ceres/cuda_sparse_cholesky.cc
new file mode 100644
index 0000000..d4a285b
--- /dev/null
+++ b/internal/ceres/cuda_sparse_cholesky.cc
@@ -0,0 +1,485 @@
+// Ceres Solver - A fast non-linear least squares minimizer
+// Copyright 2024 Google Inc. All rights reserved.
+// http://ceres-solver.org/
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// * Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// * Neither the name of Google Inc. nor the names of its contributors may be
+//   used to endorse or promote products derived from this software without
+//   specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+// POSSIBILITY OF SUCH DAMAGE.
+//
+// Author: markshachkov@gmail.com (Mark Shachkov)
+
+#include "ceres/cuda_sparse_cholesky.h"
+
+#ifndef CERES_NO_CUDSS
+
+#include <cstddef>
+#include <iostream>
+#include <memory>
+#include <string>
+#include <type_traits>
+
+#include "Eigen/Core"
+#include "ceres/compressed_row_sparse_matrix.h"
+#include "ceres/cuda_buffer.h"
+#include "ceres/linear_solver.h"
+#include "cudss.h"
+
+namespace ceres::internal {
+
+inline std::string cuDSSStatusToString(cudssStatus_t status) {
+  switch (status) {
+    case CUDSS_STATUS_SUCCESS:
+      return "CUDSS_STATUS_SUCCESS";
+    case CUDSS_STATUS_NOT_INITIALIZED:
+      return "CUDSS_STATUS_NOT_INITIALIZED";
+    case CUDSS_STATUS_ALLOC_FAILED:
+      return "CUDSS_STATUS_ALLOC_FAILED";
+    case CUDSS_STATUS_INVALID_VALUE:
+      return "CUDSS_STATUS_INVALID_VALUE";
+    case CUDSS_STATUS_NOT_SUPPORTED:
+      return "CUDSS_STATUS_NOT_SUPPORTED";
+    case CUDSS_STATUS_EXECUTION_FAILED:
+      return "CUDSS_STATUS_EXECUTION_FAILED";
+    case CUDSS_STATUS_INTERNAL_ERROR:
+      return "CUDSS_STATUS_INTERNAL_ERROR";
+    default:
+      return "unknown cuDSS status: " + std::to_string(status);
+  }
+}
+
+#define CUDSS_STATUS_CHECK(IN)                                     \
+  if (cudssStatus_t status = IN; status != CUDSS_STATUS_SUCCESS) { \
+    CHECK(false) << "Got error: " << cuDSSStatusToString(status);  \
+  }
+
+#define CUDSS_STATUS_OK_OR_RETURN_FATAL_ERROR(IN, additional_message)    \
+  if (cudssStatus_t status = IN; status != CUDSS_STATUS_SUCCESS) {       \
+    *message = std::string(additional_message) +                         \
+               " Got error: " + cuDSSStatusToString(status);             \
+    return factorize_result_ = LinearSolverTerminationType::FATAL_ERROR; \
+  }
+
+#define CUDSS_STATUS_OK_OR_RETURN_CUDSS_STATUS(IN)                 \
+  if (cudssStatus_t status = IN; status != CUDSS_STATUS_SUCCESS) { \
+    return status;                                                 \
+  }
+
+class CERES_NO_EXPORT CuDSSMatrixBase {
+ public:
+  CuDSSMatrixBase() = default;
+  CuDSSMatrixBase(const CuDSSMatrixBase&) = delete;
+  CuDSSMatrixBase(CuDSSMatrixBase&&) = delete;
+  CuDSSMatrixBase& operator=(const CuDSSMatrixBase&) = delete;
+  CuDSSMatrixBase& operator=(CuDSSMatrixBase&&) = delete;
+  ~CuDSSMatrixBase() { CUDSS_STATUS_CHECK(Free()); }
+
+  cudssStatus_t Free() noexcept {
+    if (matrix_) {
+      const auto status = cudssMatrixDestroy(matrix_);
+      matrix_ = nullptr;
+      return status;
+    }
+
+    return CUDSS_STATUS_SUCCESS;
+  }
+
+  cudssMatrix_t Get() const noexcept { return matrix_; }
+
+ protected:
+  cudssMatrix_t matrix_{nullptr};
+};
+
+class CERES_NO_EXPORT CuDSSMatrixCSR : public CuDSSMatrixBase {
+ public:
+  cudssStatus_t Reset(int64_t num_rows,
+                      int64_t num_cols,
+                      int64_t num_nonzeros,
+                      void* rows_start,
+                      void* rows_end,
+                      void* cols,
+                      void* values,
+                      cudaDataType_t index_type,
+                      cudaDataType_t value_type,
+                      cudssMatrixType_t matrix_type,
+                      cudssMatrixViewType_t matrix_storage_type,
+                      cudssIndexBase_t index_base) {
+    CUDSS_STATUS_OK_OR_RETURN_CUDSS_STATUS(Free());
+
+    return cudssMatrixCreateCsr(&matrix_,
+                                num_rows,
+                                num_cols,
+                                num_nonzeros,
+                                rows_start,
+                                rows_end,
+                                cols,
+                                values,
+                                index_type,
+                                value_type,
+                                matrix_type,
+                                matrix_storage_type,
+                                index_base);
+  }
+};
+
+class CERES_NO_EXPORT CuDSSMatrixDense : public CuDSSMatrixBase {
+ public:
+  cudssStatus_t Reset(int64_t num_rows,
+                      int64_t num_cols,
+                      int64_t leading_dimension_size,
+                      void* values,
+                      cudaDataType_t value_type,
+                      cudssLayout_t layout) {
+    CUDSS_STATUS_OK_OR_RETURN_CUDSS_STATUS(Free());
+
+    return cudssMatrixCreateDn(&matrix_,
+                               num_rows,
+                               num_cols,
+                               leading_dimension_size,
+                               values,
+                               value_type,
+                               layout);
+  }
+};
+
+struct CudssContext {
+  CudssContext(cudssHandle_t cudss_handle) : cudss_handle_(cudss_handle) {
+    CUDSS_STATUS_CHECK(cudssConfigCreate(&solver_config_));
+    CUDSS_STATUS_CHECK(cudssDataCreate(cudss_handle_, &solver_data_));
+  }
+  CudssContext(const CudssContext&) = delete;
+  CudssContext(CudssContext&&) = delete;
+  CudssContext& operator=(const CudssContext&) = delete;
+  CudssContext& operator=(CudssContext&&) = delete;
+  ~CudssContext() {
+    CUDSS_STATUS_CHECK(cudssDataDestroy(cudss_handle_, solver_data_));
+    CUDSS_STATUS_CHECK(cudssConfigDestroy(solver_config_));
+  }
+
+  cudssHandle_t cudss_handle_{nullptr};
+  cudssConfig_t solver_config_{nullptr};
+  cudssData_t solver_data_{nullptr};
+};
+
+template <typename Scalar>
+class CERES_NO_EXPORT CudaSparseCholeskyImpl final : public SparseCholesky {
+ public:
+  static_assert(std::is_same_v<Scalar, float> || std::is_same_v<Scalar, double>,
+                "Scalar type is unsuported by cuDSS");
+  static constexpr cudaDataType_t kCuDSSScalar =
+      std::is_same_v<Scalar, float> ? CUDA_R_32F : CUDA_R_64F;
+
+  CudaSparseCholeskyImpl(ContextImpl* context)
+      : context_(context),
+        lhs_cols_d_(context_),
+        lhs_rows_d_(context_),
+        lhs_values_d_(context_),
+        rhs_d_(context_),
+        x_d_(context_) {}
+  CudaSparseCholeskyImpl(const CudaSparseCholeskyImpl&) = delete;
+  CudaSparseCholeskyImpl(CudaSparseCholeskyImpl&&) = delete;
+  CudaSparseCholeskyImpl& operator=(const CudaSparseCholeskyImpl&) = delete;
+  CudaSparseCholeskyImpl& operator=(CudaSparseCholeskyImpl&&) = delete;
+  ~CudaSparseCholeskyImpl() = default;
+
+  CompressedRowSparseMatrix::StorageType StorageType() const {
+    return CompressedRowSparseMatrix::StorageType::LOWER_TRIANGULAR;
+  }
+
+  LinearSolverTerminationType Factorize(CompressedRowSparseMatrix* lhs,
+                                        std::string* message) {
+    if (lhs->num_rows() != lhs->num_cols()) {
+      *message = "lhs matrix must be square";
+      return factorize_result_ = LinearSolverTerminationType::FATAL_ERROR;
+    }
+    if (lhs->storage_type() != StorageType()) {
+      *message = "lhs matrix must be lower triangular";
+      return factorize_result_ = LinearSolverTerminationType::FATAL_ERROR;
+    }
+
+    // If, after previous attempt to factorize, cudssDataGet(CUDSS_DATA_INFO)
+    // returned a numerical error, such error will be preserved by cuDSS 0.3.0
+    // and returned by cudssDataGet(CUDSS_DATA_INFO) even after correctly
+    // factorizeable matrix is provided. Such behaviour forces us to reset a
+    // cudssData_t object (managed by CudssContext) and to loose a result of
+    // anylyze stage, thus we have to perform analyze one more time.
+    // TODO: do not re-perform analyze in case of failed factorization numerics
+    if (analyze_result_ != LinearSolverTerminationType::SUCCESS ||
+        factorize_result_ != LinearSolverTerminationType::SUCCESS) {
+      analyze_result_ = Analyze(lhs, message);
+      if (analyze_result_ != LinearSolverTerminationType::SUCCESS) {
+        return analyze_result_;
+      }
+    }
+    CHECK_NE(cudss_context_.get(), nullptr);
+
+    ConvertAndCopyToDevice(lhs->values(), lhs_values_h_.data(), lhs_values_d_);
+
+    CUDSS_STATUS_OK_OR_RETURN_FATAL_ERROR(
+        cudssExecute(context_->cudss_handle_,
+                     CUDSS_PHASE_FACTORIZATION,
+                     cudss_context_->solver_config_,
+                     cudss_context_->solver_data_,
+                     cudss_lhs_.Get(),
+                     cudss_x_.Get(),
+                     cudss_rhs_.Get()),
+        "cudssExecute with CUDSS_PHASE_FACTORIZATION failed");
+
+    int cudss_data_info;
+    CUDSS_STATUS_OK_OR_RETURN_FATAL_ERROR(
+        GetCudssDataInfo(cudss_data_info),
+        "cudssDataGet with CUDSS_DATA_INFO failed");
+    const auto factorization_status =
+        static_cast<cudssStatus_t>(cudss_data_info);
+
+    if (factorization_status == CUDSS_STATUS_SUCCESS) {
+      return factorize_result_ = LinearSolverTerminationType::SUCCESS;
+    }
+
+    if (cudss_data_info > 0) {
+      return factorize_result_ = LinearSolverTerminationType::FAILURE;
+    }
+
+    return factorize_result_ = LinearSolverTerminationType::FATAL_ERROR;
+  }
+
+  LinearSolverTerminationType Solve(const double* rhs,
+                                    double* solution,
+                                    std::string* message) {
+    CHECK_NE(cudss_context_.get(), nullptr);
+
+    if (factorize_result_ != LinearSolverTerminationType::SUCCESS) {
+      *message = "Factorize did not complete successfully previously.";
+      return factorize_result_;
+    }
+
+    ConvertAndCopyToDevice(rhs, rhs_h_.data(), rhs_d_);
+
+    CUDSS_STATUS_OK_OR_RETURN_FATAL_ERROR(
+        cudssExecute(context_->cudss_handle_,
+                     CUDSS_PHASE_SOLVE,
+                     cudss_context_->solver_config_,
+                     cudss_context_->solver_data_,
+                     cudss_lhs_.Get(),
+                     cudss_x_.Get(),
+                     cudss_rhs_.Get()),
+        "cudssExecute with CUDSS_PHASE_SOLVE failed");
+
+    ConvertAndCopyToHost(x_d_, x_h_.data(), solution);
+
+    int cudss_data_info;
+    CUDSS_STATUS_OK_OR_RETURN_FATAL_ERROR(
+        GetCudssDataInfo(cudss_data_info),
+        "cudssDataGet with CUDSS_DATA_INFO failed");
+    const auto solve_status = static_cast<cudssStatus_t>(cudss_data_info);
+
+    if (solve_status != CUDSS_STATUS_SUCCESS) {
+      return LinearSolverTerminationType::FAILURE;
+    }
+
+    return LinearSolverTerminationType::SUCCESS;
+  }
+
+ private:
+  cudssStatus_t GetCudssDataInfo(int& cudss_data_info) {
+    CHECK_NE(cudss_context_.get(), nullptr);
+    std::size_t size_written = 0;
+    CUDSS_STATUS_OK_OR_RETURN_CUDSS_STATUS(
+        cudssDataGet(context_->cudss_handle_,
+                     cudss_context_->solver_data_,
+                     CUDSS_DATA_INFO,
+                     &cudss_data_info,
+                     sizeof(cudss_data_info),
+                     &size_written));
+    // TODO: enable following check after cudssDataGet will be fixed
+    // CHECK_EQ(size_written, sizeof(cudss_data_info));
+
+    return CUDSS_STATUS_SUCCESS;
+  }
+
+  LinearSolverTerminationType Analyze(const CompressedRowSparseMatrix* lhs,
+                                      std::string* message) {
+    if (auto status = SetupCudssMatrices(lhs, message);
+        status != LinearSolverTerminationType::SUCCESS) {
+      return status;
+    }
+
+    lhs_rows_d_.CopyFromCpu(lhs->rows(), lhs->num_rows() + 1);
+    lhs_cols_d_.CopyFromCpu(lhs->cols(), lhs->num_nonzeros());
+
+    // Analyze and factorization results are stored in cudssData_t (managed by
+    // CudssContext). Given that cuDSS 0.3.0 does not reset it's error state in
+    // case of failed numerics at factorization stage, we have to reset
+    // cudssData_t and to recompute an analyze stage while trying to factorize a
+    // rescaled matrix with the same structure.
+    // TODO: move creation of CudssContext to ctor of CudaSparseCholeskyImpl
+    cudss_context_ = std::make_unique<CudssContext>(context_->cudss_handle_);
+
+    CUDSS_STATUS_OK_OR_RETURN_FATAL_ERROR(
+        cudssExecute(context_->cudss_handle_,
+                     CUDSS_PHASE_ANALYSIS,
+                     cudss_context_->solver_config_,
+                     cudss_context_->solver_data_,
+                     cudss_lhs_.Get(),
+                     cudss_x_.Get(),
+                     cudss_rhs_.Get()),
+        "cudssExecute with CUDSS_PHASE_ANALYSIS failed");
+
+    return LinearSolverTerminationType::SUCCESS;
+  }
+
+  // Resize buffers and setup cuDSS structs that describe the type and storage
+  // configuration of linear system operands.
+  LinearSolverTerminationType SetupCudssMatrices(
+      const CompressedRowSparseMatrix* lhs, std::string* message) {
+    const auto num_rows = lhs->num_rows();
+    const auto num_nonzeros = lhs->num_nonzeros();
+
+    if constexpr (std::is_same_v<Scalar, float>) {
+      lhs_values_h_.Reserve(num_nonzeros);
+      rhs_h_.Reserve(num_rows);
+      x_h_.Reserve(num_rows);
+    }
+
+    lhs_rows_d_.Reserve(num_rows + 1);
+    lhs_cols_d_.Reserve(num_nonzeros);
+    lhs_values_d_.Reserve(num_nonzeros);
+    rhs_d_.Reserve(num_rows);
+    x_d_.Reserve(num_rows);
+
+    static constexpr auto kFailedToCreateCuDSSMatrix =
+        "cudssMatrixCreate() call failed";
+    CUDSS_STATUS_OK_OR_RETURN_FATAL_ERROR(cudss_lhs_.Reset(num_rows,
+                                                           num_rows,
+                                                           num_nonzeros,
+                                                           lhs_rows_d_.data(),
+                                                           nullptr,
+                                                           lhs_cols_d_.data(),
+                                                           lhs_values_d_.data(),
+                                                           CUDA_R_32I,
+                                                           kCuDSSScalar,
+                                                           CUDSS_MTYPE_SPD,
+                                                           CUDSS_MVIEW_LOWER,
+                                                           CUDSS_BASE_ZERO),
+                                          kFailedToCreateCuDSSMatrix);
+
+    CUDSS_STATUS_OK_OR_RETURN_FATAL_ERROR(
+        cudss_rhs_.Reset(num_rows,
+                         1,
+                         num_rows,
+                         rhs_d_.data(),
+                         kCuDSSScalar,
+                         CUDSS_LAYOUT_COL_MAJOR),
+        kFailedToCreateCuDSSMatrix);
+
+    CUDSS_STATUS_OK_OR_RETURN_FATAL_ERROR(
+        cudss_x_.Reset(num_rows,
+                       1,
+                       num_rows,
+                       x_d_.data(),
+                       kCuDSSScalar,
+                       CUDSS_LAYOUT_COL_MAJOR),
+        kFailedToCreateCuDSSMatrix);
+
+    return LinearSolverTerminationType::SUCCESS;
+  }
+
+  template <typename S, typename D>
+  void Convert(const S* source, D* destination, size_t size) {
+    Eigen::Map<Eigen::Matrix<D, Eigen::Dynamic, 1>>(destination, size) =
+        Eigen::Map<const Eigen::Matrix<S, Eigen::Dynamic, 1>>(source, size)
+            .template cast<D>();
+  }
+
+  void ConvertAndCopyToDevice(const double* source,
+                              Scalar* intermediate,
+                              CudaBuffer<Scalar>& destination) {
+    const auto size = destination.size();
+    if constexpr (std::is_same_v<Scalar, double>) {
+      destination.CopyFromCpu(source, size);
+    } else {
+      Convert(source, intermediate, size);
+      destination.CopyFromCpu(intermediate, size);
+    }
+  }
+
+  void ConvertAndCopyToHost(const CudaBuffer<Scalar>& source,
+                            Scalar* intermediate,
+                            double* destination) {
+    const auto size = source.size();
+    if constexpr (std::is_same_v<Scalar, double>) {
+      source.CopyToCpu(destination, source.size());
+    } else {
+      source.CopyToCpu(intermediate, source.size());
+      Convert(intermediate, destination, size);
+    }
+  }
+
+  ContextImpl* context_{nullptr};
+  std::unique_ptr<CudssContext> cudss_context_;
+  CuDSSMatrixCSR cudss_lhs_;
+  CuDSSMatrixDense cudss_rhs_;
+  CuDSSMatrixDense cudss_x_;
+
+  CudaPinnedHostBuffer<Scalar> lhs_values_h_;
+  CudaPinnedHostBuffer<Scalar> rhs_h_;
+  CudaPinnedHostBuffer<Scalar> x_h_;
+  CudaBuffer<int> lhs_rows_d_;
+  CudaBuffer<int> lhs_cols_d_;
+  CudaBuffer<Scalar> lhs_values_d_;
+  CudaBuffer<Scalar> rhs_d_;
+  CudaBuffer<Scalar> x_d_;
+
+  LinearSolverTerminationType analyze_result_ =
+      LinearSolverTerminationType::FATAL_ERROR;
+  LinearSolverTerminationType factorize_result_ =
+      LinearSolverTerminationType::FATAL_ERROR;
+};
+
+template <typename Scalar>
+std::unique_ptr<SparseCholesky> CudaSparseCholesky<Scalar>::Create(
+    ContextImpl* context, const OrderingType ordering_type) {
+  if (ordering_type == OrderingType::NESDIS) {
+    LOG(FATAL)
+        << "Congratulations you have found a bug in Ceres Solver. Please "
+           "report it to the Ceres Solver developers.";
+    return nullptr;
+  }
+
+  if (context == nullptr || !context->IsCudaInitialized()) {
+    LOG(FATAL) << "CudaSparseCholesky requires CUDA context to be initialized";
+    return nullptr;
+  }
+
+  return std::make_unique<CudaSparseCholeskyImpl<Scalar>>(context);
+}
+
+template class CudaSparseCholesky<float>;
+template class CudaSparseCholesky<double>;
+
+}  // namespace ceres::internal
+
+#undef CUDSS_STATUS_CHECK
+#undef CUDSS_STATUS_OK_OR_RETURN_FATAL_ERROR
+#undef CUDSS_STATUS_OK_OR_RETURN_CUDSS_STATUS
+
+#endif  // CERES_NO_CUDSS
diff --git a/internal/ceres/cuda_sparse_cholesky.h b/internal/ceres/cuda_sparse_cholesky.h
new file mode 100644
index 0000000..e147633
--- /dev/null
+++ b/internal/ceres/cuda_sparse_cholesky.h
@@ -0,0 +1,66 @@
+// Ceres Solver - A fast non-linear least squares minimizer
+// Copyright 2024 Google Inc. All rights reserved.
+// http://ceres-solver.org/
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// * Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// * Neither the name of Google Inc. nor the names of its contributors may be
+//   used to endorse or promote products derived from this software without
+//   specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+// POSSIBILITY OF SUCH DAMAGE.
+//
+// Author: markshachkov@gmail.com (Mark Shachkov)
+//
+// A C++ interface to cuDSS.
+
+#ifndef CERES_INTERNAL_CUDA_SPARSE_CHOLESKY_H_
+#define CERES_INTERNAL_CUDA_SPARSE_CHOLESKY_H_
+
+// This include must come before any #ifndef check on Ceres compile options.
+#include "ceres/internal/config.h"
+
+#ifndef CERES_NO_CUDSS
+
+#include <memory>
+
+#include "ceres/internal/export.h"
+#include "ceres/linear_solver.h"
+#include "ceres/sparse_cholesky.h"
+
+namespace ceres::internal {
+
+// This class is a factory for implementation of sparse cholesky that uses cuDSS
+// on CUDA capable GPU's to solve sparse linear systems. Scalar controls the
+// precision used during computations, currently float and double are supported.
+// Details of implementation are incapsulated into cuda_sparse_cholesky.cc
+template <typename Scalar = double>
+class CERES_NO_EXPORT CudaSparseCholesky : public SparseCholesky {
+ public:
+  static constexpr bool IsNestedDissectionAvailable() noexcept { return false; }
+
+  static std::unique_ptr<SparseCholesky> Create(
+      ContextImpl* context, const OrderingType ordering_type);
+};
+
+}  // namespace ceres::internal
+
+#endif  // CERES_NO_CUDSS
+
+#endif  // CERES_INTERNAL_CUDA_SPARSE_CHOLESKY_H_
diff --git a/internal/ceres/dynamic_sparse_normal_cholesky_solver.cc b/internal/ceres/dynamic_sparse_normal_cholesky_solver.cc
index d77d7f7..6dd0adf 100644
--- a/internal/ceres/dynamic_sparse_normal_cholesky_solver.cc
+++ b/internal/ceres/dynamic_sparse_normal_cholesky_solver.cc
@@ -46,6 +46,7 @@
 #include "ceres/triplet_sparse_matrix.h"
 #include "ceres/types.h"
 #include "ceres/wall_time.h"
+#include "cuda_sparse_cholesky.h"
 
 #ifdef CERES_USE_EIGEN_SPARSE
 #include "Eigen/SparseCholesky"
@@ -88,6 +89,9 @@
     case EIGEN_SPARSE:
       summary = SolveImplUsingEigen(A, x);
       break;
+    case CUDA_SPARSE:
+      summary = SolveImplUsingCuda(A, x);
+      break;
     default:
       LOG(FATAL) << "Unsupported sparse linear algebra library for "
                  << "dynamic sparsity: "
@@ -231,4 +235,63 @@
 #endif
 }
 
+LinearSolver::Summary DynamicSparseNormalCholeskySolver::SolveImplUsingCuda(
+    CompressedRowSparseMatrix* A, double* rhs_and_solution) {
+#ifdef CERES_NO_CUDSS
+  (void)A;
+  (void)rhs_and_solution;
+
+  LinearSolver::Summary summary;
+  summary.num_iterations = 0;
+  summary.termination_type = LinearSolverTerminationType::FATAL_ERROR;
+  summary.message =
+      "SPARSE_NORMAL_CHOLESKY cannot be used with CUDA_SPARSE "
+      "because Ceres was not built with support for cuDSS. "
+      "This requires enabling building with -DUSE_CUDA=ON and ensuring that "
+      "cuDSS is found.";
+  return summary;
+#else
+
+  EventLogger event_logger("DynamicSparseNormalCholeskySolver::cuDSS::Solve");
+
+  // TODO: Consider computing A^T*A on device via cuSPARSE
+  // https://github.com/ceres-solver/ceres-solver/issues/1066
+  Eigen::Map<Eigen::SparseMatrix<double, Eigen::RowMajor>> a(
+      A->num_rows(),
+      A->num_cols(),
+      A->num_nonzeros(),
+      A->mutable_rows(),
+      A->mutable_cols(),
+      A->mutable_values());
+  Eigen::SparseMatrix<double, Eigen::RowMajor> ata =
+      (a.transpose() * a).triangularView<Eigen::Lower>();
+
+  CompressedRowSparseMatrix lhs(ata.rows(), ata.cols(), ata.nonZeros());
+  std::copy_n(ata.outerIndexPtr(), lhs.num_rows() + 1, lhs.mutable_rows());
+  std::copy_n(ata.innerIndexPtr(), lhs.num_nonzeros(), lhs.mutable_cols());
+  std::copy_n(ata.valuePtr(), lhs.num_nonzeros(), lhs.mutable_values());
+  lhs.set_storage_type(
+      CompressedRowSparseMatrix::StorageType::LOWER_TRIANGULAR);
+  event_logger.AddEvent("Compute A^T * A");
+
+  auto sparse_cholesky = CudaSparseCholesky<double>::Create(
+      options_.context, options_.ordering_type);
+
+  LinearSolver::Summary summary;
+  summary.num_iterations = 1;
+  summary.termination_type = sparse_cholesky->Factorize(&lhs, &summary.message);
+  if (summary.termination_type != LinearSolverTerminationType::SUCCESS) {
+    return summary;
+  }
+  event_logger.AddEvent("Analyze");
+
+  const Vector rhs = ConstVectorRef(rhs_and_solution, A->num_cols());
+  summary.termination_type =
+      sparse_cholesky->Solve(rhs.data(), rhs_and_solution, &summary.message);
+  event_logger.AddEvent("Solve");
+
+  return summary;
+#endif
+}
+
 }  // namespace ceres::internal
diff --git a/internal/ceres/dynamic_sparse_normal_cholesky_solver.h b/internal/ceres/dynamic_sparse_normal_cholesky_solver.h
index 394ba2a..e5e8a28 100644
--- a/internal/ceres/dynamic_sparse_normal_cholesky_solver.h
+++ b/internal/ceres/dynamic_sparse_normal_cholesky_solver.h
@@ -70,6 +70,9 @@
   LinearSolver::Summary SolveImplUsingEigen(CompressedRowSparseMatrix* A,
                                             double* rhs_and_solution);
 
+  LinearSolver::Summary SolveImplUsingCuda(CompressedRowSparseMatrix* A,
+                                            double* rhs_and_solution);
+
   const LinearSolver::Options options_;
 };
 
diff --git a/internal/ceres/dynamic_sparsity_test.cc b/internal/ceres/dynamic_sparsity_test.cc
index 0c29595..ae56037 100644
--- a/internal/ceres/dynamic_sparsity_test.cc
+++ b/internal/ceres/dynamic_sparsity_test.cc
@@ -426,38 +426,41 @@
 
   Solver::Options options;
   options.max_num_iterations = 100;
-  options.linear_solver_type = SPARSE_NORMAL_CHOLESKY;
-  // Only SuiteSparse & EigenSparse currently support dynamic sparsity.
-  options.sparse_linear_algebra_library_type =
+  SparseLinearAlgebraLibraryType sparse_library_types[] = {
 #if !defined(CERES_NO_SUITESPARSE)
-      ceres::SUITE_SPARSE;
-#elif defined(CERES_USE_EIGEN_SPARSE)
-      ceres::EIGEN_SPARSE;
+      ceres::SUITE_SPARSE,
 #endif
-
-  // First, solve `X` and `t` jointly with dynamic_sparsity = true.
-  Matrix X0 = X;
-  Vector t0 = t;
-  options.dynamic_sparsity = false;
-  Solver::Summary static_summary;
-  Solve(options, &problem, &static_summary);
-  EXPECT_EQ(static_summary.termination_type, CONVERGENCE)
-      << static_summary.FullReport();
-
-  X = X0;
-  t = t0;
-  options.dynamic_sparsity = true;
-  Solver::Summary dynamic_summary;
-  Solve(options, &problem, &dynamic_summary);
-  EXPECT_EQ(dynamic_summary.termination_type, CONVERGENCE)
-      << dynamic_summary.FullReport();
-
-  EXPECT_NEAR(static_summary.final_cost,
-              dynamic_summary.final_cost,
-              std::numeric_limits<double>::epsilon())
-      << "Static: \n"
-      << static_summary.FullReport() << "\nDynamic: \n"
-      << dynamic_summary.FullReport();
+#if defined(CERES_USE_EIGEN_SPARSE)
+      ceres::EIGEN_SPARSE,
+#endif
+#if !defined(CERES_NO_CUDSS)
+      ceres::CUDA_SPARSE,
+#endif
+  };
+  for (auto type : sparse_library_types) {
+    options.sparse_linear_algebra_library_type = type;
+    // First, solve `X` and `t` jointly with dynamic_sparsity = true.
+    Matrix X0 = X;
+    Vector t0 = t;
+    options.dynamic_sparsity = false;
+    Solver::Summary static_summary;
+    Solve(options, &problem, &static_summary);
+    EXPECT_EQ(static_summary.termination_type, CONVERGENCE)
+        << static_summary.FullReport();
+    X = X0;
+    t = t0;
+    options.dynamic_sparsity = true;
+    Solver::Summary dynamic_summary;
+    Solve(options, &problem, &dynamic_summary);
+    EXPECT_EQ(dynamic_summary.termination_type, CONVERGENCE)
+        << dynamic_summary.FullReport();
+    EXPECT_NEAR(static_summary.final_cost,
+                dynamic_summary.final_cost,
+                std::numeric_limits<double>::epsilon())
+        << "Static: \n"
+        << static_summary.FullReport() << "\nDynamic: \n"
+        << dynamic_summary.FullReport();
+  }
 }
 
 }  // namespace ceres::internal
diff --git a/internal/ceres/evaluator.cc b/internal/ceres/evaluator.cc
index 64eb4c5..9226dc5 100644
--- a/internal/ceres/evaluator.cc
+++ b/internal/ceres/evaluator.cc
@@ -63,7 +63,11 @@
           options, program);
     case DENSE_SCHUR:
     case SPARSE_SCHUR:
-    case ITERATIVE_SCHUR:
+    case ITERATIVE_SCHUR: {
+      return std::make_unique<
+            ProgramEvaluator<BlockEvaluatePreparer, BlockJacobianWriter>>(
+            options, program);
+    }
     case CGNR: {
       if (options.sparse_linear_algebra_library_type == CUDA_SPARSE) {
         return std::make_unique<ProgramEvaluator<ScratchEvaluatePreparer,
diff --git a/internal/ceres/generate_bundle_adjustment_tests.py b/internal/ceres/generate_bundle_adjustment_tests.py
index ac83bc3..0359b40 100644
--- a/internal/ceres/generate_bundle_adjustment_tests.py
+++ b/internal/ceres/generate_bundle_adjustment_tests.py
@@ -53,9 +53,11 @@
     ('SPARSE_NORMAL_CHOLESKY', 'SUITE_SPARSE'),
     ('SPARSE_NORMAL_CHOLESKY', 'EIGEN_SPARSE'),
     ('SPARSE_NORMAL_CHOLESKY', 'ACCELERATE_SPARSE'),
+    ('SPARSE_NORMAL_CHOLESKY', 'CUDA_SPARSE'),
     ('SPARSE_SCHUR',           'SUITE_SPARSE'),
     ('SPARSE_SCHUR',           'EIGEN_SPARSE'),
     ('SPARSE_SCHUR',           'ACCELERATE_SPARSE'),
+    ('SPARSE_SCHUR',           'CUDA_SPARSE')
 ]
 
 ITERATIVE_SOLVER_CONFIGS = [
@@ -66,9 +68,11 @@
     ('ITERATIVE_SCHUR',        'SUITE_SPARSE',     'CLUSTER_JACOBI'),
     ('ITERATIVE_SCHUR',        'EIGEN_SPARSE',     'CLUSTER_JACOBI'),
     ('ITERATIVE_SCHUR',        'ACCELERATE_SPARSE','CLUSTER_JACOBI'),
+    ('ITERATIVE_SCHUR',        'CUDA_SPARSE',      'CLUSTER_JACOBI'),
     ('ITERATIVE_SCHUR',        'SUITE_SPARSE',     'CLUSTER_TRIDIAGONAL'),
     ('ITERATIVE_SCHUR',        'EIGEN_SPARSE',     'CLUSTER_TRIDIAGONAL'),
     ('ITERATIVE_SCHUR',        'ACCELERATE_SPARSE','CLUSTER_TRIDIAGONAL'),
+    ('ITERATIVE_SCHUR',        'CUDA_SPARSE',      'CLUSTER_TRIDIAGONAL'),
 ]
 
 FILENAME_SHORTENING_MAP = dict(
@@ -83,6 +87,7 @@
   SUITE_SPARSE='suitesparse',
   EIGEN_SPARSE='eigensparse',
   ACCELERATE_SPARSE='acceleratesparse',
+  CUDA_SPARSE='cudasparse',
   IDENTITY='identity',
   JACOBI='jacobi',
   SCHUR_JACOBI='schurjacobi',
@@ -217,6 +222,14 @@
   elif sparse_backend == 'EIGEN_SPARSE':
     preprocessor_conditions_begin.append('#ifdef CERES_USE_EIGEN_SPARSE')
     preprocessor_conditions_end.insert(0, '#endif  // CERES_USE_EIGEN_SPARSE')
+  elif sparse_backend == 'CUDA_SPARSE':
+    preprocessor_conditions_begin.append('#ifndef CERES_NO_CUDA')
+    preprocessor_conditions_end.insert(0, '#endif  // CERES_NO_CUDA')
+    if linear_solver == 'SPARSE_SCHUR' or linear_solver == 'SPARSE_NORMAL_CHOLESKY' or (
+            linear_solver == 'ITERATIVE_SCHUR' and (
+            preconditioner == 'CLUSTER_JACOBI' or preconditioner == 'CLUSTER_TRIDIAGONAL')):
+      preprocessor_conditions_begin.append('#ifndef CERES_NO_CUDSS')
+      preprocessor_conditions_end.insert(0, '#endif  // CERES_NO_CUDSS')
 
   if dense_backend == "LAPACK":
     preprocessor_conditions_begin.append('#ifndef CERES_NO_LAPACK')
diff --git a/internal/ceres/generated_bundle_adjustment_tests/CMakeLists.txt b/internal/ceres/generated_bundle_adjustment_tests/CMakeLists.txt
index 5f4f65e..8da71b4 100644
--- a/internal/ceres/generated_bundle_adjustment_tests/CMakeLists.txt
+++ b/internal/ceres/generated_bundle_adjustment_tests/CMakeLists.txt
@@ -41,69 +41,85 @@
 ceres_test(ba_sparsecholesky_suitesparse_auto)
 ceres_test(ba_sparsecholesky_eigensparse_auto)
 ceres_test(ba_sparsecholesky_acceleratesparse_auto)
+ceres_test(ba_sparsecholesky_cudasparse_auto)
 ceres_test(ba_sparseschur_suitesparse_auto)
 ceres_test(ba_sparseschur_eigensparse_auto)
 ceres_test(ba_sparseschur_acceleratesparse_auto)
+ceres_test(ba_sparseschur_cudasparse_auto)
 ceres_test(ba_iterschur_jacobi_auto)
 ceres_test(ba_iterschur_schurjacobi_auto)
 ceres_test(ba_iterschur_spse_auto)
 ceres_test(ba_iterschur_suitesparse_clustjacobi_auto)
 ceres_test(ba_iterschur_eigensparse_clustjacobi_auto)
 ceres_test(ba_iterschur_acceleratesparse_clustjacobi_auto)
+ceres_test(ba_iterschur_cudasparse_clustjacobi_auto)
 ceres_test(ba_iterschur_suitesparse_clusttri_auto)
 ceres_test(ba_iterschur_eigensparse_clusttri_auto)
 ceres_test(ba_iterschur_acceleratesparse_clusttri_auto)
+ceres_test(ba_iterschur_cudasparse_clusttri_auto)
 ceres_test(ba_denseschur_eigen_auto_threads)
 ceres_test(ba_denseschur_lapack_auto_threads)
 ceres_test(ba_denseschur_cuda_auto_threads)
 ceres_test(ba_sparsecholesky_suitesparse_auto_threads)
 ceres_test(ba_sparsecholesky_eigensparse_auto_threads)
 ceres_test(ba_sparsecholesky_acceleratesparse_auto_threads)
+ceres_test(ba_sparsecholesky_cudasparse_auto_threads)
 ceres_test(ba_sparseschur_suitesparse_auto_threads)
 ceres_test(ba_sparseschur_eigensparse_auto_threads)
 ceres_test(ba_sparseschur_acceleratesparse_auto_threads)
+ceres_test(ba_sparseschur_cudasparse_auto_threads)
 ceres_test(ba_iterschur_jacobi_auto_threads)
 ceres_test(ba_iterschur_schurjacobi_auto_threads)
 ceres_test(ba_iterschur_spse_auto_threads)
 ceres_test(ba_iterschur_suitesparse_clustjacobi_auto_threads)
 ceres_test(ba_iterschur_eigensparse_clustjacobi_auto_threads)
 ceres_test(ba_iterschur_acceleratesparse_clustjacobi_auto_threads)
+ceres_test(ba_iterschur_cudasparse_clustjacobi_auto_threads)
 ceres_test(ba_iterschur_suitesparse_clusttri_auto_threads)
 ceres_test(ba_iterschur_eigensparse_clusttri_auto_threads)
 ceres_test(ba_iterschur_acceleratesparse_clusttri_auto_threads)
+ceres_test(ba_iterschur_cudasparse_clusttri_auto_threads)
 ceres_test(ba_denseschur_eigen_user)
 ceres_test(ba_denseschur_lapack_user)
 ceres_test(ba_denseschur_cuda_user)
 ceres_test(ba_sparsecholesky_suitesparse_user)
 ceres_test(ba_sparsecholesky_eigensparse_user)
 ceres_test(ba_sparsecholesky_acceleratesparse_user)
+ceres_test(ba_sparsecholesky_cudasparse_user)
 ceres_test(ba_sparseschur_suitesparse_user)
 ceres_test(ba_sparseschur_eigensparse_user)
 ceres_test(ba_sparseschur_acceleratesparse_user)
+ceres_test(ba_sparseschur_cudasparse_user)
 ceres_test(ba_iterschur_jacobi_user)
 ceres_test(ba_iterschur_schurjacobi_user)
 ceres_test(ba_iterschur_spse_user)
 ceres_test(ba_iterschur_suitesparse_clustjacobi_user)
 ceres_test(ba_iterschur_eigensparse_clustjacobi_user)
 ceres_test(ba_iterschur_acceleratesparse_clustjacobi_user)
+ceres_test(ba_iterschur_cudasparse_clustjacobi_user)
 ceres_test(ba_iterschur_suitesparse_clusttri_user)
 ceres_test(ba_iterschur_eigensparse_clusttri_user)
 ceres_test(ba_iterschur_acceleratesparse_clusttri_user)
+ceres_test(ba_iterschur_cudasparse_clusttri_user)
 ceres_test(ba_denseschur_eigen_user_threads)
 ceres_test(ba_denseschur_lapack_user_threads)
 ceres_test(ba_denseschur_cuda_user_threads)
 ceres_test(ba_sparsecholesky_suitesparse_user_threads)
 ceres_test(ba_sparsecholesky_eigensparse_user_threads)
 ceres_test(ba_sparsecholesky_acceleratesparse_user_threads)
+ceres_test(ba_sparsecholesky_cudasparse_user_threads)
 ceres_test(ba_sparseschur_suitesparse_user_threads)
 ceres_test(ba_sparseschur_eigensparse_user_threads)
 ceres_test(ba_sparseschur_acceleratesparse_user_threads)
+ceres_test(ba_sparseschur_cudasparse_user_threads)
 ceres_test(ba_iterschur_jacobi_user_threads)
 ceres_test(ba_iterschur_schurjacobi_user_threads)
 ceres_test(ba_iterschur_spse_user_threads)
 ceres_test(ba_iterschur_suitesparse_clustjacobi_user_threads)
 ceres_test(ba_iterschur_eigensparse_clustjacobi_user_threads)
 ceres_test(ba_iterschur_acceleratesparse_clustjacobi_user_threads)
+ceres_test(ba_iterschur_cudasparse_clustjacobi_user_threads)
 ceres_test(ba_iterschur_suitesparse_clusttri_user_threads)
 ceres_test(ba_iterschur_eigensparse_clusttri_user_threads)
 ceres_test(ba_iterschur_acceleratesparse_clusttri_user_threads)
+ceres_test(ba_iterschur_cudasparse_clusttri_user_threads)
diff --git a/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clustjacobi_auto_test.cc b/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clustjacobi_auto_test.cc
new file mode 100644
index 0000000..1719aee
--- /dev/null
+++ b/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clustjacobi_auto_test.cc
@@ -0,0 +1,67 @@
+// Ceres Solver - A fast non-linear least squares minimizer
+// Copyright 2023 Google Inc. All rights reserved.
+// http://ceres-solver.org/
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// * Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// * Neither the name of Google Inc. nor the names of its contributors may be
+//   used to endorse or promote products derived from this software without
+//   specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+// POSSIBILITY OF SUCH DAMAGE.
+//
+// ========================================
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// ========================================
+//
+// This file is generated using generate_bundle_adjustment_tests.py.
+
+#include "ceres/bundle_adjustment_test_util.h"
+#include "ceres/internal/config.h"
+#include "gtest/gtest.h"
+
+#ifndef CERES_NO_CUDA
+#ifndef CERES_NO_CUDSS
+
+namespace ceres::internal {
+
+TEST_F(BundleAdjustmentTest,
+       IterativeSchur_CudaSparse_ClusterJacobi_AutomaticOrdering) {  // NOLINT
+  BundleAdjustmentProblem bundle_adjustment_problem;
+  Solver::Options* options = bundle_adjustment_problem.mutable_solver_options();
+  options->eta = 0.01;
+  options->num_threads = 1;
+  options->linear_solver_type = ITERATIVE_SCHUR;
+  options->dense_linear_algebra_library_type = EIGEN;
+  options->sparse_linear_algebra_library_type = CUDA_SPARSE;
+  options->preconditioner_type = CLUSTER_JACOBI;
+  if (kAutomaticOrdering) {
+    options->linear_solver_ordering = nullptr;
+  }
+  Problem* problem = bundle_adjustment_problem.mutable_problem();
+  RunSolverForConfigAndExpectResidualsMatch(*options, problem);
+}
+
+}  // namespace ceres::internal
+
+#endif  // CERES_NO_CUDSS
+#endif  // CERES_NO_CUDA
diff --git a/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clustjacobi_auto_threads_test.cc b/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clustjacobi_auto_threads_test.cc
new file mode 100644
index 0000000..c0c3489
--- /dev/null
+++ b/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clustjacobi_auto_threads_test.cc
@@ -0,0 +1,67 @@
+// Ceres Solver - A fast non-linear least squares minimizer
+// Copyright 2023 Google Inc. All rights reserved.
+// http://ceres-solver.org/
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// * Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// * Neither the name of Google Inc. nor the names of its contributors may be
+//   used to endorse or promote products derived from this software without
+//   specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+// POSSIBILITY OF SUCH DAMAGE.
+//
+// ========================================
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// ========================================
+//
+// This file is generated using generate_bundle_adjustment_tests.py.
+
+#include "ceres/bundle_adjustment_test_util.h"
+#include "ceres/internal/config.h"
+#include "gtest/gtest.h"
+
+#ifndef CERES_NO_CUDA
+#ifndef CERES_NO_CUDSS
+
+namespace ceres::internal {
+
+TEST_F(BundleAdjustmentTest,
+       IterativeSchur_CudaSparse_ClusterJacobi_AutomaticOrdering_Threads) {  // NOLINT
+  BundleAdjustmentProblem bundle_adjustment_problem;
+  Solver::Options* options = bundle_adjustment_problem.mutable_solver_options();
+  options->eta = 0.01;
+  options->num_threads = 4;
+  options->linear_solver_type = ITERATIVE_SCHUR;
+  options->dense_linear_algebra_library_type = EIGEN;
+  options->sparse_linear_algebra_library_type = CUDA_SPARSE;
+  options->preconditioner_type = CLUSTER_JACOBI;
+  if (kAutomaticOrdering) {
+    options->linear_solver_ordering = nullptr;
+  }
+  Problem* problem = bundle_adjustment_problem.mutable_problem();
+  RunSolverForConfigAndExpectResidualsMatch(*options, problem);
+}
+
+}  // namespace ceres::internal
+
+#endif  // CERES_NO_CUDSS
+#endif  // CERES_NO_CUDA
diff --git a/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clustjacobi_user_test.cc b/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clustjacobi_user_test.cc
new file mode 100644
index 0000000..554a432
--- /dev/null
+++ b/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clustjacobi_user_test.cc
@@ -0,0 +1,67 @@
+// Ceres Solver - A fast non-linear least squares minimizer
+// Copyright 2023 Google Inc. All rights reserved.
+// http://ceres-solver.org/
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// * Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// * Neither the name of Google Inc. nor the names of its contributors may be
+//   used to endorse or promote products derived from this software without
+//   specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+// POSSIBILITY OF SUCH DAMAGE.
+//
+// ========================================
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// ========================================
+//
+// This file is generated using generate_bundle_adjustment_tests.py.
+
+#include "ceres/bundle_adjustment_test_util.h"
+#include "ceres/internal/config.h"
+#include "gtest/gtest.h"
+
+#ifndef CERES_NO_CUDA
+#ifndef CERES_NO_CUDSS
+
+namespace ceres::internal {
+
+TEST_F(BundleAdjustmentTest,
+       IterativeSchur_CudaSparse_ClusterJacobi_UserOrdering) {  // NOLINT
+  BundleAdjustmentProblem bundle_adjustment_problem;
+  Solver::Options* options = bundle_adjustment_problem.mutable_solver_options();
+  options->eta = 0.01;
+  options->num_threads = 1;
+  options->linear_solver_type = ITERATIVE_SCHUR;
+  options->dense_linear_algebra_library_type = EIGEN;
+  options->sparse_linear_algebra_library_type = CUDA_SPARSE;
+  options->preconditioner_type = CLUSTER_JACOBI;
+  if (kUserOrdering) {
+    options->linear_solver_ordering = nullptr;
+  }
+  Problem* problem = bundle_adjustment_problem.mutable_problem();
+  RunSolverForConfigAndExpectResidualsMatch(*options, problem);
+}
+
+}  // namespace ceres::internal
+
+#endif  // CERES_NO_CUDSS
+#endif  // CERES_NO_CUDA
diff --git a/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clustjacobi_user_threads_test.cc b/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clustjacobi_user_threads_test.cc
new file mode 100644
index 0000000..1e39f65
--- /dev/null
+++ b/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clustjacobi_user_threads_test.cc
@@ -0,0 +1,67 @@
+// Ceres Solver - A fast non-linear least squares minimizer
+// Copyright 2023 Google Inc. All rights reserved.
+// http://ceres-solver.org/
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// * Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// * Neither the name of Google Inc. nor the names of its contributors may be
+//   used to endorse or promote products derived from this software without
+//   specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+// POSSIBILITY OF SUCH DAMAGE.
+//
+// ========================================
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// ========================================
+//
+// This file is generated using generate_bundle_adjustment_tests.py.
+
+#include "ceres/bundle_adjustment_test_util.h"
+#include "ceres/internal/config.h"
+#include "gtest/gtest.h"
+
+#ifndef CERES_NO_CUDA
+#ifndef CERES_NO_CUDSS
+
+namespace ceres::internal {
+
+TEST_F(BundleAdjustmentTest,
+       IterativeSchur_CudaSparse_ClusterJacobi_UserOrdering_Threads) {  // NOLINT
+  BundleAdjustmentProblem bundle_adjustment_problem;
+  Solver::Options* options = bundle_adjustment_problem.mutable_solver_options();
+  options->eta = 0.01;
+  options->num_threads = 4;
+  options->linear_solver_type = ITERATIVE_SCHUR;
+  options->dense_linear_algebra_library_type = EIGEN;
+  options->sparse_linear_algebra_library_type = CUDA_SPARSE;
+  options->preconditioner_type = CLUSTER_JACOBI;
+  if (kUserOrdering) {
+    options->linear_solver_ordering = nullptr;
+  }
+  Problem* problem = bundle_adjustment_problem.mutable_problem();
+  RunSolverForConfigAndExpectResidualsMatch(*options, problem);
+}
+
+}  // namespace ceres::internal
+
+#endif  // CERES_NO_CUDSS
+#endif  // CERES_NO_CUDA
diff --git a/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clusttri_auto_test.cc b/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clusttri_auto_test.cc
new file mode 100644
index 0000000..6addad5
--- /dev/null
+++ b/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clusttri_auto_test.cc
@@ -0,0 +1,67 @@
+// Ceres Solver - A fast non-linear least squares minimizer
+// Copyright 2023 Google Inc. All rights reserved.
+// http://ceres-solver.org/
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// * Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// * Neither the name of Google Inc. nor the names of its contributors may be
+//   used to endorse or promote products derived from this software without
+//   specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+// POSSIBILITY OF SUCH DAMAGE.
+//
+// ========================================
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// ========================================
+//
+// This file is generated using generate_bundle_adjustment_tests.py.
+
+#include "ceres/bundle_adjustment_test_util.h"
+#include "ceres/internal/config.h"
+#include "gtest/gtest.h"
+
+#ifndef CERES_NO_CUDA
+#ifndef CERES_NO_CUDSS
+
+namespace ceres::internal {
+
+TEST_F(BundleAdjustmentTest,
+       IterativeSchur_CudaSparse_ClusterTridiagonal_AutomaticOrdering) {  // NOLINT
+  BundleAdjustmentProblem bundle_adjustment_problem;
+  Solver::Options* options = bundle_adjustment_problem.mutable_solver_options();
+  options->eta = 0.01;
+  options->num_threads = 1;
+  options->linear_solver_type = ITERATIVE_SCHUR;
+  options->dense_linear_algebra_library_type = EIGEN;
+  options->sparse_linear_algebra_library_type = CUDA_SPARSE;
+  options->preconditioner_type = CLUSTER_TRIDIAGONAL;
+  if (kAutomaticOrdering) {
+    options->linear_solver_ordering = nullptr;
+  }
+  Problem* problem = bundle_adjustment_problem.mutable_problem();
+  RunSolverForConfigAndExpectResidualsMatch(*options, problem);
+}
+
+}  // namespace ceres::internal
+
+#endif  // CERES_NO_CUDSS
+#endif  // CERES_NO_CUDA
diff --git a/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clusttri_auto_threads_test.cc b/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clusttri_auto_threads_test.cc
new file mode 100644
index 0000000..d2fefb8
--- /dev/null
+++ b/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clusttri_auto_threads_test.cc
@@ -0,0 +1,67 @@
+// Ceres Solver - A fast non-linear least squares minimizer
+// Copyright 2023 Google Inc. All rights reserved.
+// http://ceres-solver.org/
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// * Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// * Neither the name of Google Inc. nor the names of its contributors may be
+//   used to endorse or promote products derived from this software without
+//   specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+// POSSIBILITY OF SUCH DAMAGE.
+//
+// ========================================
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// ========================================
+//
+// This file is generated using generate_bundle_adjustment_tests.py.
+
+#include "ceres/bundle_adjustment_test_util.h"
+#include "ceres/internal/config.h"
+#include "gtest/gtest.h"
+
+#ifndef CERES_NO_CUDA
+#ifndef CERES_NO_CUDSS
+
+namespace ceres::internal {
+
+TEST_F(BundleAdjustmentTest,
+       IterativeSchur_CudaSparse_ClusterTridiagonal_AutomaticOrdering_Threads) {  // NOLINT
+  BundleAdjustmentProblem bundle_adjustment_problem;
+  Solver::Options* options = bundle_adjustment_problem.mutable_solver_options();
+  options->eta = 0.01;
+  options->num_threads = 4;
+  options->linear_solver_type = ITERATIVE_SCHUR;
+  options->dense_linear_algebra_library_type = EIGEN;
+  options->sparse_linear_algebra_library_type = CUDA_SPARSE;
+  options->preconditioner_type = CLUSTER_TRIDIAGONAL;
+  if (kAutomaticOrdering) {
+    options->linear_solver_ordering = nullptr;
+  }
+  Problem* problem = bundle_adjustment_problem.mutable_problem();
+  RunSolverForConfigAndExpectResidualsMatch(*options, problem);
+}
+
+}  // namespace ceres::internal
+
+#endif  // CERES_NO_CUDSS
+#endif  // CERES_NO_CUDA
diff --git a/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clusttri_user_test.cc b/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clusttri_user_test.cc
new file mode 100644
index 0000000..ae04eec
--- /dev/null
+++ b/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clusttri_user_test.cc
@@ -0,0 +1,67 @@
+// Ceres Solver - A fast non-linear least squares minimizer
+// Copyright 2023 Google Inc. All rights reserved.
+// http://ceres-solver.org/
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// * Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// * Neither the name of Google Inc. nor the names of its contributors may be
+//   used to endorse or promote products derived from this software without
+//   specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+// POSSIBILITY OF SUCH DAMAGE.
+//
+// ========================================
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// ========================================
+//
+// This file is generated using generate_bundle_adjustment_tests.py.
+
+#include "ceres/bundle_adjustment_test_util.h"
+#include "ceres/internal/config.h"
+#include "gtest/gtest.h"
+
+#ifndef CERES_NO_CUDA
+#ifndef CERES_NO_CUDSS
+
+namespace ceres::internal {
+
+TEST_F(BundleAdjustmentTest,
+       IterativeSchur_CudaSparse_ClusterTridiagonal_UserOrdering) {  // NOLINT
+  BundleAdjustmentProblem bundle_adjustment_problem;
+  Solver::Options* options = bundle_adjustment_problem.mutable_solver_options();
+  options->eta = 0.01;
+  options->num_threads = 1;
+  options->linear_solver_type = ITERATIVE_SCHUR;
+  options->dense_linear_algebra_library_type = EIGEN;
+  options->sparse_linear_algebra_library_type = CUDA_SPARSE;
+  options->preconditioner_type = CLUSTER_TRIDIAGONAL;
+  if (kUserOrdering) {
+    options->linear_solver_ordering = nullptr;
+  }
+  Problem* problem = bundle_adjustment_problem.mutable_problem();
+  RunSolverForConfigAndExpectResidualsMatch(*options, problem);
+}
+
+}  // namespace ceres::internal
+
+#endif  // CERES_NO_CUDSS
+#endif  // CERES_NO_CUDA
diff --git a/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clusttri_user_threads_test.cc b/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clusttri_user_threads_test.cc
new file mode 100644
index 0000000..cad2051
--- /dev/null
+++ b/internal/ceres/generated_bundle_adjustment_tests/ba_iterschur_cudasparse_clusttri_user_threads_test.cc
@@ -0,0 +1,67 @@
+// Ceres Solver - A fast non-linear least squares minimizer
+// Copyright 2023 Google Inc. All rights reserved.
+// http://ceres-solver.org/
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// * Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// * Neither the name of Google Inc. nor the names of its contributors may be
+//   used to endorse or promote products derived from this software without
+//   specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+// POSSIBILITY OF SUCH DAMAGE.
+//
+// ========================================
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// ========================================
+//
+// This file is generated using generate_bundle_adjustment_tests.py.
+
+#include "ceres/bundle_adjustment_test_util.h"
+#include "ceres/internal/config.h"
+#include "gtest/gtest.h"
+
+#ifndef CERES_NO_CUDA
+#ifndef CERES_NO_CUDSS
+
+namespace ceres::internal {
+
+TEST_F(BundleAdjustmentTest,
+       IterativeSchur_CudaSparse_ClusterTridiagonal_UserOrdering_Threads) {  // NOLINT
+  BundleAdjustmentProblem bundle_adjustment_problem;
+  Solver::Options* options = bundle_adjustment_problem.mutable_solver_options();
+  options->eta = 0.01;
+  options->num_threads = 4;
+  options->linear_solver_type = ITERATIVE_SCHUR;
+  options->dense_linear_algebra_library_type = EIGEN;
+  options->sparse_linear_algebra_library_type = CUDA_SPARSE;
+  options->preconditioner_type = CLUSTER_TRIDIAGONAL;
+  if (kUserOrdering) {
+    options->linear_solver_ordering = nullptr;
+  }
+  Problem* problem = bundle_adjustment_problem.mutable_problem();
+  RunSolverForConfigAndExpectResidualsMatch(*options, problem);
+}
+
+}  // namespace ceres::internal
+
+#endif  // CERES_NO_CUDSS
+#endif  // CERES_NO_CUDA
diff --git a/internal/ceres/generated_bundle_adjustment_tests/ba_sparsecholesky_cudasparse_auto_test.cc b/internal/ceres/generated_bundle_adjustment_tests/ba_sparsecholesky_cudasparse_auto_test.cc
new file mode 100644
index 0000000..77f2fc1
--- /dev/null
+++ b/internal/ceres/generated_bundle_adjustment_tests/ba_sparsecholesky_cudasparse_auto_test.cc
@@ -0,0 +1,67 @@
+// Ceres Solver - A fast non-linear least squares minimizer
+// Copyright 2023 Google Inc. All rights reserved.
+// http://ceres-solver.org/
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// * Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// * Neither the name of Google Inc. nor the names of its contributors may be
+//   used to endorse or promote products derived from this software without
+//   specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+// POSSIBILITY OF SUCH DAMAGE.
+//
+// ========================================
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// ========================================
+//
+// This file is generated using generate_bundle_adjustment_tests.py.
+
+#include "ceres/bundle_adjustment_test_util.h"
+#include "ceres/internal/config.h"
+#include "gtest/gtest.h"
+
+#ifndef CERES_NO_CUDA
+#ifndef CERES_NO_CUDSS
+
+namespace ceres::internal {
+
+TEST_F(BundleAdjustmentTest,
+       SparseNormalCholesky_CudaSparse_AutomaticOrdering) {  // NOLINT
+  BundleAdjustmentProblem bundle_adjustment_problem;
+  Solver::Options* options = bundle_adjustment_problem.mutable_solver_options();
+  options->eta = 0.01;
+  options->num_threads = 1;
+  options->linear_solver_type = SPARSE_NORMAL_CHOLESKY;
+  options->dense_linear_algebra_library_type = EIGEN;
+  options->sparse_linear_algebra_library_type = CUDA_SPARSE;
+  options->preconditioner_type = IDENTITY;
+  if (kAutomaticOrdering) {
+    options->linear_solver_ordering = nullptr;
+  }
+  Problem* problem = bundle_adjustment_problem.mutable_problem();
+  RunSolverForConfigAndExpectResidualsMatch(*options, problem);
+}
+
+}  // namespace ceres::internal
+
+#endif  // CERES_NO_CUDSS
+#endif  // CERES_NO_CUDA
diff --git a/internal/ceres/generated_bundle_adjustment_tests/ba_sparsecholesky_cudasparse_auto_threads_test.cc b/internal/ceres/generated_bundle_adjustment_tests/ba_sparsecholesky_cudasparse_auto_threads_test.cc
new file mode 100644
index 0000000..f436544
--- /dev/null
+++ b/internal/ceres/generated_bundle_adjustment_tests/ba_sparsecholesky_cudasparse_auto_threads_test.cc
@@ -0,0 +1,67 @@
+// Ceres Solver - A fast non-linear least squares minimizer
+// Copyright 2023 Google Inc. All rights reserved.
+// http://ceres-solver.org/
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// * Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// * Neither the name of Google Inc. nor the names of its contributors may be
+//   used to endorse or promote products derived from this software without
+//   specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+// POSSIBILITY OF SUCH DAMAGE.
+//
+// ========================================
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// ========================================
+//
+// This file is generated using generate_bundle_adjustment_tests.py.
+
+#include "ceres/bundle_adjustment_test_util.h"
+#include "ceres/internal/config.h"
+#include "gtest/gtest.h"
+
+#ifndef CERES_NO_CUDA
+#ifndef CERES_NO_CUDSS
+
+namespace ceres::internal {
+
+TEST_F(BundleAdjustmentTest,
+       SparseNormalCholesky_CudaSparse_AutomaticOrdering_Threads) {  // NOLINT
+  BundleAdjustmentProblem bundle_adjustment_problem;
+  Solver::Options* options = bundle_adjustment_problem.mutable_solver_options();
+  options->eta = 0.01;
+  options->num_threads = 4;
+  options->linear_solver_type = SPARSE_NORMAL_CHOLESKY;
+  options->dense_linear_algebra_library_type = EIGEN;
+  options->sparse_linear_algebra_library_type = CUDA_SPARSE;
+  options->preconditioner_type = IDENTITY;
+  if (kAutomaticOrdering) {
+    options->linear_solver_ordering = nullptr;
+  }
+  Problem* problem = bundle_adjustment_problem.mutable_problem();
+  RunSolverForConfigAndExpectResidualsMatch(*options, problem);
+}
+
+}  // namespace ceres::internal
+
+#endif  // CERES_NO_CUDSS
+#endif  // CERES_NO_CUDA
diff --git a/internal/ceres/generated_bundle_adjustment_tests/ba_sparsecholesky_cudasparse_user_test.cc b/internal/ceres/generated_bundle_adjustment_tests/ba_sparsecholesky_cudasparse_user_test.cc
new file mode 100644
index 0000000..dea6202
--- /dev/null
+++ b/internal/ceres/generated_bundle_adjustment_tests/ba_sparsecholesky_cudasparse_user_test.cc
@@ -0,0 +1,67 @@
+// Ceres Solver - A fast non-linear least squares minimizer
+// Copyright 2023 Google Inc. All rights reserved.
+// http://ceres-solver.org/
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// * Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// * Neither the name of Google Inc. nor the names of its contributors may be
+//   used to endorse or promote products derived from this software without
+//   specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+// POSSIBILITY OF SUCH DAMAGE.
+//
+// ========================================
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// ========================================
+//
+// This file is generated using generate_bundle_adjustment_tests.py.
+
+#include "ceres/bundle_adjustment_test_util.h"
+#include "ceres/internal/config.h"
+#include "gtest/gtest.h"
+
+#ifndef CERES_NO_CUDA
+#ifndef CERES_NO_CUDSS
+
+namespace ceres::internal {
+
+TEST_F(BundleAdjustmentTest,
+       SparseNormalCholesky_CudaSparse_UserOrdering) {  // NOLINT
+  BundleAdjustmentProblem bundle_adjustment_problem;
+  Solver::Options* options = bundle_adjustment_problem.mutable_solver_options();
+  options->eta = 0.01;
+  options->num_threads = 1;
+  options->linear_solver_type = SPARSE_NORMAL_CHOLESKY;
+  options->dense_linear_algebra_library_type = EIGEN;
+  options->sparse_linear_algebra_library_type = CUDA_SPARSE;
+  options->preconditioner_type = IDENTITY;
+  if (kUserOrdering) {
+    options->linear_solver_ordering = nullptr;
+  }
+  Problem* problem = bundle_adjustment_problem.mutable_problem();
+  RunSolverForConfigAndExpectResidualsMatch(*options, problem);
+}
+
+}  // namespace ceres::internal
+
+#endif  // CERES_NO_CUDSS
+#endif  // CERES_NO_CUDA
diff --git a/internal/ceres/generated_bundle_adjustment_tests/ba_sparsecholesky_cudasparse_user_threads_test.cc b/internal/ceres/generated_bundle_adjustment_tests/ba_sparsecholesky_cudasparse_user_threads_test.cc
new file mode 100644
index 0000000..90a69e8
--- /dev/null
+++ b/internal/ceres/generated_bundle_adjustment_tests/ba_sparsecholesky_cudasparse_user_threads_test.cc
@@ -0,0 +1,67 @@
+// Ceres Solver - A fast non-linear least squares minimizer
+// Copyright 2023 Google Inc. All rights reserved.
+// http://ceres-solver.org/
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// * Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// * Neither the name of Google Inc. nor the names of its contributors may be
+//   used to endorse or promote products derived from this software without
+//   specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+// POSSIBILITY OF SUCH DAMAGE.
+//
+// ========================================
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// ========================================
+//
+// This file is generated using generate_bundle_adjustment_tests.py.
+
+#include "ceres/bundle_adjustment_test_util.h"
+#include "ceres/internal/config.h"
+#include "gtest/gtest.h"
+
+#ifndef CERES_NO_CUDA
+#ifndef CERES_NO_CUDSS
+
+namespace ceres::internal {
+
+TEST_F(BundleAdjustmentTest,
+       SparseNormalCholesky_CudaSparse_UserOrdering_Threads) {  // NOLINT
+  BundleAdjustmentProblem bundle_adjustment_problem;
+  Solver::Options* options = bundle_adjustment_problem.mutable_solver_options();
+  options->eta = 0.01;
+  options->num_threads = 4;
+  options->linear_solver_type = SPARSE_NORMAL_CHOLESKY;
+  options->dense_linear_algebra_library_type = EIGEN;
+  options->sparse_linear_algebra_library_type = CUDA_SPARSE;
+  options->preconditioner_type = IDENTITY;
+  if (kUserOrdering) {
+    options->linear_solver_ordering = nullptr;
+  }
+  Problem* problem = bundle_adjustment_problem.mutable_problem();
+  RunSolverForConfigAndExpectResidualsMatch(*options, problem);
+}
+
+}  // namespace ceres::internal
+
+#endif  // CERES_NO_CUDSS
+#endif  // CERES_NO_CUDA
diff --git a/internal/ceres/generated_bundle_adjustment_tests/ba_sparseschur_cudasparse_auto_test.cc b/internal/ceres/generated_bundle_adjustment_tests/ba_sparseschur_cudasparse_auto_test.cc
new file mode 100644
index 0000000..5c45977
--- /dev/null
+++ b/internal/ceres/generated_bundle_adjustment_tests/ba_sparseschur_cudasparse_auto_test.cc
@@ -0,0 +1,67 @@
+// Ceres Solver - A fast non-linear least squares minimizer
+// Copyright 2023 Google Inc. All rights reserved.
+// http://ceres-solver.org/
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// * Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// * Neither the name of Google Inc. nor the names of its contributors may be
+//   used to endorse or promote products derived from this software without
+//   specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+// POSSIBILITY OF SUCH DAMAGE.
+//
+// ========================================
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// ========================================
+//
+// This file is generated using generate_bundle_adjustment_tests.py.
+
+#include "ceres/bundle_adjustment_test_util.h"
+#include "ceres/internal/config.h"
+#include "gtest/gtest.h"
+
+#ifndef CERES_NO_CUDA
+#ifndef CERES_NO_CUDSS
+
+namespace ceres::internal {
+
+TEST_F(BundleAdjustmentTest,
+       SparseSchur_CudaSparse_AutomaticOrdering) {  // NOLINT
+  BundleAdjustmentProblem bundle_adjustment_problem;
+  Solver::Options* options = bundle_adjustment_problem.mutable_solver_options();
+  options->eta = 0.01;
+  options->num_threads = 1;
+  options->linear_solver_type = SPARSE_SCHUR;
+  options->dense_linear_algebra_library_type = EIGEN;
+  options->sparse_linear_algebra_library_type = CUDA_SPARSE;
+  options->preconditioner_type = IDENTITY;
+  if (kAutomaticOrdering) {
+    options->linear_solver_ordering = nullptr;
+  }
+  Problem* problem = bundle_adjustment_problem.mutable_problem();
+  RunSolverForConfigAndExpectResidualsMatch(*options, problem);
+}
+
+}  // namespace ceres::internal
+
+#endif  // CERES_NO_CUDSS
+#endif  // CERES_NO_CUDA
diff --git a/internal/ceres/generated_bundle_adjustment_tests/ba_sparseschur_cudasparse_auto_threads_test.cc b/internal/ceres/generated_bundle_adjustment_tests/ba_sparseschur_cudasparse_auto_threads_test.cc
new file mode 100644
index 0000000..b866acb
--- /dev/null
+++ b/internal/ceres/generated_bundle_adjustment_tests/ba_sparseschur_cudasparse_auto_threads_test.cc
@@ -0,0 +1,67 @@
+// Ceres Solver - A fast non-linear least squares minimizer
+// Copyright 2023 Google Inc. All rights reserved.
+// http://ceres-solver.org/
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// * Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// * Neither the name of Google Inc. nor the names of its contributors may be
+//   used to endorse or promote products derived from this software without
+//   specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+// POSSIBILITY OF SUCH DAMAGE.
+//
+// ========================================
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// ========================================
+//
+// This file is generated using generate_bundle_adjustment_tests.py.
+
+#include "ceres/bundle_adjustment_test_util.h"
+#include "ceres/internal/config.h"
+#include "gtest/gtest.h"
+
+#ifndef CERES_NO_CUDA
+#ifndef CERES_NO_CUDSS
+
+namespace ceres::internal {
+
+TEST_F(BundleAdjustmentTest,
+       SparseSchur_CudaSparse_AutomaticOrdering_Threads) {  // NOLINT
+  BundleAdjustmentProblem bundle_adjustment_problem;
+  Solver::Options* options = bundle_adjustment_problem.mutable_solver_options();
+  options->eta = 0.01;
+  options->num_threads = 4;
+  options->linear_solver_type = SPARSE_SCHUR;
+  options->dense_linear_algebra_library_type = EIGEN;
+  options->sparse_linear_algebra_library_type = CUDA_SPARSE;
+  options->preconditioner_type = IDENTITY;
+  if (kAutomaticOrdering) {
+    options->linear_solver_ordering = nullptr;
+  }
+  Problem* problem = bundle_adjustment_problem.mutable_problem();
+  RunSolverForConfigAndExpectResidualsMatch(*options, problem);
+}
+
+}  // namespace ceres::internal
+
+#endif  // CERES_NO_CUDSS
+#endif  // CERES_NO_CUDA
diff --git a/internal/ceres/generated_bundle_adjustment_tests/ba_sparseschur_cudasparse_user_test.cc b/internal/ceres/generated_bundle_adjustment_tests/ba_sparseschur_cudasparse_user_test.cc
new file mode 100644
index 0000000..a7e5e14
--- /dev/null
+++ b/internal/ceres/generated_bundle_adjustment_tests/ba_sparseschur_cudasparse_user_test.cc
@@ -0,0 +1,67 @@
+// Ceres Solver - A fast non-linear least squares minimizer
+// Copyright 2023 Google Inc. All rights reserved.
+// http://ceres-solver.org/
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// * Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// * Neither the name of Google Inc. nor the names of its contributors may be
+//   used to endorse or promote products derived from this software without
+//   specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+// POSSIBILITY OF SUCH DAMAGE.
+//
+// ========================================
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// ========================================
+//
+// This file is generated using generate_bundle_adjustment_tests.py.
+
+#include "ceres/bundle_adjustment_test_util.h"
+#include "ceres/internal/config.h"
+#include "gtest/gtest.h"
+
+#ifndef CERES_NO_CUDA
+#ifndef CERES_NO_CUDSS
+
+namespace ceres::internal {
+
+TEST_F(BundleAdjustmentTest,
+       SparseSchur_CudaSparse_UserOrdering) {  // NOLINT
+  BundleAdjustmentProblem bundle_adjustment_problem;
+  Solver::Options* options = bundle_adjustment_problem.mutable_solver_options();
+  options->eta = 0.01;
+  options->num_threads = 1;
+  options->linear_solver_type = SPARSE_SCHUR;
+  options->dense_linear_algebra_library_type = EIGEN;
+  options->sparse_linear_algebra_library_type = CUDA_SPARSE;
+  options->preconditioner_type = IDENTITY;
+  if (kUserOrdering) {
+    options->linear_solver_ordering = nullptr;
+  }
+  Problem* problem = bundle_adjustment_problem.mutable_problem();
+  RunSolverForConfigAndExpectResidualsMatch(*options, problem);
+}
+
+}  // namespace ceres::internal
+
+#endif  // CERES_NO_CUDSS
+#endif  // CERES_NO_CUDA
diff --git a/internal/ceres/generated_bundle_adjustment_tests/ba_sparseschur_cudasparse_user_threads_test.cc b/internal/ceres/generated_bundle_adjustment_tests/ba_sparseschur_cudasparse_user_threads_test.cc
new file mode 100644
index 0000000..7319b28
--- /dev/null
+++ b/internal/ceres/generated_bundle_adjustment_tests/ba_sparseschur_cudasparse_user_threads_test.cc
@@ -0,0 +1,67 @@
+// Ceres Solver - A fast non-linear least squares minimizer
+// Copyright 2023 Google Inc. All rights reserved.
+// http://ceres-solver.org/
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// * Redistributions of source code must retain the above copyright notice,
+//   this list of conditions and the following disclaimer.
+// * Redistributions in binary form must reproduce the above copyright notice,
+//   this list of conditions and the following disclaimer in the documentation
+//   and/or other materials provided with the distribution.
+// * Neither the name of Google Inc. nor the names of its contributors may be
+//   used to endorse or promote products derived from this software without
+//   specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE
+// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+// POSSIBILITY OF SUCH DAMAGE.
+//
+// ========================================
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// THIS FILE IS AUTOGENERATED. DO NOT EDIT.
+// ========================================
+//
+// This file is generated using generate_bundle_adjustment_tests.py.
+
+#include "ceres/bundle_adjustment_test_util.h"
+#include "ceres/internal/config.h"
+#include "gtest/gtest.h"
+
+#ifndef CERES_NO_CUDA
+#ifndef CERES_NO_CUDSS
+
+namespace ceres::internal {
+
+TEST_F(BundleAdjustmentTest,
+       SparseSchur_CudaSparse_UserOrdering_Threads) {  // NOLINT
+  BundleAdjustmentProblem bundle_adjustment_problem;
+  Solver::Options* options = bundle_adjustment_problem.mutable_solver_options();
+  options->eta = 0.01;
+  options->num_threads = 4;
+  options->linear_solver_type = SPARSE_SCHUR;
+  options->dense_linear_algebra_library_type = EIGEN;
+  options->sparse_linear_algebra_library_type = CUDA_SPARSE;
+  options->preconditioner_type = IDENTITY;
+  if (kUserOrdering) {
+    options->linear_solver_ordering = nullptr;
+  }
+  Problem* problem = bundle_adjustment_problem.mutable_problem();
+  RunSolverForConfigAndExpectResidualsMatch(*options, problem);
+}
+
+}  // namespace ceres::internal
+
+#endif  // CERES_NO_CUDSS
+#endif  // CERES_NO_CUDA
diff --git a/internal/ceres/reorder_program.cc b/internal/ceres/reorder_program.cc
index 44c4e46..268c477 100644
--- a/internal/ceres/reorder_program.cc
+++ b/internal/ceres/reorder_program.cc
@@ -576,7 +576,9 @@
     // would involve performing two symbolic factorisations instead of one
     // which would have a negative overall impact on performance.
     return true;
-
+  } else if (sparse_linear_algebra_library_type == CUDA_SPARSE) {
+    // cuDSS has the same limitation as Accelerate
+    return true;
   } else if (sparse_linear_algebra_library_type == EIGEN_SPARSE) {
     OrderingForSparseNormalCholeskyUsingEigenSparse(
         linear_solver_ordering_type,
diff --git a/internal/ceres/solver.cc b/internal/ceres/solver.cc
index 611e465..fadb17b 100644
--- a/internal/ceres/solver.cc
+++ b/internal/ceres/solver.cc
@@ -41,6 +41,7 @@
 #include "ceres/casts.h"
 #include "ceres/context.h"
 #include "ceres/context_impl.h"
+#include "ceres/cuda_sparse_cholesky.h"
 #include "ceres/detect_structure.h"
 #include "ceres/eigensparse.h"
 #include "ceres/gradient_checking_cost_function.h"
@@ -113,7 +114,12 @@
            internal::SuiteSparse::IsNestedDissectionAvailable()) ||
           (type == ACCELERATE_SPARSE) ||
           ((type == EIGEN_SPARSE) &&
-           internal::EigenSparse::IsNestedDissectionAvailable()));
+           internal::EigenSparse::IsNestedDissectionAvailable())
+#ifndef CERES_NO_CUDSS
+          || ((type == CUDA_SPARSE) &&
+              internal::CudaSparseCholesky<>::IsNestedDissectionAvailable())
+#endif
+  );
 }
 
 bool IsIterativeSolver(LinearSolverType type) {
@@ -168,8 +174,15 @@
     return false;
   }
 
-  if (!IsSparseLinearAlgebraLibraryTypeAvailable(
+  if (IsSparseLinearAlgebraLibraryTypeAvailable(
           options.sparse_linear_algebra_library_type)) {
+    if (options.sparse_linear_algebra_library_type == CUDA_SPARSE) {
+#if defined(CERES_NO_CUDSS)
+      *error = StringPrintf(kNoLibraryFormat, solver_name, library_name);
+      return false;
+#endif
+    }
+  } else {
     *error = StringPrintf(kNoLibraryFormat, solver_name, library_name);
     return false;
   }
@@ -678,7 +691,12 @@
       options.linear_solver_type == DENSE_QR) {
     return (options.dense_linear_algebra_library_type == CUDA);
   }
-  if (options.linear_solver_type == CGNR) {
+  if (options.linear_solver_type == CGNR ||
+      options.linear_solver_type == SPARSE_SCHUR ||
+      options.linear_solver_type == SPARSE_NORMAL_CHOLESKY ||
+      (options.linear_solver_type == ITERATIVE_SCHUR &&
+       (options.preconditioner_type == CLUSTER_JACOBI ||
+        options.preconditioner_type == CLUSTER_TRIDIAGONAL))) {
     return (options.sparse_linear_algebra_library_type == CUDA_SPARSE);
   }
   return false;
diff --git a/internal/ceres/solver_utils.cc b/internal/ceres/solver_utils.cc
index f5fbf05..ac22210 100644
--- a/internal/ceres/solver_utils.cc
+++ b/internal/ceres/solver_utils.cc
@@ -36,6 +36,9 @@
 #include "ceres/version.h"
 #ifndef CERES_NO_CUDA
 #include "cuda_runtime.h"
+#ifndef CERES_NO_CUDSS
+#include "cudss.h"
+#endif  // CERES_NO_CUDSS
 #endif  // CERES_NO_CUDA
 
 namespace ceres::internal {
@@ -79,6 +82,11 @@
 
 #ifndef CERES_NO_CUDA
   "-cuda-(" CERES_TO_STRING(CUDART_VERSION) ")"
+#ifndef CERES_NO_CUDSS
+  "-cudss-(" CERES_SEMVER_VERSION(CUDSS_VERSION_MAJOR,
+                                  CUDSS_VERSION_MINOR,
+                                  CUDSS_VERSION_PATCH) ")"
+#endif // CERES_NO_CUDSS
 #endif
   ;
 // clang-format on
diff --git a/internal/ceres/sparse_cholesky.cc b/internal/ceres/sparse_cholesky.cc
index 4f1bf87..ecfeb29 100644
--- a/internal/ceres/sparse_cholesky.cc
+++ b/internal/ceres/sparse_cholesky.cc
@@ -34,6 +34,7 @@
 #include <utility>
 
 #include "ceres/accelerate_sparse.h"
+#include "ceres/cuda_sparse_cholesky.h"
 #include "ceres/eigensparse.h"
 #include "ceres/float_suitesparse.h"
 #include "ceres/iterative_refiner.h"
@@ -87,6 +88,19 @@
       LOG(FATAL) << "Ceres was compiled without support for Apple's Accelerate "
                  << "framework solvers.";
 #endif
+    case CUDA_SPARSE:
+#ifndef CERES_NO_CUDSS
+      if (options.use_mixed_precision_solves) {
+        sparse_cholesky = CudaSparseCholesky<float>::Create(
+            options.context, options.ordering_type);
+      } else {
+        sparse_cholesky = CudaSparseCholesky<double>::Create(
+            options.context, options.ordering_type);
+      }
+      break;
+#else   // CERES_NO_CUDSS
+      LOG(FATAL) << "Ceres was compiled without support for cuDSS.";
+#endif  // CERES_NO_CUDSS
 
     default:
       LOG(FATAL) << "Unknown sparse linear algebra library type : "
diff --git a/internal/ceres/sparse_cholesky_test.cc b/internal/ceres/sparse_cholesky_test.cc
index d0d962e..bb5c7c8 100644
--- a/internal/ceres/sparse_cholesky_test.cc
+++ b/internal/ceres/sparse_cholesky_test.cc
@@ -39,6 +39,7 @@
 #include "Eigen/SparseCore"
 #include "ceres/block_sparse_matrix.h"
 #include "ceres/compressed_row_sparse_matrix.h"
+#include "ceres/cuda_sparse_cholesky.h"
 #include "ceres/inner_product_computer.h"
 #include "ceres/internal/config.h"
 #include "ceres/internal/eigen.h"
@@ -117,6 +118,12 @@
   sparse_cholesky_options.sparse_linear_algebra_library_type =
       sparse_linear_algebra_library_type;
   sparse_cholesky_options.ordering_type = ordering_type;
+#ifndef CERES_NO_CUDSS
+  ContextImpl context;
+  sparse_cholesky_options.context = &context;
+  std::string error;
+  CHECK(context.InitCuda(&error)) << error;
+#endif  // CERES_NO_CUDSS
   auto sparse_cholesky = SparseCholesky::Create(sparse_cholesky_options);
   const CompressedRowSparseMatrix::StorageType storage_type =
       sparse_cholesky->StorageType();
@@ -251,7 +258,7 @@
     ParamInfoToString);
 
 INSTANTIATE_TEST_SUITE_P(
-    EigenSparseCholeskySingle,
+    FloatEigenSparseCholesky,
     SparseCholeskyTest,
     ::testing::Combine(::testing::Values(EIGEN_SPARSE),
                        ::testing::Values(OrderingType::AMD,
@@ -260,6 +267,26 @@
     ParamInfoToString);
 #endif  // CERES_USE_EIGEN_SPARSE
 
+#ifndef CERES_NO_CUDSS
+using CuDSSCholeskyDouble = CudaSparseCholesky<double>;
+INSTANTIATE_TEST_SUITE_P(
+    CuDSSCholeskyDouble,
+    SparseCholeskyTest,
+    ::testing::Combine(::testing::Values(CUDA_SPARSE),
+                       ::testing::Values(OrderingType::AMD),
+                       ::testing::Values(true, false)),
+    ParamInfoToString);
+
+using CuDSSCholeskySingle = CudaSparseCholesky<float>;
+INSTANTIATE_TEST_SUITE_P(
+    CuDSSCholeskySingle,
+    SparseCholeskyTest,
+    ::testing::Combine(::testing::Values(CUDA_SPARSE),
+                       ::testing::Values(OrderingType::AMD),
+                       ::testing::Values(true, false)),
+    ParamInfoToString);
+#endif  // CERES_NO_CUDSS
+
 #if defined(CERES_USE_EIGEN_SPARSE) && !defined(CERES_NO_EIGEN_METIS)
 INSTANTIATE_TEST_SUITE_P(
     EigenSparseCholeskyMETIS,
diff --git a/internal/ceres/sparse_normal_cholesky_solver_test.cc b/internal/ceres/sparse_normal_cholesky_solver_test.cc
index 3396e34..2ff3f6e 100644
--- a/internal/ceres/sparse_normal_cholesky_solver_test.cc
+++ b/internal/ceres/sparse_normal_cholesky_solver_test.cc
@@ -88,8 +88,9 @@
 
     EXPECT_EQ(summary.termination_type, LinearSolverTerminationType::SUCCESS);
 
+    const double eps = options.use_mixed_precision_solves ? 2e-6 : 1e-8;
     for (int i = 0; i < A_->num_cols(); ++i) {
-      EXPECT_NEAR(expected_solution(i), actual_solution(i), 1e-8)
+      EXPECT_NEAR(expected_solution(i), actual_solution(i), eps)
           << "\nExpected: " << expected_solution.transpose()
           << "\nActual: " << actual_solution.transpose();
     }
@@ -177,4 +178,31 @@
 }
 #endif  // CERES_USE_EIGEN_SPARSE
 
+#ifndef CERES_NO_CUDSS
+TEST_F(SparseNormalCholeskySolverTest, SparseNormalCholeskyUsingCuDSSSingle) {
+  LinearSolver::Options options;
+  options.sparse_linear_algebra_library_type = CUDA_SPARSE;
+  options.type = SPARSE_NORMAL_CHOLESKY;
+  options.ordering_type = OrderingType::AMD;
+  options.use_mixed_precision_solves = true;
+  ContextImpl context;
+  options.context = &context;
+  std::string error;
+  CHECK(context.InitCuda(&error)) << error;
+  TestSolver(options);
+}
+
+TEST_F(SparseNormalCholeskySolverTest, SparseNormalCholeskyUsingCuDSSDouble) {
+  LinearSolver::Options options;
+  options.sparse_linear_algebra_library_type = CUDA_SPARSE;
+  options.type = SPARSE_NORMAL_CHOLESKY;
+  options.ordering_type = OrderingType::AMD;
+  ContextImpl context;
+  options.context = &context;
+  std::string error;
+  CHECK(context.InitCuda(&error)) << error;
+  TestSolver(options);
+}
+#endif  // CERES_USE_EIGEN_SPARSE
+
 }  // namespace ceres::internal
diff --git a/internal/ceres/visibility_based_preconditioner.cc b/internal/ceres/visibility_based_preconditioner.cc
index 42e8a6e..34d6e0c 100644
--- a/internal/ceres/visibility_based_preconditioner.cc
+++ b/internal/ceres/visibility_based_preconditioner.cc
@@ -100,6 +100,8 @@
   sparse_cholesky_options.sparse_linear_algebra_library_type =
       options_.sparse_linear_algebra_library_type;
   sparse_cholesky_options.ordering_type = options_.ordering_type;
+  sparse_cholesky_options.context = options_.context;
+
   sparse_cholesky_ = SparseCholesky::Create(sparse_cholesky_options);
 
   const time_t init_time = time(nullptr);