Unify allocate/delete_device_data in base/cuda.h

diff --git a/include/deal.II/base/cuda.h b/include/deal.II/base/cuda.h
index 9a4d57dd4771ade59f5a7f0793aad39430ef5c11..e701af54e4d602968a5e82254af375f764269ac9 100644 (file)
--- a/include/deal.II/base/cuda.h
+++ b/include/deal.II/base/cuda.h
@@ -103,6 +103,39 @@ namespace Utilities
 #endif
     }
 
+    /**
+     * Deleter to be used for `std::unique_ptr` pointing to device memory.
+     */
+    template <typename Number>
+    void
+    delete_device_data(Number *device_ptr) noexcept
+    {
+#ifdef DEAL_II_COMPILER_CUDA_AWARE
+      const cudaError_t error_code = cudaFree(device_ptr);
+      (void)error_code;
+      AssertNothrow(error_code == cudaSuccess,
+                    dealii::ExcCudaError(cudaGetErrorString(error_code)));
+#else
+      (void)device_ptr;
+#endif
+    }
+
+    /**
+     * Allocator to be used for `std::unique_ptr` pointing to device memory.
+     */
+    template <typename Number>
+    Number *
+    allocate_device_data(const std::size_t size)
+    {
+#ifdef DEAL_II_COMPILER_CUDA_AWARE
+      Number *device_ptr;
+      Utilities::CUDA::malloc(device_ptr, size);
+      return device_ptr;
+#else
+      (void)size;
+#endif
+    }
+
     /**
      * Copy the elements in @p pointer_dev to the host in @p vector_host.
      */

diff --git a/source/lac/cuda_precondition.cu b/source/lac/cuda_precondition.cu
index 0e8c7db8d347cd6734fb85a891b22d3a46603223..01343ce43b0a552326d9f86c71f230be069d42de 100644 (file)
--- a/source/lac/cuda_precondition.cu
+++ b/source/lac/cuda_precondition.cu
@@ -1184,27 +1184,6 @@ namespace
                                        pBufferSizeInBytes);
   }
   */
-
-
-
-  template <typename Number>
-  void
-  delete_device_vector(Number *device_ptr) noexcept
-  {
-    const cudaError_t error_code = cudaFree(device_ptr);
-    (void)error_code;
-    AssertNothrow(error_code == cudaSuccess,
-                  dealii::ExcCudaError(cudaGetErrorString(error_code)));
-  }
-
-  template <typename Number>
-  Number *
-  allocate_device_vector(const std::size_t size)
-  {
-    Number *device_ptr;
-    Utilities::CUDA::malloc(device_ptr, size);
-    return device_ptr;
-  }
 } // namespace
 
 namespace CUDAWrappers
@@ -1220,11 +1199,11 @@ namespace CUDAWrappers
   template <typename Number>
   PreconditionIC<Number>::PreconditionIC(const Utilities::CUDA::Handle &handle)
     : cusparse_handle(handle.cusparse_handle)
-    , P_val_dev(nullptr, delete_device_vector<Number>)
+    , P_val_dev(nullptr, Utilities::CUDA::delete_device_data<Number>)
     , P_row_ptr_dev(nullptr)
     , P_column_index_dev(nullptr)
-    , tmp_dev(nullptr, delete_device_vector<Number>)
-    , buffer_dev(nullptr, delete_device_vector<void>)
+    , tmp_dev(nullptr, Utilities::CUDA::delete_device_data<Number>)
+    , buffer_dev(nullptr, Utilities::CUDA::delete_device_data<void>)
     , policy_L(CUSPARSE_SOLVE_POLICY_USE_LEVEL)
     , policy_Lt(CUSPARSE_SOLVE_POLICY_USE_LEVEL)
     , policy_M(CUSPARSE_SOLVE_POLICY_USE_LEVEL)
@@ -1316,7 +1295,8 @@ namespace CUDAWrappers
     const Number *const A_val_dev       = std::get<0>(cusparse_matrix);
 
     // create a copy of the matrix entries since the algorithm works in-place.
-    P_val_dev.reset(allocate_device_vector<Number>(n_nonzero_elements));
+    P_val_dev.reset(
+      Utilities::CUDA::allocate_device_data<Number>(n_nonzero_elements));
     cudaError_t cuda_status = cudaMemcpy(P_val_dev.get(),
                                          A_val_dev,
                                          n_nonzero_elements * sizeof(Number),
@@ -1327,7 +1307,7 @@ namespace CUDAWrappers
     const cusparseMatDescr_t mat_descr = std::get<3>(cusparse_matrix);
 
     // initialize an internal buffer we need later on
-    tmp_dev.reset(allocate_device_vector<Number>(n_rows));
+    tmp_dev.reset(Utilities::CUDA::allocate_device_data<Number>(n_rows));
 
     // step 3: query how much memory used in csric02 and csrsv2, and allocate
     // the buffer
@@ -1371,10 +1351,10 @@ namespace CUDAWrappers
 
     const int BufferSize =
       std::max(BufferSize_M, std::max(BufferSize_L, BufferSize_Lt));
-    // workaround: since allocate_device_vector needs a type, we pass char
+    // workaround: since allocate_device_data needs a type, we pass char
     // which is required to have size 1.
     buffer_dev.reset(static_cast<void *>(
-      allocate_device_vector<char>(BufferSize / sizeof(char))));
+      Utilities::CUDA::allocate_device_data<char>(BufferSize / sizeof(char))));
 
     // step 4: perform analysis of incomplete Cholesky on M
     //         perform analysis of triangular solve on L
