Rename CudaBuffer to CudaDeviceBuffer Clarify that CudaBuffer manages device memory in contrast to CudaPinnedHostBuffer and CudaStreamedBuffer. Thanks to @S-o-T for suggesting this. Fixes https://github.com/ceres-solver/ceres-solver/issues/1077 Change-Id: Ifa19d556984d47611c38e33fb2024f94eca6fd29
diff --git a/internal/ceres/context_impl.h b/internal/ceres/context_impl.h index 86381f4..423fd54 100644 --- a/internal/ceres/context_impl.h +++ b/internal/ceres/context_impl.h
@@ -90,7 +90,7 @@ // - Only use default stream // - Return control to the callee without synchronization whenever possible // - Stream synchronization occurs only after GPU to CPU transfers, and is - // handled by CudaBuffer + // handled by CudaDeviceBuffer // // b. Multi-stream workloads // Multi-stream workloads are more restricted in order to make it harder to
diff --git a/internal/ceres/cuda_block_sparse_crs_view.cc b/internal/ceres/cuda_block_sparse_crs_view.cc index 7b8cda7..9b04d08 100644 --- a/internal/ceres/cuda_block_sparse_crs_view.cc +++ b/internal/ceres/cuda_block_sparse_crs_view.cc
@@ -43,8 +43,8 @@ : context_(context) { block_structure_ = std::make_unique<CudaBlockSparseStructure>( *bsm.block_structure(), context); - CudaBuffer<int32_t> rows(context, bsm.num_rows() + 1); - CudaBuffer<int32_t> cols(context, bsm.num_nonzeros()); + CudaDeviceBuffer<int32_t> rows(context, bsm.num_rows() + 1); + CudaDeviceBuffer<int32_t> cols(context, bsm.num_nonzeros()); FillCRSStructure(block_structure_->num_row_blocks(), bsm.num_rows(), block_structure_->first_cell_in_row_block(),
diff --git a/internal/ceres/cuda_block_structure.h b/internal/ceres/cuda_block_structure.h index 6da6fdd..b2d0035 100644 --- a/internal/ceres/cuda_block_structure.h +++ b/internal/ceres/cuda_block_structure.h
@@ -107,11 +107,11 @@ int num_row_blocks_e_; int num_col_blocks_; bool is_crs_compatible_; - CudaBuffer<int> first_cell_in_row_block_; - CudaBuffer<int> value_offset_row_block_f_; - CudaBuffer<Cell> cells_; - CudaBuffer<Block> row_blocks_; - CudaBuffer<Block> col_blocks_; + CudaDeviceBuffer<int> first_cell_in_row_block_; + CudaDeviceBuffer<int> value_offset_row_block_f_; + CudaDeviceBuffer<Cell> cells_; + CudaDeviceBuffer<Block> row_blocks_; + CudaDeviceBuffer<Block> col_blocks_; friend class CudaBlockStructureTest; }; } // namespace ceres::internal
diff --git a/internal/ceres/cuda_buffer.h b/internal/ceres/cuda_buffer.h index c5756ea..543c19a 100644 --- a/internal/ceres/cuda_buffer.h +++ b/internal/ceres/cuda_buffer.h
@@ -51,23 +51,23 @@ // threads, since active Cuda devices are determined by the cuda runtime on a // per-thread basis. template <typename T> -class CudaBuffer { +class CudaDeviceBuffer { public: - explicit CudaBuffer(ContextImpl* context) : context_(context) {} - CudaBuffer(ContextImpl* context, int size) : context_(context) { + explicit CudaDeviceBuffer(ContextImpl* context) : context_(context) {} + CudaDeviceBuffer(ContextImpl* context, int size) : context_(context) { Reserve(size); } - CudaBuffer(CudaBuffer&& other) + CudaDeviceBuffer(CudaDeviceBuffer&& other) : data_(other.data_), size_(other.size_), context_(other.context_) { other.data_ = nullptr; other.size_ = 0; } - CudaBuffer(const CudaBuffer&) = delete; - CudaBuffer& operator=(const CudaBuffer&) = delete; + CudaDeviceBuffer(const CudaDeviceBuffer&) = delete; + CudaDeviceBuffer& operator=(const CudaDeviceBuffer&) = delete; - ~CudaBuffer() { + ~CudaDeviceBuffer() { if (data_ != nullptr) { CHECK_EQ(cudaFree(data_), cudaSuccess); } @@ -88,7 +88,7 @@ } // Perform an asynchronous copy from CPU memory to GPU memory managed by this - // CudaBuffer instance using the stream provided. + // CudaDeviceBuffer instance using the stream provided. void CopyFromCpu(const T* data, const size_t size) { Reserve(size); CHECK_EQ(cudaMemcpyAsync(data_, @@ -100,7 +100,7 @@ } // Perform an asynchronous copy from a vector in CPU memory to GPU memory - // managed by this CudaBuffer instance. + // managed by this CudaDeviceBuffer instance. void CopyFromCpuVector(const std::vector<T>& data) { Reserve(data.size()); CHECK_EQ(cudaMemcpyAsync(data_, @@ -112,7 +112,7 @@ } // Perform an asynchronous copy from another GPU memory array to the GPU - // memory managed by this CudaBuffer instance using the stream provided. + // memory managed by this CudaDeviceBuffer instance using the stream provided. void CopyFromGPUArray(const T* data, const size_t size) { Reserve(size); CHECK_EQ(cudaMemcpyAsync(data_, @@ -123,7 +123,7 @@ cudaSuccess); } - // Copy data from the GPU memory managed by this CudaBuffer instance to CPU + // Copy data from the GPU memory managed by this CudaDeviceBuffer instance to CPU // memory. It is the caller's responsibility to ensure that the CPU memory // pointer is valid, i.e. it is not null, and that it points to memory of // at least this->size() size. This method ensures all previously dispatched @@ -141,9 +141,9 @@ } // Copy N items from another GPU memory array to the GPU memory managed by - // this CudaBuffer instance, growing this buffer's size if needed. This copy + // this CudaDeviceBuffer instance, growing this buffer's size if needed. This copy // is asynchronous, and operates on the stream provided. - void CopyNItemsFrom(int n, const CudaBuffer<T>& other) { + void CopyNItemsFrom(int n, const CudaDeviceBuffer<T>& other) { Reserve(n); CHECK(other.data_ != nullptr); CHECK(data_ != nullptr); @@ -155,11 +155,11 @@ cudaSuccess); } - // Return a pointer to the GPU memory managed by this CudaBuffer instance. + // Return a pointer to the GPU memory managed by this CudaDeviceBuffer instance. T* data() { return data_; } const T* data() const { return data_; } // Return the number of items of type T that can fit in the GPU memory - // allocated so far by this CudaBuffer instance. + // allocated so far by this CudaDeviceBuffer instance. size_t size() const { return size_; } private:
diff --git a/internal/ceres/cuda_kernels_vector_ops_test.cc b/internal/ceres/cuda_kernels_vector_ops_test.cc index f32d1d2..26fe65b 100644 --- a/internal/ceres/cuda_kernels_vector_ops_test.cc +++ b/internal/ceres/cuda_kernels_vector_ops_test.cc
@@ -52,9 +52,9 @@ std::string cuda_error; EXPECT_TRUE(context.InitCuda(&cuda_error)) << cuda_error; std::vector<double> fp64_cpu = {1.0, 1.5, 2.0, 2.5, 3.0, 3.5, 4.0, 4.5, 5.0}; - CudaBuffer<double> fp64_gpu(&context); + CudaDeviceBuffer<double> fp64_gpu(&context); fp64_gpu.CopyFromCpuVector(fp64_cpu); - CudaBuffer<float> fp32_gpu(&context); + CudaDeviceBuffer<float> fp32_gpu(&context); fp32_gpu.Reserve(fp64_cpu.size()); CudaFP64ToFP32(fp64_gpu.data(), fp32_gpu.data(), @@ -79,9 +79,9 @@ EXPECT_GT(fp64_cpu[1], 0.0); EXPECT_TRUE(std::isfinite(fp64_cpu[2])); EXPECT_TRUE(std::isfinite(fp64_cpu[3])); - CudaBuffer<double> fp64_gpu(&context); + CudaDeviceBuffer<double> fp64_gpu(&context); fp64_gpu.CopyFromCpuVector(fp64_cpu); - CudaBuffer<float> fp32_gpu(&context); + CudaDeviceBuffer<float> fp32_gpu(&context); fp32_gpu.Reserve(fp64_cpu.size()); CudaFP64ToFP32(fp64_gpu.data(), fp32_gpu.data(), @@ -100,9 +100,9 @@ std::string cuda_error; EXPECT_TRUE(context.InitCuda(&cuda_error)) << cuda_error; std::vector<float> fp32_cpu = {1.0, 1.5, 2.0, 2.5, 3.0, 3.5, 4.0, 4.5, 5.0}; - CudaBuffer<float> fp32_gpu(&context); + CudaDeviceBuffer<float> fp32_gpu(&context); fp32_gpu.CopyFromCpuVector(fp32_cpu); - CudaBuffer<double> fp64_gpu(&context); + CudaDeviceBuffer<double> fp64_gpu(&context); fp64_gpu.Reserve(fp32_cpu.size()); CudaFP32ToFP64(fp32_gpu.data(), fp64_gpu.data(), @@ -120,7 +120,7 @@ std::string cuda_error; EXPECT_TRUE(context.InitCuda(&cuda_error)) << cuda_error; std::vector<float> fp32_cpu = {1.0, 1.5, 2.0, 2.5, 3.0, 3.5, 4.0, 4.5, 5.0}; - CudaBuffer<float> fp32_gpu(&context); + CudaDeviceBuffer<float> fp32_gpu(&context); fp32_gpu.CopyFromCpuVector(fp32_cpu); CudaSetZeroFP32(fp32_gpu.data(), fp32_cpu.size(), context.DefaultStream()); std::vector<float> fp32_cpu_zero(fp32_cpu.size()); @@ -135,7 +135,7 @@ std::string cuda_error; EXPECT_TRUE(context.InitCuda(&cuda_error)) << cuda_error; std::vector<double> fp64_cpu = {1.0, 1.5, 2.0, 2.5, 3.0, 3.5, 4.0, 4.5, 5.0}; - CudaBuffer<double> fp64_gpu(&context); + CudaDeviceBuffer<double> fp64_gpu(&context); fp64_gpu.CopyFromCpuVector(fp64_cpu); CudaSetZeroFP64(fp64_gpu.data(), fp64_cpu.size(), context.DefaultStream()); std::vector<double> fp64_cpu_zero(fp64_cpu.size()); @@ -152,9 +152,9 @@ std::vector<float> fp32_cpu_a = {1.0, 1.5, 2.0, 2.5, 3.0, 3.5, 4.0, 4.5, 5.0}; std::vector<double> fp64_cpu_b = { 1.0, 1.5, 2.0, 2.5, 3.0, 3.5, 4.0, 4.5, 5.0}; - CudaBuffer<float> fp32_gpu_a(&context); + CudaDeviceBuffer<float> fp32_gpu_a(&context); fp32_gpu_a.CopyFromCpuVector(fp32_cpu_a); - CudaBuffer<double> fp64_gpu_b(&context); + CudaDeviceBuffer<double> fp64_gpu_b(&context); fp64_gpu_b.CopyFromCpuVector(fp64_cpu_b); CudaDsxpy(fp64_gpu_b.data(), fp32_gpu_a.data(), @@ -173,11 +173,11 @@ std::vector<double> x_cpu = {1, 2, 3, 4}; std::vector<double> y_cpu = {4, 3, 2, 1}; std::vector<double> d_cpu = {10, 20, 30, 40}; - CudaBuffer<double> x_gpu(&context); + CudaDeviceBuffer<double> x_gpu(&context); x_gpu.CopyFromCpuVector(x_cpu); - CudaBuffer<double> y_gpu(&context); + CudaDeviceBuffer<double> y_gpu(&context); y_gpu.CopyFromCpuVector(y_cpu); - CudaBuffer<double> d_gpu(&context); + CudaDeviceBuffer<double> d_gpu(&context); d_gpu.CopyFromCpuVector(d_cpu); CudaDtDxpy(y_gpu.data(), d_gpu.data(),
diff --git a/internal/ceres/cuda_partitioned_block_sparse_crs_view.cc b/internal/ceres/cuda_partitioned_block_sparse_crs_view.cc index 4d75031..4b648f6 100644 --- a/internal/ceres/cuda_partitioned_block_sparse_crs_view.cc +++ b/internal/ceres/cuda_partitioned_block_sparse_crs_view.cc
@@ -61,10 +61,10 @@ : bsm.num_cols(); const int num_cols_f = bsm.num_cols() - num_cols_e; - CudaBuffer<int32_t> rows_e(context, num_rows + 1); - CudaBuffer<int32_t> cols_e(context, num_nonzeros_e); - CudaBuffer<int32_t> rows_f(context, num_rows + 1); - CudaBuffer<int32_t> cols_f(context, num_nonzeros_f); + CudaDeviceBuffer<int32_t> rows_e(context, num_rows + 1); + CudaDeviceBuffer<int32_t> cols_e(context, num_nonzeros_e); + CudaDeviceBuffer<int32_t> rows_f(context, num_rows + 1); + CudaDeviceBuffer<int32_t> cols_f(context, num_nonzeros_f); num_row_blocks_e_ = block_structure_->num_row_blocks_e(); FillCRSStructurePartitioned(block_structure_->num_row_blocks(),
diff --git a/internal/ceres/cuda_sparse_cholesky.cc b/internal/ceres/cuda_sparse_cholesky.cc index 689ab08..bb39ec1 100644 --- a/internal/ceres/cuda_sparse_cholesky.cc +++ b/internal/ceres/cuda_sparse_cholesky.cc
@@ -441,7 +441,7 @@ void ConvertAndCopyToDevice(const double* source, Scalar* intermediate, - CudaBuffer<Scalar>& destination) { + CudaDeviceBuffer<Scalar>& destination) { const auto size = destination.size(); if constexpr (std::is_same_v<Scalar, double>) { destination.CopyFromCpu(source, size); @@ -451,7 +451,7 @@ } } - void ConvertAndCopyToHost(const CudaBuffer<Scalar>& source, + void ConvertAndCopyToHost(const CudaDeviceBuffer<Scalar>& source, Scalar* intermediate, double* destination) { const auto size = source.size(); @@ -472,11 +472,11 @@ 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_; + CudaDeviceBuffer<int> lhs_rows_d_; + CudaDeviceBuffer<int> lhs_cols_d_; + CudaDeviceBuffer<Scalar> lhs_values_d_; + CudaDeviceBuffer<Scalar> rhs_d_; + CudaDeviceBuffer<Scalar> x_d_; LinearSolverTerminationType analyze_result_ = LinearSolverTerminationType::FATAL_ERROR;
diff --git a/internal/ceres/cuda_sparse_matrix.cc b/internal/ceres/cuda_sparse_matrix.cc index 9cb25d5..a9c22c4 100644 --- a/internal/ceres/cuda_sparse_matrix.cc +++ b/internal/ceres/cuda_sparse_matrix.cc
@@ -102,8 +102,8 @@ } // namespace CudaSparseMatrix::CudaSparseMatrix(int num_cols, - CudaBuffer<int32_t>&& rows, - CudaBuffer<int32_t>&& cols, + CudaDeviceBuffer<int32_t>&& rows, + CudaDeviceBuffer<int32_t>&& cols, ContextImpl* context) : num_rows_(rows.size() - 1), num_cols_(num_cols),
diff --git a/internal/ceres/cuda_sparse_matrix.h b/internal/ceres/cuda_sparse_matrix.h index 2940d1d..4313b8d 100644 --- a/internal/ceres/cuda_sparse_matrix.h +++ b/internal/ceres/cuda_sparse_matrix.h
@@ -67,8 +67,8 @@ // Create matrix from existing row and column index buffers. // Values are left uninitialized. CudaSparseMatrix(int num_cols, - CudaBuffer<int32_t>&& rows, - CudaBuffer<int32_t>&& cols, + CudaDeviceBuffer<int32_t>&& rows, + CudaDeviceBuffer<int32_t>&& cols, ContextImpl* context); ~CudaSparseMatrix(); @@ -121,11 +121,11 @@ ContextImpl* context_ = nullptr; // CSR row indices. - CudaBuffer<int32_t> rows_; + CudaDeviceBuffer<int32_t> rows_; // CSR column indices. - CudaBuffer<int32_t> cols_; + CudaDeviceBuffer<int32_t> cols_; // CSR values. - CudaBuffer<double> values_; + CudaDeviceBuffer<double> values_; // CuSparse object that describes this matrix. cusparseSpMatDescr_t descr_ = nullptr; @@ -134,7 +134,7 @@ cusparseDnVecDescr_t descr_vec_left_ = nullptr; cusparseDnVecDescr_t descr_vec_right_ = nullptr; - mutable CudaBuffer<uint8_t> spmv_buffer_; + mutable CudaDeviceBuffer<uint8_t> spmv_buffer_; }; } // namespace ceres::internal
diff --git a/internal/ceres/cuda_streamed_buffer.h b/internal/ceres/cuda_streamed_buffer.h index a8532d0..d43ab7f 100644 --- a/internal/ceres/cuda_streamed_buffer.h +++ b/internal/ceres/cuda_streamed_buffer.h
@@ -328,7 +328,7 @@ const int kValuesPerBatch; ContextImpl* context_ = nullptr; - CudaBuffer<T> values_gpu_; + CudaDeviceBuffer<T> values_gpu_; T* values_cpu_pinned_ = nullptr; cudaEvent_t copy_finished_[kNumBatches] = {nullptr}; };
diff --git a/internal/ceres/cuda_vector.h b/internal/ceres/cuda_vector.h index 8db5649..0351d2d 100644 --- a/internal/ceres/cuda_vector.h +++ b/internal/ceres/cuda_vector.h
@@ -119,7 +119,7 @@ int num_rows_ = 0; ContextImpl* context_ = nullptr; - CudaBuffer<double> data_; + CudaDeviceBuffer<double> data_; // CuSparse object that describes this dense vector. cusparseDnVecDescr_t descr_ = nullptr; };
diff --git a/internal/ceres/dense_cholesky.h b/internal/ceres/dense_cholesky.h index 4f55c58..7f2e2ca 100644 --- a/internal/ceres/dense_cholesky.h +++ b/internal/ceres/dense_cholesky.h
@@ -216,13 +216,13 @@ // and *Solve. size_t num_cols_ = 0; // GPU memory allocated for the A matrix (lhs matrix). - CudaBuffer<double> lhs_; + CudaDeviceBuffer<double> lhs_; // GPU memory allocated for the B matrix (rhs vector). - CudaBuffer<double> rhs_; + CudaDeviceBuffer<double> rhs_; // Scratch space for cuSOLVER on the GPU. - CudaBuffer<double> device_workspace_; + CudaDeviceBuffer<double> device_workspace_; // Required for error handling with cuSOLVER. - CudaBuffer<int> error_; + CudaDeviceBuffer<int> error_; // Cache the result of Factorize to ensure that when Solve is called, the // factorization of lhs is valid. LinearSolverTerminationType factorize_result_ = @@ -277,21 +277,21 @@ // Number of columns in the A matrix, to be cached between calls to *Factorize // and *Solve. size_t num_cols_ = 0; - CudaBuffer<double> lhs_fp64_; - CudaBuffer<double> rhs_fp64_; - CudaBuffer<float> lhs_fp32_; + CudaDeviceBuffer<double> lhs_fp64_; + CudaDeviceBuffer<double> rhs_fp64_; + CudaDeviceBuffer<float> lhs_fp32_; // Scratch space for cuSOLVER on the GPU. - CudaBuffer<float> device_workspace_; + CudaDeviceBuffer<float> device_workspace_; // Required for error handling with cuSOLVER. - CudaBuffer<int> error_; + CudaDeviceBuffer<int> error_; // Solution to lhs * x = rhs. - CudaBuffer<double> x_fp64_; + CudaDeviceBuffer<double> x_fp64_; // Incremental correction to x. - CudaBuffer<float> correction_fp32_; + CudaDeviceBuffer<float> correction_fp32_; // Residual to iterative refinement. - CudaBuffer<float> residual_fp32_; - CudaBuffer<double> residual_fp64_; + CudaDeviceBuffer<float> residual_fp32_; + CudaDeviceBuffer<double> residual_fp64_; // Number of inner refinement iterations to perform. int max_num_refinement_iterations_ = 0;
diff --git a/internal/ceres/dense_qr.h b/internal/ceres/dense_qr.h index af35827..246db69 100644 --- a/internal/ceres/dense_qr.h +++ b/internal/ceres/dense_qr.h
@@ -174,15 +174,15 @@ // and *Solve. size_t num_cols_ = 0; // GPU memory allocated for the A matrix (lhs matrix). - CudaBuffer<double> lhs_; + CudaDeviceBuffer<double> lhs_; // GPU memory allocated for the B matrix (rhs vector). - CudaBuffer<double> rhs_; + CudaDeviceBuffer<double> rhs_; // GPU memory allocated for the TAU matrix (scaling of householder vectors). - CudaBuffer<double> tau_; + CudaDeviceBuffer<double> tau_; // Scratch space for cuSOLVER on the GPU. - CudaBuffer<double> device_workspace_; + CudaDeviceBuffer<double> device_workspace_; // Required for error handling with cuSOLVER. - CudaBuffer<int> error_; + CudaDeviceBuffer<int> error_; // Cache the result of Factorize to ensure that when Solve is called, the // factiorization of lhs is valid. LinearSolverTerminationType factorize_result_ =