From: Timo Heister Date: Sat, 23 May 2020 00:44:33 +0000 (-0400) Subject: cuda: deprecate atomicAdd_wrapper X-Git-Tag: v9.3.0-rc1~1581^2 X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=9b6656ec8dd3a8fdcefdd8fd2eada374a8ceabec;p=dealii.git cuda: deprecate atomicAdd_wrapper --- diff --git a/include/deal.II/lac/cuda_atomic.h b/include/deal.II/lac/cuda_atomic.h index 5bfb113508..368e1b3c0b 100644 --- a/include/deal.II/lac/cuda_atomic.h +++ b/include/deal.II/lac/cuda_atomic.h @@ -29,8 +29,10 @@ namespace LinearAlgebra /** * Provide atomicAdd for floats. * + * @deprecated Use atomicAdd(address, val) directly. * @ingroup CUDAWrappers */ + DEAL_II_DEPRECATED inline __device__ float atomicAdd_wrapper(float *address, float val) { @@ -42,8 +44,10 @@ namespace LinearAlgebra /** * Provide atomicAdd for doubles. * + * @deprecated Use atomicAdd(address, val) directly. * @ingroup CUDAWrappers */ + DEAL_II_DEPRECATED inline __device__ double atomicAdd_wrapper(double *address, double val) { diff --git a/include/deal.II/lac/cuda_kernels.templates.h b/include/deal.II/lac/cuda_kernels.templates.h index b9a478845e..856d73da40 100644 --- a/include/deal.II/lac/cuda_kernels.templates.h +++ b/include/deal.II/lac/cuda_kernels.templates.h @@ -93,7 +93,7 @@ namespace LinearAlgebra __device__ Number ElemSum::atomic_op(Number *dst, const Number a) { - return atomicAdd_wrapper(dst, a); + return atomicAdd(dst, a); } @@ -129,7 +129,7 @@ namespace LinearAlgebra __device__ Number L1Norm::atomic_op(Number *dst, const Number a) { - return atomicAdd_wrapper(dst, a); + return atomicAdd(dst, a); } @@ -269,7 +269,7 @@ namespace LinearAlgebra __device__ Number DotProduct::atomic_op(Number *dst, const Number a) { - return atomicAdd_wrapper(dst, a); + return atomicAdd(dst, a); } diff --git a/include/deal.II/matrix_free/cuda_fe_evaluation.h b/include/deal.II/matrix_free/cuda_fe_evaluation.h index 44a3d4cb00..f42b87409a 100644 --- a/include/deal.II/matrix_free/cuda_fe_evaluation.h +++ b/include/deal.II/matrix_free/cuda_fe_evaluation.h @@ -403,8 +403,7 @@ namespace CUDAWrappers if (use_coloring) dst[destination_idx] += values[idx]; else - LinearAlgebra::CUDAWrappers::atomicAdd_wrapper(&dst[destination_idx], - values[idx]); + atomicAdd(&dst[destination_idx], values[idx]); } diff --git a/source/lac/cuda_sparse_matrix.cu b/source/lac/cuda_sparse_matrix.cu index 547e29c53a..27a08e59f4 100644 --- a/source/lac/cuda_sparse_matrix.cu +++ b/source/lac/cuda_sparse_matrix.cu @@ -135,8 +135,7 @@ namespace CUDAWrappers if (row < n_rows) { for (int j = row_ptr_dev[row]; j < row_ptr_dev[row + 1]; ++j) - dealii::LinearAlgebra::CUDAWrappers::atomicAdd_wrapper( - &sums[column_index_dev[j]], abs(val_dev[j])); + atomicAdd(&sums[column_index_dev[j]], abs(val_dev[j])); } }