@@ -1524,11 +1504,11 @@ namespace CUDAWrappers
   PreconditionILU<Number>::PreconditionILU(
     const Utilities::CUDA::Handle &handle)
     : cusparse_handle(handle.cusparse_handle)
-    , P_val_dev(nullptr, delete_device_vector<Number>)
+    , P_val_dev(nullptr, Utilities::CUDA::delete_device_data<Number>)
     , P_row_ptr_dev(nullptr)
     , P_column_index_dev(nullptr)
-    , tmp_dev(nullptr, delete_device_vector<Number>)
-    , buffer_dev(nullptr, delete_device_vector<void>)
+    , tmp_dev(nullptr, Utilities::CUDA::delete_device_data<Number>)
+    , buffer_dev(nullptr, Utilities::CUDA::delete_device_data<void>)
     , policy_L(CUSPARSE_SOLVE_POLICY_USE_LEVEL)
     , policy_U(CUSPARSE_SOLVE_POLICY_USE_LEVEL)
     , policy_M(CUSPARSE_SOLVE_POLICY_USE_LEVEL)
@@ -1637,7 +1617,8 @@ namespace CUDAWrappers
     const Number *const A_val_dev       = std::get<0>(cusparse_matrix);
 
     // create a copy of the matrix entries since the algorithm works in-place.
-    P_val_dev.reset(allocate_device_vector<Number>(n_nonzero_elements));
+    P_val_dev.reset(
+      Utilities::CUDA::allocate_device_data<Number>(n_nonzero_elements));
     cudaError_t cuda_status = cudaMemcpy(P_val_dev.get(),
                                          A_val_dev,
                                          n_nonzero_elements * sizeof(Number),
@@ -1648,7 +1629,7 @@ namespace CUDAWrappers
     const cusparseMatDescr_t mat_descr = std::get<3>(cusparse_matrix);
 
     // initialize an internal buffer we need later on
-    tmp_dev.reset(allocate_device_vector<Number>(n_rows));
+    tmp_dev.reset(Utilities::CUDA::allocate_device_data<Number>(n_rows));
 
     // step 3: query how much memory used in csrilu02 and csrsv2, and allocate
     // the buffer
@@ -1692,10 +1673,10 @@ namespace CUDAWrappers
 
     const int BufferSize =
       std::max(BufferSize_M, std::max(BufferSize_L, BufferSize_U));
-    // workaround: since allocate_device_vector needs a type, we pass char
+    // workaround: since allocate_device_data needs a type, we pass char
     // which is required to have size 1.
     buffer_dev.reset(static_cast<void *>(
-      allocate_device_vector<char>(BufferSize / sizeof(char))));
+      Utilities::CUDA::allocate_device_data<char>(BufferSize / sizeof(char))));
 
     // step 4: perform analysis of incomplete Cholesky on M
     //         perform analysis of triangular solve on L

