mirror of
https://gitlab.com/libeigen/eigen.git
synced 2025-08-12 19:59:05 +08:00
Make use of atomicExch for atomicExchCustom
This commit is contained in:
parent
eeb0d880ee
commit
9537e8b118
@ -74,15 +74,9 @@ __device__ inline Type atomicExchCustom(Type* address, Type val) {
|
|||||||
}
|
}
|
||||||
|
|
||||||
template <>
|
template <>
|
||||||
__device__ inline double atomicExchCustom<double>(double* address, double val) {
|
__device__ inline double atomicExchCustom(double* address, double val) {
|
||||||
unsigned long long int* address_as_ull = (unsigned long long int*)address;
|
unsigned long long int* address_as_ull = reinterpret_cast<unsigned long long int*>(address);
|
||||||
unsigned long long int old = *address_as_ull;
|
return __longlong_as_double(atomicExch(address_as_ull, __double_as_longlong(val)));
|
||||||
unsigned long long int assumed;
|
|
||||||
do {
|
|
||||||
assumed = old;
|
|
||||||
old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val));
|
|
||||||
} while (assumed != old);
|
|
||||||
return __longlong_as_double(old);
|
|
||||||
}
|
}
|
||||||
|
|
||||||
#ifdef EIGEN_HAS_CUDA_FP16
|
#ifdef EIGEN_HAS_CUDA_FP16
|
||||||
|
Loading…
x
Reference in New Issue
Block a user