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

@@ -156,8 +156,8 @@ if (USE_CUDA)
         COMMAND ${CUDA_TOOLKIT_ROOT_DIR}/bin/cuda-memcheck --leak-check full
             $<TARGET_FILE:cuda_dense_cholesky_test>)
   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)

internal/ceres/cuda_kernels.cu.cc

@@ -40,7 +40,7 @@ namespace ceres::internal {
 // toolkit documentation.
 constexpr int kCudaBlockSize = 256;
 
-template<typename SrcType, typename DstType>
+template <typename SrcType, typename DstType>
 __global__ void TypeConversionKernel(const SrcType* __restrict__ input,
                                      DstType* __restrict__ output,
                                      const int size) {
@@ -68,7 +68,7 @@ void CudaFP32ToFP64(const float* input,
       <<<num_blocks, kCudaBlockSize, 0, stream>>>(input, output, size);
 }
 
-template<typename T>
+template <typename T>
 __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<double><<<num_blocks, kCudaBlockSize, 0, stream>>>(
-      output, size);
+  SetZeroKernel<double>
+      <<<num_blocks, kCudaBlockSize, 0, stream>>>(output, size);
 }
 
 template <typename SrcType, typename DstType>
 __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<float, double>
       <<<num_blocks, kCudaBlockSize, 0, stream>>>(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<<<num_blocks, kCudaBlockSize, 0, stream>>>(
-      y, D, x, size);
+  CudaDtDxpyKernel<<<num_blocks, kCudaBlockSize, 0, stream>>>(y, D, x, size);
 }
 
-} // namespace ceres_cuda_kernels
+}  // namespace ceres::internal

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_

internal/ceres/cuda_kernels_test.cc

@@ -28,13 +28,14 @@
 //
 // Author: joydeepb@cs.utexas.edu (Joydeep Biswas)
 
+#include "ceres/cuda_kernels.h"
+
 #include <math.h>
 
 #include <limits>
 #include <string>
 #include <vector>
 
-#include "ceres/ceres_cuda_kernels.h"
 #include "ceres/cuda_buffer.h"
 #include "ceres/internal/config.h"
 #include "ceres/internal/eigen.h"

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
+#endif  // CERES_NO_CUDA

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
+#endif  // CERES_NO_CUDA

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"

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

internal/ceres/dense_qr.cc

@@ -33,6 +33,7 @@
 #include <algorithm>
 #include <memory>
 #include <string>
+
 #ifndef CERES_NO_CUDA
 #include "ceres/context_impl.h"
 #include "cublas_v2.h"