From: Timo Heister <timo.heister@gmail.com>
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<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 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]));
         }
     }