From: Timo Heister <timo.heister@gmail.com>
Date: Thu, 21 May 2020 22:44:09 +0000 (-0400)
Subject: clean up cuda_atomic
X-Git-Tag: v9.3.0-rc1~1589^2~2
X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=9c8bac67a7be9ba9ddcaac6775f6be406ee59c22;p=dealii.git

clean up cuda_atomic
---

diff --git a/include/deal.II/lac/cuda_atomic.h b/include/deal.II/lac/cuda_atomic.h
index 7a6cc60496..5bfb113508 100644
--- a/include/deal.II/lac/cuda_atomic.h
+++ b/include/deal.II/lac/cuda_atomic.h
@@ -47,27 +47,7 @@ namespace LinearAlgebra
     inline __device__ double
     atomicAdd_wrapper(double *address, double val)
     {
-      // Use native instruction for CUDA 8 on Pascal or newer architecture
-#  if __CUDACC_VER_MAJOR__ >= 8 && \
-    (!defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 600)
       return atomicAdd(address, val);
-#  else
-
-      unsigned long long int *address_as_ull =
-        reinterpret_cast<unsigned long long int *>(address);
-      unsigned long long int old = *address_as_ull, assumed;
-      do
-        {
-          assumed = old;
-          old     = atomicCAS(address_as_ull,
-                          assumed,
-                          __double_as_longlong(val +
-                                               __longlong_as_double(assumed)));
-        }
-      while (assumed != old);
-
-      return __longlong_as_double(old);
-#  endif
     }