From 9c8bac67a7be9ba9ddcaac6775f6be406ee59c22 Mon Sep 17 00:00:00 2001 From: Timo Heister Date: Thu, 21 May 2020 18:44:09 -0400 Subject: [PATCH] clean up cuda_atomic --- include/deal.II/lac/cuda_atomic.h | 20 -------------------- 1 file changed, 20 deletions(-) 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(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 } -- 2.39.5