19e201c85SYohann // Copyright (c) 2017-2022, 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 10*94b7b29bSJeremy L Thompson #ifndef CEED_CUDA_ATOMIC_ADD_FALLBACK_H 11*94b7b29bSJeremy L Thompson #define CEED_CUDA_ATOMIC_ADD_FALLBACK_H 129e201c85SYohann 139e201c85SYohann #include <ceed/types.h> 149e201c85SYohann 159e201c85SYohann //------------------------------------------------------------------------------ 169e201c85SYohann // Atomic add, for older CUDA 179e201c85SYohann //------------------------------------------------------------------------------ 189e201c85SYohann __device__ CeedScalar atomicAdd(CeedScalar *address, CeedScalar val) { 199e201c85SYohann unsigned long long int *address_as_ull = (unsigned long long int *)address; 209e201c85SYohann unsigned long long int old = *address_as_ull, assumed; 219e201c85SYohann do { 229e201c85SYohann assumed = old; 232b730f8bSJeremy L Thompson old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed))); 249e201c85SYohann // Note: uses integer comparison to avoid hang in case of NaN 259e201c85SYohann // (since NaN != NaN) 269e201c85SYohann } while (assumed != old); 279e201c85SYohann return __longlong_as_double(old); 289e201c85SYohann } 299e201c85SYohann 309e201c85SYohann //------------------------------------------------------------------------------ 319e201c85SYohann 32*94b7b29bSJeremy L Thompson #endif // CEED_CUDA_ATOMIC_ADD_FALLBACK_H 33