From 4707c3aa86fb3d429eec1b85e5415e1a01609f5a Mon Sep 17 00:00:00 2001 From: nluehr Date: Tue, 21 Nov 2017 10:47:00 -0800 Subject: [PATCH] Fix incorrect integer cast in predux(). Bug corrupts results on Maxwell and earlier GPU architectures. (cherry picked from commit dd6de618c3fda4275aff3a57c590f82b6e628ac1) --- Eigen/src/Core/arch/CUDA/PacketMathHalf.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/Eigen/src/Core/arch/CUDA/PacketMathHalf.h b/Eigen/src/Core/arch/CUDA/PacketMathHalf.h index c66d38469..f749c573f 100644 --- a/Eigen/src/Core/arch/CUDA/PacketMathHalf.h +++ b/Eigen/src/Core/arch/CUDA/PacketMathHalf.h @@ -230,7 +230,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 } @@ -264,7 +264,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 }