From 6085e45be71b4e472d76dfac1a58e1ed28b93a83 Mon Sep 17 00:00:00 2001 From: Sameer Agarwal Date: Wed, 17 Aug 2022 16:27:13 -0700 Subject: [PATCH] Minor CUDA cleanup. 1. ceres_cuda_kernels.cu -> cuda_kernels.cu.cc 2. Add missing ifdef guards. 3. Fix an errant namespace Change-Id: I81a5bfbe3c795ff0b3ef41c3bcaa037b99d5b254 --- internal/ceres/CMakeLists.txt | 4 ++-- ...res_cuda_kernels.cu => cuda_kernels.cu.cc} | 20 ++++++++----------- .../{ceres_cuda_kernels.h => cuda_kernels.h} | 5 +++++ internal/ceres/cuda_kernels_test.cc | 3 ++- internal/ceres/cuda_sparse_matrix.cc | 4 ++-- internal/ceres/cuda_vector.cc | 4 ++-- internal/ceres/cuda_vector.h | 2 +- internal/ceres/dense_cholesky.cc | 2 +- internal/ceres/dense_qr.cc | 1 + 9 files changed, 24 insertions(+), 21 deletions(-) rename internal/ceres/{ceres_cuda_kernels.cu => cuda_kernels.cu.cc} (91%) rename internal/ceres/{ceres_cuda_kernels.h => cuda_kernels.h} (96%) diff --git a/internal/ceres/CMakeLists.txt b/internal/ceres/CMakeLists.txt index dae16b532..947003dae 100644 --- a/internal/ceres/CMakeLists.txt +++ b/internal/ceres/CMakeLists.txt @@ -156,8 +156,8 @@ if (USE_CUDA) COMMAND ${CUDA_TOOLKIT_ROOT_DIR}/bin/cuda-memcheck --leak-check full $) endif (BUILD_TESTING AND GFLAGS) - set_source_files_properties(ceres_cuda_kernels.cu PROPERTIES LANGUAGE CUDA) - add_library(ceres_cuda_kernels ceres_cuda_kernels.cu) + set_source_files_properties(cuda_kernels.cu.cc PROPERTIES LANGUAGE CUDA) + add_library(ceres_cuda_kernels cuda_kernels.cu.cc) target_compile_features(ceres_cuda_kernels PRIVATE cxx_std_14) list(APPEND CERES_LIBRARY_PRIVATE_DEPENDENCIES ceres_cuda_kernels) endif (USE_CUDA) diff --git a/internal/ceres/ceres_cuda_kernels.cu b/internal/ceres/cuda_kernels.cu.cc similarity index 91% rename from internal/ceres/ceres_cuda_kernels.cu rename to internal/ceres/cuda_kernels.cu.cc index cd045e33a..2a1e75574 100644 --- a/internal/ceres/ceres_cuda_kernels.cu +++ b/internal/ceres/cuda_kernels.cu.cc @@ -40,7 +40,7 @@ namespace ceres::internal { // toolkit documentation. constexpr int kCudaBlockSize = 256; -template +template __global__ void TypeConversionKernel(const SrcType* __restrict__ input, DstType* __restrict__ output, const int size) { @@ -68,7 +68,7 @@ void CudaFP32ToFP64(const float* input, <<>>(input, output, size); } -template +template __global__ void SetZeroKernel(T* __restrict__ output, const int size) { const int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < size) { @@ -83,24 +83,21 @@ void CudaSetZeroFP32(float* output, const int size, cudaStream_t stream) { void CudaSetZeroFP64(double* output, const int size, cudaStream_t stream) { const int num_blocks = (size + kCudaBlockSize - 1) / kCudaBlockSize; - SetZeroKernel<<>>( - output, size); + SetZeroKernel + <<>>(output, size); } template __global__ void XPlusEqualsYKernel(DstType* __restrict__ x, const SrcType* __restrict__ y, - const int size) { + const int size) { const int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < size) { x[i] = x[i] + DstType(y[i]); } } -void CudaDsxpy(double* x, - float* y, - const int size, - cudaStream_t stream) { +void CudaDsxpy(double* x, float* y, const int size, cudaStream_t stream) { const int num_blocks = (size + kCudaBlockSize - 1) / kCudaBlockSize; XPlusEqualsYKernel <<>>(x, y, size); @@ -122,8 +119,7 @@ void CudaDtDxpy(double* y, const int size, cudaStream_t stream) { const int num_blocks = (size + kCudaBlockSize - 1) / kCudaBlockSize; - CudaDtDxpyKernel<<>>( - y, D, x, size); + CudaDtDxpyKernel<<>>(y, D, x, size); } -} // namespace ceres_cuda_kernels \ No newline at end of file +} // namespace ceres::internal diff --git a/internal/ceres/ceres_cuda_kernels.h b/internal/ceres/cuda_kernels.h similarity index 96% rename from internal/ceres/ceres_cuda_kernels.h rename to internal/ceres/cuda_kernels.h index 91675e80b..d347c206c 100644 --- a/internal/ceres/ceres_cuda_kernels.h +++ b/internal/ceres/cuda_kernels.h @@ -28,6 +28,9 @@ // // Author: joydeepb@cs.utexas.edu (Joydeep Biswas) +#ifndef CERES_INTERNAL_CUDA_KERNELS_H_ +#define CERES_INTERNAL_CUDA_KERNELS_H_ + #include "ceres/internal/config.h" #ifndef CERES_NO_CUDA @@ -72,3 +75,5 @@ void CudaDtDxpy(double* y, } // namespace ceres::internal #endif // CERES_NO_CUDA + +#endif // CERES_INTERNAL_CUDA_KERNELS_H_ diff --git a/internal/ceres/cuda_kernels_test.cc b/internal/ceres/cuda_kernels_test.cc index 83b922d12..053b4428a 100644 --- a/internal/ceres/cuda_kernels_test.cc +++ b/internal/ceres/cuda_kernels_test.cc @@ -28,13 +28,14 @@ // // Author: joydeepb@cs.utexas.edu (Joydeep Biswas) +#include "ceres/cuda_kernels.h" + #include #include #include #include -#include "ceres/ceres_cuda_kernels.h" #include "ceres/cuda_buffer.h" #include "ceres/internal/config.h" #include "ceres/internal/eigen.h" diff --git a/internal/ceres/cuda_sparse_matrix.cc b/internal/ceres/cuda_sparse_matrix.cc index da6498151..e366112dd 100644 --- a/internal/ceres/cuda_sparse_matrix.cc +++ b/internal/ceres/cuda_sparse_matrix.cc @@ -51,8 +51,8 @@ #ifndef CERES_NO_CUDA -#include "ceres/ceres_cuda_kernels.h" #include "ceres/cuda_buffer.h" +#include "ceres/cuda_kernels.h" #include "ceres/cuda_vector.h" #include "cusparse.h" @@ -146,4 +146,4 @@ void CudaSparseMatrix::LeftMultiplyAndAccumulate(const CudaVector& x, } // namespace ceres::internal -#endif // CERES_NO_CUDA \ No newline at end of file +#endif // CERES_NO_CUDA diff --git a/internal/ceres/cuda_vector.cc b/internal/ceres/cuda_vector.cc index 7debeba5c..7bac13a5a 100644 --- a/internal/ceres/cuda_vector.cc +++ b/internal/ceres/cuda_vector.cc @@ -43,8 +43,8 @@ #ifndef CERES_NO_CUDA -#include "ceres/ceres_cuda_kernels.h" #include "ceres/cuda_buffer.h" +#include "ceres/cuda_kernels.h" #include "ceres/cuda_vector.h" #include "cublas_v2.h" @@ -178,4 +178,4 @@ void CudaVector::Scale(double s) { } // namespace ceres::internal -#endif // CERES_NO_CUDA \ No newline at end of file +#endif // CERES_NO_CUDA diff --git a/internal/ceres/cuda_vector.h b/internal/ceres/cuda_vector.h index e7c4b81bd..985eac0cd 100644 --- a/internal/ceres/cuda_vector.h +++ b/internal/ceres/cuda_vector.h @@ -49,8 +49,8 @@ #ifndef CERES_NO_CUDA -#include "ceres/ceres_cuda_kernels.h" #include "ceres/cuda_buffer.h" +#include "ceres/cuda_kernels.h" #include "ceres/internal/eigen.h" #include "cublas_v2.h" #include "cusparse.h" diff --git a/internal/ceres/dense_cholesky.cc b/internal/ceres/dense_cholesky.cc index af42e73bf..a3e578f54 100644 --- a/internal/ceres/dense_cholesky.cc +++ b/internal/ceres/dense_cholesky.cc @@ -40,8 +40,8 @@ #include "ceres/iterative_refiner.h" #ifndef CERES_NO_CUDA -#include "ceres/ceres_cuda_kernels.h" #include "ceres/context_impl.h" +#include "ceres/cuda_kernels.h" #include "cuda_runtime.h" #include "cusolverDn.h" #endif // CERES_NO_CUDA diff --git a/internal/ceres/dense_qr.cc b/internal/ceres/dense_qr.cc index 775073d88..fb3c22821 100644 --- a/internal/ceres/dense_qr.cc +++ b/internal/ceres/dense_qr.cc @@ -33,6 +33,7 @@ #include #include #include + #ifndef CERES_NO_CUDA #include "ceres/context_impl.h" #include "cublas_v2.h"