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