diff --git a/source/lac/cuda_sparse_matrix.cu b/source/lac/cuda_sparse_matrix.cu
index c409b68ff7e1cc9a9bb4b2f615721b9199ee50ae..07b5cae62543a8169c42fd16d89aaff94fce2db9 100644 (file)
--- a/source/lac/cuda_sparse_matrix.cu
+++ b/source/lac/cuda_sparse_matrix.cu
@@ -27,28 +27,6 @@ DEAL_II_NAMESPACE_OPEN
 
 namespace CUDAWrappers
 {
-  namespace
-  {
-    template <typename Number>
-    void
-    delete_device_data(Number *device_ptr) noexcept
-    {
-      const cudaError_t error_code = cudaFree(device_ptr);
-      (void)error_code;
-      AssertNothrow(error_code == cudaSuccess,
-                    dealii::ExcCudaError(cudaGetErrorString(error_code)));
-    }
-
-    template <typename Number>
-    Number *
-    allocate_device_data(const std::size_t size)
-    {
-      Number *device_ptr;
-      Utilities::CUDA::malloc(device_ptr, size);
-      return device_ptr;
-    }
-  } // namespace
-
   namespace internal
   {
     template <typename Number>
@@ -190,9 +168,9 @@ namespace CUDAWrappers
   SparseMatrix<Number>::SparseMatrix()
     : nnz(0)
     , n_rows(0)
-    , val_dev(nullptr, delete_device_data<Number>)
-    , column_index_dev(nullptr, delete_device_data<int>)
-    , row_ptr_dev(nullptr, delete_device_data<int>)
+    , val_dev(nullptr, Utilities::CUDA::delete_device_data<Number>)
+    , column_index_dev(nullptr, Utilities::CUDA::delete_device_data<int>)
+    , row_ptr_dev(nullptr, Utilities::CUDA::delete_device_data<int>)
     , descr(nullptr)
   {}
 
@@ -202,9 +180,9 @@ namespace CUDAWrappers
   SparseMatrix<Number>::SparseMatrix(
     Utilities::CUDA::Handle &             handle,
     const ::dealii::SparseMatrix<Number> &sparse_matrix_host)
-    : val_dev(nullptr, delete_device_data<Number>)
-    , column_index_dev(nullptr, delete_device_data<int>)
-    , row_ptr_dev(nullptr, delete_device_data<int>)
+    : val_dev(nullptr, Utilities::CUDA::delete_device_data<Number>)
+    , column_index_dev(nullptr, Utilities::CUDA::delete_device_data<int>)
+    , row_ptr_dev(nullptr, Utilities::CUDA::delete_device_data<int>)
     , descr(nullptr)
   {
     reinit(handle, sparse_matrix_host);
@@ -318,7 +296,7 @@ namespace CUDAWrappers
       }
 
     // Copy the elements to the gpu
-    val_dev.reset(allocate_device_data<Number>(nnz));
+    val_dev.reset(Utilities::CUDA::allocate_device_data<Number>(nnz));
     cudaError_t error_code = cudaMemcpy(val_dev.get(),
                                         &val[0],
                                         nnz * sizeof(Number),
@@ -326,7 +304,7 @@ namespace CUDAWrappers
     AssertCuda(error_code);
 
     // Copy the column indices to the gpu
-    column_index_dev.reset(allocate_device_data<int>(nnz));
+    column_index_dev.reset(Utilities::CUDA::allocate_device_data<int>(nnz));
     AssertCuda(error_code);
     error_code = cudaMemcpy(column_index_dev.get(),
                             &column_index[0],
@@ -335,7 +313,7 @@ namespace CUDAWrappers
     AssertCuda(error_code);
 
     // Copy the row pointer to the gpu
-    row_ptr_dev.reset(allocate_device_data<int>(row_ptr_size));
+    row_ptr_dev.reset(Utilities::CUDA::allocate_device_data<int>(row_ptr_size));
     AssertCuda(error_code);
     error_code = cudaMemcpy(row_ptr_dev.get(),
                             &row_ptr[0],

diff --git a/source/lac/cuda_vector.cu b/source/lac/cuda_vector.cu
index 0cb54b675b05559b700a3496320d8272eadb3432..4f12a49740dbe8185066c28798a95ba8253bffad 100644 (file)
--- a/source/lac/cuda_vector.cu
+++ b/source/lac/cuda_vector.cu
@@ -35,33 +35,11 @@ namespace LinearAlgebra
     using ::dealii::CUDAWrappers::block_size;
     using ::dealii::CUDAWrappers::chunk_size;
 
-    namespace
-    {
-      template <typename Number>
-      void
-      delete_device_vector(Number *device_ptr) noexcept
-      {
-        const cudaError_t error_code = cudaFree(device_ptr);
-        (void)error_code;
-        AssertNothrow(error_code == cudaSuccess,
-                      dealii::ExcCudaError(cudaGetErrorString(error_code)));
-      }
-
-      template <typename Number>
-      Number *
-      allocate_device_vector(const std::size_t size)
-      {
-        Number *device_ptr;
-        Utilities::CUDA::malloc(device_ptr, size);
-        return device_ptr;
-      }
-    } // namespace
-
 
 
     template <typename Number>
     Vector<Number>::Vector()
-      : val(nullptr, delete_device_vector<Number>)
+      : val(nullptr, Utilities::CUDA::delete_device_data<Number>)
       , n_elements(0)
     {}
 
@@ -69,8 +47,8 @@ namespace LinearAlgebra
 
     template <typename Number>
     Vector<Number>::Vector(const Vector<Number> &V)
-      : val(allocate_device_vector<Number>(V.n_elements),
-            delete_device_vector<Number>)
+      : val(Utilities::CUDA::allocate_device_data<Number>(V.n_elements),
+            Utilities::CUDA::delete_device_data<Number>)
       , n_elements(V.n_elements)
     {
       // Copy the values.
@@ -106,7 +84,7 @@ namespace LinearAlgebra
 
     template <typename Number>
     Vector<Number>::Vector(const size_type n)
-      : val(nullptr, delete_device_vector<Number>)
+      : val(nullptr, Utilities::CUDA::delete_device_data<Number>)
       , n_elements(0)
     {
       reinit(n, false);
@@ -122,7 +100,7 @@ namespace LinearAlgebra
       if (n == 0)
         val.reset();
       else if (n != n_elements)
-        val.reset(allocate_device_vector<Number>(n));
+        val.reset(Utilities::CUDA::allocate_device_data<Number>(n));
 
       // If necessary set the elements to zero
       if (omit_zeroing_entries == false)