Add cuDSS as sparse Cholesky solver
cuDSS could be used as an alternative for SuiteSparse and EigenSparse
in case if CUDA capable GPU is available.
Change-Id: I7a567093ce91363478118153e181134ed5804573
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);
