Hendrik/xgboost
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Restore the custom double atomic add. (#7198)

Browse Source
This commit is contained in:
Jiaming Yuan 2021-08-28 18:30:42 +08:00 committed by GitHub
parent 7a1d67f9cb
commit ba69244a94
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
1 changed files with 21 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

21
src/common/device_helpers.cuh
View File

@ -53,6 +53,27 @@
#endif // defined(XGBOOST_USE_RMM) && XGBOOST_USE_RMM == 1 #endif // defined(XGBOOST_USE_RMM) && XGBOOST_USE_RMM == 1
#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 600 || defined(__clang__)
#else // In device code and CUDA < 600
__device__ __forceinline__ double atomicAdd(double* address, double val) { // NOLINT
unsigned long long int* address_as_ull =
(unsigned long long int*)address; // NOLINT
unsigned long long int old = *address_as_ull, assumed; // NOLINT
do {
assumed = old;
old = atomicCAS(address_as_ull, assumed,
__double_as_longlong(val + __longlong_as_double(assumed)));
// Note: uses integer comparison to avoid hang in case of NaN (since NaN !=
// NaN)
} while (assumed != old);
return __longlong_as_double(old);
}
#endif
namespace dh { namespace dh {
namespace detail { namespace detail {
template <size_t size> template <size_t size>