aboutsummaryrefslogtreecommitdiffhomepage
path: root/Eigen/src/Core/arch/CUDA
diff options
context:
space:
mode:
authorGravatar nluehr <nluehr@nvidia.com>2017-11-21 10:47:00 -0800
committerGravatar nluehr <nluehr@nvidia.com>2017-11-21 10:47:00 -0800
commitdd6de618c3fda4275aff3a57c590f82b6e628ac1 (patch)
tree1399b7172c8dcae0c5de08389e30a7750cddc8a8 /Eigen/src/Core/arch/CUDA
parent3dc6ff73cae64da97244e493905de7be40d1bcde (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.h4
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
}