diff options
author | nluehr <nluehr@nvidia.com> | 2017-11-21 10:47:00 -0800 |
---|---|---|
committer | nluehr <nluehr@nvidia.com> | 2017-11-21 10:47:00 -0800 |
commit | dd6de618c3fda4275aff3a57c590f82b6e628ac1 (patch) | |
tree | 1399b7172c8dcae0c5de08389e30a7750cddc8a8 /Eigen/src/Core/arch/CUDA | |
parent | 3dc6ff73cae64da97244e493905de7be40d1bcde (diff) |
Fix incorrect integer cast in predux<half2>().
Bug corrupts results on Maxwell and earlier GPU architectures.
Diffstat (limited to 'Eigen/src/Core/arch/CUDA')
-rw-r--r-- | Eigen/src/Core/arch/CUDA/PacketMathHalf.h | 4 |
1 files changed, 2 insertions, 2 deletions
diff --git a/Eigen/src/Core/arch/CUDA/PacketMathHalf.h b/Eigen/src/Core/arch/CUDA/PacketMathHalf.h index ce48e4b31..e3e45b03a 100644 --- a/Eigen/src/Core/arch/CUDA/PacketMathHalf.h +++ b/Eigen/src/Core/arch/CUDA/PacketMathHalf.h @@ -231,7 +231,7 @@ template<> __device__ EIGEN_STRONG_INLINE Eigen::half predux<half2>(const half2& #else float a1 = __low2float(a); float a2 = __high2float(a); - return Eigen::half(half_impl::raw_uint16_to_half(__float2half_rn(a1 + a2))); + return Eigen::half(__float2half_rn(a1 + a2)); #endif } @@ -265,7 +265,7 @@ template<> __device__ EIGEN_STRONG_INLINE Eigen::half predux_mul<half2>(const ha #else float a1 = __low2float(a); float a2 = __high2float(a); - return Eigen::half(half_impl::raw_uint16_to_half(__float2half_rn(a1 * a2))); + return Eigen::half(__float2half_rn(a1 * a2)); #endif } |