From: Daniel Arndt Date: Mon, 3 Sep 2018 15:52:59 +0000 (+0200) Subject: Unify allocate/delete_device_data in base/cuda.h X-Git-Tag: v9.1.0-rc1~754^2~2 X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=717351d59af26e72a2039ba823ba181584f58acd;p=dealii.git 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 9a4d57dd47..e701af54e4 100644 --- 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 + 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 + 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 0e8c7db8d3..01343ce43b 100644 --- a/source/lac/cuda_precondition.cu +++ b/source/lac/cuda_precondition.cu @@ -1184,27 +1184,6 @@ namespace pBufferSizeInBytes); } */ - - - - template - 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 - 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 PreconditionIC::PreconditionIC(const Utilities::CUDA::Handle &handle) : cusparse_handle(handle.cusparse_handle) - , P_val_dev(nullptr, delete_device_vector) + , P_val_dev(nullptr, Utilities::CUDA::delete_device_data) , P_row_ptr_dev(nullptr) , P_column_index_dev(nullptr) - , tmp_dev(nullptr, delete_device_vector) - , buffer_dev(nullptr, delete_device_vector) + , tmp_dev(nullptr, Utilities::CUDA::delete_device_data) + , buffer_dev(nullptr, Utilities::CUDA::delete_device_data) , 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(n_nonzero_elements)); + P_val_dev.reset( + Utilities::CUDA::allocate_device_data(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(n_rows)); + tmp_dev.reset(Utilities::CUDA::allocate_device_data(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( - allocate_device_vector(BufferSize / sizeof(char)))); + Utilities::CUDA::allocate_device_data(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::PreconditionILU( const Utilities::CUDA::Handle &handle) : cusparse_handle(handle.cusparse_handle) - , P_val_dev(nullptr, delete_device_vector) + , P_val_dev(nullptr, Utilities::CUDA::delete_device_data) , P_row_ptr_dev(nullptr) , P_column_index_dev(nullptr) - , tmp_dev(nullptr, delete_device_vector) - , buffer_dev(nullptr, delete_device_vector) + , tmp_dev(nullptr, Utilities::CUDA::delete_device_data) + , buffer_dev(nullptr, Utilities::CUDA::delete_device_data) , 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(n_nonzero_elements)); + P_val_dev.reset( + Utilities::CUDA::allocate_device_data(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(n_rows)); + tmp_dev.reset(Utilities::CUDA::allocate_device_data(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( - allocate_device_vector(BufferSize / sizeof(char)))); + Utilities::CUDA::allocate_device_data(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 c409b68ff7..07b5cae625 100644 --- 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 - 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 - 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 @@ -190,9 +168,9 @@ namespace CUDAWrappers SparseMatrix::SparseMatrix() : nnz(0) , n_rows(0) - , val_dev(nullptr, delete_device_data) - , column_index_dev(nullptr, delete_device_data) - , row_ptr_dev(nullptr, delete_device_data) + , val_dev(nullptr, Utilities::CUDA::delete_device_data) + , column_index_dev(nullptr, Utilities::CUDA::delete_device_data) + , row_ptr_dev(nullptr, Utilities::CUDA::delete_device_data) , descr(nullptr) {} @@ -202,9 +180,9 @@ namespace CUDAWrappers SparseMatrix::SparseMatrix( Utilities::CUDA::Handle & handle, const ::dealii::SparseMatrix &sparse_matrix_host) - : val_dev(nullptr, delete_device_data) - , column_index_dev(nullptr, delete_device_data) - , row_ptr_dev(nullptr, delete_device_data) + : val_dev(nullptr, Utilities::CUDA::delete_device_data) + , column_index_dev(nullptr, Utilities::CUDA::delete_device_data) + , row_ptr_dev(nullptr, Utilities::CUDA::delete_device_data) , 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(nnz)); + val_dev.reset(Utilities::CUDA::allocate_device_data(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(nnz)); + column_index_dev.reset(Utilities::CUDA::allocate_device_data(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(row_ptr_size)); + row_ptr_dev.reset(Utilities::CUDA::allocate_device_data(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 0cb54b675b..4f12a49740 100644 --- 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 - 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 - Number * - allocate_device_vector(const std::size_t size) - { - Number *device_ptr; - Utilities::CUDA::malloc(device_ptr, size); - return device_ptr; - } - } // namespace - template Vector::Vector() - : val(nullptr, delete_device_vector) + : val(nullptr, Utilities::CUDA::delete_device_data) , n_elements(0) {} @@ -69,8 +47,8 @@ namespace LinearAlgebra template Vector::Vector(const Vector &V) - : val(allocate_device_vector(V.n_elements), - delete_device_vector) + : val(Utilities::CUDA::allocate_device_data(V.n_elements), + Utilities::CUDA::delete_device_data) , n_elements(V.n_elements) { // Copy the values. @@ -106,7 +84,7 @@ namespace LinearAlgebra template Vector::Vector(const size_type n) - : val(nullptr, delete_device_vector) + : val(nullptr, Utilities::CUDA::delete_device_data) , 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(n)); + val.reset(Utilities::CUDA::allocate_device_data(n)); // If necessary set the elements to zero if (omit_zeroing_entries == false)