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