From 048c4d6efd34ae26cebf5a6b662d4480dfe61f0e Mon Sep 17 00:00:00 2001 From: Benoit Steiner Date: Fri, 11 Mar 2016 17:21:42 -0800 Subject: Made half floats usable on hardware that doesn't support them natively. --- Eigen/src/Core/arch/CUDA/PacketMathHalf.h | 46 ------------------------------- 1 file changed, 46 deletions(-) (limited to 'Eigen/src/Core/arch/CUDA/PacketMathHalf.h') diff --git a/Eigen/src/Core/arch/CUDA/PacketMathHalf.h b/Eigen/src/Core/arch/CUDA/PacketMathHalf.h index 9e3c51d49..9e1d87062 100644 --- a/Eigen/src/Core/arch/CUDA/PacketMathHalf.h +++ b/Eigen/src/Core/arch/CUDA/PacketMathHalf.h @@ -19,55 +19,9 @@ #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 300 -// The following operations require arch >= 5.3 -#if __CUDA_ARCH__ >= 530 -__device__ half operator + (const half& a, const half& b) { - return __hadd(a, b); -} -__device__ half operator * (const half& a, const half& b) { - return __hmul(a, b); -} -__device__ half operator - (const half& a, const half& b) { - return __hsub(a, b); -} -__device__ half operator / (const half& a, const half& b) { - float num = __half2float(a); - float denom = __half2float(b); - return __float2half(num / denom); -} -__device__ half operator - (const half& a) { - return __hneg(a); -} -__device__ half& operator += (half& a, const half& b) { - a = a + b; - return a; -} -__device__ half& operator *= (half& a, const half& b) { - a = a * b; - return a; -} -__device__ half& operator -= (half& a, const half& b) { - a = a - b; - return a; -} -__device__ half& operator /= (half& a, const half& b) { - a = a / b; - return a; -} - -namespace std { -__device__ half abs(const half& a) { - half result; - result.x = a.x & 0x7FFF; - return result; -} -} -#endif - namespace Eigen { namespace internal { -template<> struct is_arithmetic { enum { value = true }; }; template<> struct is_arithmetic { enum { value = true }; }; template<> struct packet_traits : default_packet_traits -- cgit v1.2.3