From dd6de618c3fda4275aff3a57c590f82b6e628ac1 Mon Sep 17 00:00:00 2001 From: nluehr Date: Tue, 21 Nov 2017 10:47:00 -0800 Subject: Fix incorrect integer cast in predux(). Bug corrupts results on Maxwell and earlier GPU architectures. --- Eigen/src/Core/arch/CUDA/PacketMathHalf.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) (limited to 'Eigen/src/Core/arch/CUDA') 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(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(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 } -- cgit v1.2.3