15aed82e4SJeremy L Thompson // Copyright (c) 2017-2024, 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 109e201c85SYohann 11*c0b5abf0SJeremy L Thompson #include <ceed/types.h> 129e201c85SYohann 139e201c85SYohann //------------------------------------------------------------------------------ 149e201c85SYohann // Atomic add, for older CUDA 159e201c85SYohann //------------------------------------------------------------------------------ 169e201c85SYohann __device__ CeedScalar atomicAdd(CeedScalar *address, CeedScalar val) { 179e201c85SYohann unsigned long long int *address_as_ull = (unsigned long long int *)address; 189e201c85SYohann unsigned long long int old = *address_as_ull, assumed; 199e201c85SYohann do { 209e201c85SYohann assumed = old; 212b730f8bSJeremy L Thompson old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed))); 229e201c85SYohann // Note: uses integer comparison to avoid hang in case of NaN 239e201c85SYohann // (since NaN != NaN) 249e201c85SYohann } while (assumed != old); 259e201c85SYohann return __longlong_as_double(old); 269e201c85SYohann } 27