1 // Copyright (c) 2017-2024, 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 11 #include <ceed.h> 12 13 //------------------------------------------------------------------------------ 14 // Atomic add, for older CUDA 15 //------------------------------------------------------------------------------ 16 __device__ CeedScalar atomicAdd(CeedScalar *address, CeedScalar val) { 17 unsigned long long int *address_as_ull = (unsigned long long int *)address; 18 unsigned long long int old = *address_as_ull, assumed; 19 do { 20 assumed = old; 21 old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed))); 22 // Note: uses integer comparison to avoid hang in case of NaN 23 // (since NaN != NaN) 24 } while (assumed != old); 25 return __longlong_as_double(old); 26 } 27