diff options
author | Igor Babuschkin <igor@babuschk.in> | 2016-08-05 14:29:58 +0100 |
---|---|---|
committer | Igor Babuschkin <igor@babuschk.in> | 2016-08-05 14:29:58 +0100 |
commit | 9537e8b118062191b0b3f7ac58fb237dbbc6587b (patch) | |
tree | 1eb06a2d9b4c02cffcc4d04237f85207de139a2c | |
parent | eeb0d880ee6dac1b28ac820566017f641f45fcf9 (diff) |
Make use of atomicExch for atomicExchCustom
-rw-r--r-- | unsupported/Eigen/CXX11/src/Tensor/TensorReductionCuda.h | 12 |
1 files changed, 3 insertions, 9 deletions
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorReductionCuda.h b/unsupported/Eigen/CXX11/src/Tensor/TensorReductionCuda.h index 4e2e416e6..c3bdc2783 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorReductionCuda.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorReductionCuda.h @@ -74,15 +74,9 @@ __device__ inline Type atomicExchCustom(Type* address, Type val) { } template <> -__device__ inline double atomicExchCustom<double>(double* address, double val) { - unsigned long long int* address_as_ull = (unsigned long long int*)address; - unsigned long long int old = *address_as_ull; - 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); +__device__ inline double atomicExchCustom(double* address, double val) { + unsigned long long int* address_as_ull = reinterpret_cast<unsigned long long int*>(address); + return __longlong_as_double(atomicExch(address_as_ull, __double_as_longlong(val))); } #ifdef EIGEN_HAS_CUDA_FP16 |