cuda: deprecate atomicAdd_wrapper

diff --git a/include/deal.II/lac/cuda_atomic.h b/include/deal.II/lac/cuda_atomic.h
index 5bfb1135086078b0256d1440fab0ff53a2bbc29c..368e1b3c0b9ebeaa53ddd8cc1bab235acae95e6c 100644 (file)
--- 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 b9a478845e91eaf4606b1aa508f780f98170964b..856d73da40717c5191fc4f3af538b8fd53f13f4a 100644 (file)
--- 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<Number>::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<Number>::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<Number>::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 44a3d4cb00a3c2d4936fd053fa749dad31b1f751..f42b87409a330d94e80351910278beb39fe26fea 100644 (file)
--- 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 547e29c53ab526a26e91fd3bbc1128c9fd1f42c5..27a08e59f491a3b02785dbdcbbcd817ff20c2453 100644 (file)
--- 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]));
         }
     }