1*9ba83ac0SJeremy L Thompson // Copyright (c) 2017-2026, Lawrence Livermore National Security, LLC and other CEED contributors. 29e201c85SYohann // All Rights Reserved. See the top-level LICENSE and NOTICE files for details. 39e201c85SYohann // 49e201c85SYohann // SPDX-License-Identifier: BSD-2-Clause 59e201c85SYohann // 69e201c85SYohann // This file is part of CEED: http://github.com/ceed 79e201c85SYohann 89e201c85SYohann /// @file 99e201c85SYohann /// Internal header for CUDA atomic add fallback definition 10c0b5abf0SJeremy L Thompson #include <ceed/types.h> 119e201c85SYohann 129e201c85SYohann //------------------------------------------------------------------------------ 139e201c85SYohann // Atomic add, for older CUDA 149e201c85SYohann //------------------------------------------------------------------------------ 159e201c85SYohann __device__ CeedScalar atomicAdd(CeedScalar *address, CeedScalar val) { 169e201c85SYohann unsigned long long int *address_as_ull = (unsigned long long int *)address; 179e201c85SYohann unsigned long long int old = *address_as_ull, assumed; 189e201c85SYohann do { 199e201c85SYohann assumed = old; 202b730f8bSJeremy L Thompson old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed))); 219e201c85SYohann // Note: uses integer comparison to avoid hang in case of NaN 229e201c85SYohann // (since NaN != NaN) 239e201c85SYohann } while (assumed != old); 249e201c85SYohann return __longlong_as_double(old); 259e201c85SYohann } 26