1 // Copyright (c) 2017-2022, Lawrence Livermore National Security, LLC and other CEED contributors. 2 // All Rights Reserved. See the top-level LICENSE and NOTICE files for details. 3 // 4 // SPDX-License-Identifier: BSD-2-Clause 5 // 6 // This file is part of CEED: http://github.com/ceed 7 8 /// @file 9 /// Internal header for CUDA atomic add fallback definition 10 #ifndef _ceed_cuda_atomic_add_fallback_h 11 #define _ceed_cuda_atomic_add_fallback_h 12 13 #include <ceed/types.h> 14 15 //------------------------------------------------------------------------------ 16 // Atomic add, for older CUDA 17 //------------------------------------------------------------------------------ 18 __device__ CeedScalar atomicAdd(CeedScalar *address, CeedScalar val) { 19 unsigned long long int *address_as_ull = (unsigned long long int *)address; 20 unsigned long long int old = *address_as_ull, assumed; 21 do { 22 assumed = old; 23 old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed))); 24 // Note: uses integer comparison to avoid hang in case of NaN 25 // (since NaN != NaN) 26 } while (assumed != old); 27 return __longlong_as_double(old); 28 } 29 30 //------------------------------------------------------------------------------ 31 32 #endif 33