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 = 24 atomicCAS(address_as_ull, assumed, 25 __double_as_longlong(val + 26 __longlong_as_double(assumed))); 27 // Note: uses integer comparison to avoid hang in case of NaN 28 // (since NaN != NaN) 29 } while (assumed != old); 30 return __longlong_as_double(old); 31 } 32 33 //------------------------------------------------------------------------------ 34 35 #endif 36