aboutsummaryrefslogtreecommitdiffhomepage
diff options
context:
space:
mode:
Diffstat
-rw-r--r--Eigen/src/Core/arch/GPU/PacketMath.h895
-rw-r--r--Eigen/src/Core/arch/GPU/TypeCasting.h33
-rw-r--r--unsupported/Eigen/CXX11/src/Tensor/TensorMeta.h8
-rw-r--r--unsupported/Eigen/CXX11/src/Tensor/TensorReduction.h8
-rw-r--r--unsupported/Eigen/CXX11/src/Tensor/TensorReductionGpu.h290
5 files changed, 1068 insertions, 166 deletions
diff --git a/Eigen/src/Core/arch/GPU/PacketMath.h b/Eigen/src/Core/arch/GPU/PacketMath.h
index 9e18c5145..1f6a562c5 100644
--- a/Eigen/src/Core/arch/GPU/PacketMath.h
+++ b/Eigen/src/Core/arch/GPU/PacketMath.h
@@ -476,24 +476,29 @@ ptranspose(PacketBlock<double2,2>& kernel) {
kernel.packet[1].x = tmp;
}
-#endif
+#endif // defined(EIGEN_GPUCC) && defined(EIGEN_USE_GPU)
-// Packet math for Eigen::half
-// Most of the following operations require arch >= 3.0
-#if (defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDACC) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 300) || \
+// Packet4h2 must be defined in the macro without EIGEN_CUDA_ARCH, meaning
+// its corresponding packet_traits<Eigen::half> must be visible on host.
+#if (defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDACC)) || \
(defined(EIGEN_HAS_HIP_FP16) && defined(EIGEN_HIPCC) && defined(EIGEN_HIP_DEVICE_COMPILE)) || \
(defined(EIGEN_HAS_CUDA_FP16) && defined(__clang__) && defined(__CUDA__))
+typedef ulonglong2 Packet4h2;
+template<> struct unpacket_traits<Packet4h2> { typedef Eigen::half type; enum {size=8, alignment=Aligned16, vectorizable=true, masked_load_available=false, masked_store_available=false}; typedef Packet4h2 half; };
+template<> struct is_arithmetic<Packet4h2> { enum { value = true }; };
+
+template<> struct unpacket_traits<half2> { typedef Eigen::half type; enum {size=2, alignment=Aligned16, vectorizable=true, masked_load_available=false, masked_store_available=false}; typedef half2 half; };
template<> struct is_arithmetic<half2> { enum { value = true }; };
template<> struct packet_traits<Eigen::half> : default_packet_traits
{
- typedef half2 type;
- typedef half2 half;
+ typedef Packet4h2 type;
+ typedef Packet4h2 half;
enum {
Vectorizable = 1,
AlignedOnScalar = 1,
- size=2,
+ size=8,
HasHalfPacket = 0,
HasAdd = 1,
HasSub = 1,
@@ -508,9 +513,8 @@ template<> struct packet_traits<Eigen::half> : default_packet_traits
};
};
-template<> struct unpacket_traits<half2> { typedef Eigen::half type; enum {size=2, alignment=Aligned16, vectorizable=true, masked_load_available=false, masked_store_available=false}; typedef half2 half; };
-
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pset1<half2>(const Eigen::half& from) {
+template<>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pset1<half2>(const Eigen::half& from) {
#if !defined(EIGEN_CUDA_ARCH) && !defined(EIGEN_HIP_DEVICE_COMPILE)
half2 r;
r.x = from;
@@ -521,23 +525,40 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pset1<half2>(const Eigen:
#endif
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pload<half2>(const Eigen::half* from) {
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2
+pset1<Packet4h2>(const Eigen::half& from) {
+ Packet4h2 r;
+ half2* p_alias = reinterpret_cast<half2*>(&r);
+ p_alias[0] = pset1<half2>(from);
+ p_alias[1] = pset1<half2>(from);
+ p_alias[2] = pset1<half2>(from);
+ p_alias[3] = pset1<half2>(from);
+ return r;
+}
+
+#if defined(EIGEN_CUDA_ARCH) || defined(EIGEN_HIP_DEVICE_COMPILE) || (defined(EIGEN_CUDACC) && EIGEN_COMP_CLANG && !EIGEN_COMP_NVCC)
+namespace {
+
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pload(const Eigen::half* from) {
return *reinterpret_cast<const half2*>(from);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 ploadu<half2>(const Eigen::half* from) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 ploadu(const Eigen::half* from) {
return __halves2half2(from[0], from[1]);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 ploaddup<half2>(const Eigen::half* from) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 ploaddup(const Eigen::half* from) {
return __halves2half2(from[0], from[0]);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstore<Eigen::half>(Eigen::half* to, const half2& from) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstore(Eigen::half* to,
+ const half2& from) {
*reinterpret_cast<half2*>(to) = from;
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstoreu<Eigen::half>(Eigen::half* to, const half2& from) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstoreu(Eigen::half* to,
+ const half2& from) {
#if !defined(EIGEN_CUDA_ARCH) && !defined(EIGEN_HIP_DEVICE_COMPILE)
to[0] = from.x;
to[1] = from.y;
@@ -547,8 +568,9 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstoreu<Eigen::half>(Eigen
#endif
}
-template<>
- EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE half2 ploadt_ro<half2, Aligned>(const Eigen::half* from) {
+
+EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE half2 ploadt_ro_aligned(
+ const Eigen::half* from) {
#if defined(EIGEN_HIP_DEVICE_COMPILE)
@@ -565,8 +587,8 @@ template<>
#endif
}
-template<>
-EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE half2 ploadt_ro<half2, Unaligned>(const Eigen::half* from) {
+EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE half2 ploadt_ro_unaligned(
+ const Eigen::half* from) {
#if defined(EIGEN_HIP_DEVICE_COMPILE)
@@ -583,20 +605,22 @@ EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE half2 ploadt_ro<half2, Unaligned>(const Ei
#endif
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pgather<Eigen::half, half2>(const Eigen::half* from, Index stride) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pgather(const Eigen::half* from,
+ Index stride) {
return __halves2half2(from[0*stride], from[1*stride]);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pscatter<Eigen::half, half2>(Eigen::half* to, const half2& from, Index stride) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pscatter(
+ Eigen::half* to, const half2& from, Index stride) {
to[stride*0] = __low2half(from);
to[stride*1] = __high2half(from);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half pfirst<half2>(const half2& a) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half pfirst(const half2& a) {
return __low2half(a);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pabs<half2>(const half2& a) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pabs(const half2& a) {
half a1 = __low2half(a);
half a2 = __high2half(a);
half result1 = half_impl::raw_uint16_to_half(a1.x & 0x7FFF);
@@ -604,12 +628,12 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pabs<half2>(const half2&
return __halves2half2(result1, result2);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 ptrue<half2>(const half2& a) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 ptrue(const half2& a) {
half true_half = half_impl::raw_uint16_to_half(0xffffu);
return pset1<half2>(true_half);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pzero<half2>(const half2& a) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pzero(const half2& a) {
half false_half = half_impl::raw_uint16_to_half(0x0000u);
return pset1<half2>(false_half);
}
@@ -624,7 +648,7 @@ ptranspose(PacketBlock<half2,2>& kernel) {
kernel.packet[1] = __halves2half2(a2, b2);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 plset<half2>(const Eigen::half& a) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 plset(const Eigen::half& a) {
#if defined(EIGEN_HIP_DEVICE_COMPILE)
return __halves2half2(a, __hadd(a, __float2half(1.0f)));
@@ -641,10 +665,9 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 plset<half2>(const Eigen:
#endif
}
-template <>
-EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pselect<half2>(const half2& mask,
- const half2& a,
- const half2& b) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pselect(const half2& mask,
+ const half2& a,
+ const half2& b) {
half mask_low = __low2half(mask);
half mask_high = __high2half(mask);
half result_low = mask_low == half(0) ? __low2half(b) : __low2half(a);
@@ -652,9 +675,8 @@ EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pselect<half2>(const half2& mask,
return __halves2half2(result_low, result_high);
}
-template <>
-EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pcmp_eq<half2>(const half2& a,
- const half2& b) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pcmp_eq(const half2& a,
+ const half2& b) {
half true_half = half_impl::raw_uint16_to_half(0xffffu);
half false_half = half_impl::raw_uint16_to_half(0x0000u);
half a1 = __low2half(a);
@@ -666,9 +688,8 @@ EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pcmp_eq<half2>(const half2& a,
return __halves2half2(eq1, eq2);
}
-template <>
-EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pcmp_lt<half2>(const half2& a,
- const half2& b) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pcmp_lt(const half2& a,
+ const half2& b) {
half true_half = half_impl::raw_uint16_to_half(0xffffu);
half false_half = half_impl::raw_uint16_to_half(0x0000u);
half a1 = __low2half(a);
@@ -680,9 +701,8 @@ EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pcmp_lt<half2>(const half2& a,
return __halves2half2(eq1, eq2);
}
-template <>
-EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pand<half2>(const half2& a,
- const half2& b) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pand(const half2& a,
+ const half2& b) {
half a1 = __low2half(a);
half a2 = __high2half(a);
half b1 = __low2half(b);
@@ -692,9 +712,8 @@ EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pand<half2>(const half2& a,
return __halves2half2(result1, result2);
}
-template <>
-EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 por<half2>(const half2& a,
- const half2& b) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 por(const half2& a,
+ const half2& b) {
half a1 = __low2half(a);
half a2 = __high2half(a);
half b1 = __low2half(b);
@@ -704,9 +723,8 @@ EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 por<half2>(const half2& a,
return __halves2half2(result1, result2);
}
-template <>
-EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pxor<half2>(const half2& a,
- const half2& b) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pxor(const half2& a,
+ const half2& b) {
half a1 = __low2half(a);
half a2 = __high2half(a);
half b1 = __low2half(b);
@@ -716,9 +734,8 @@ EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pxor<half2>(const half2& a,
return __halves2half2(result1, result2);
}
-template <>
-EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pandnot<half2>(const half2& a,
- const half2& b) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pandnot(const half2& a,
+ const half2& b) {
half a1 = __low2half(a);
half a2 = __high2half(a);
half b1 = __low2half(b);
@@ -728,7 +745,8 @@ EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pandnot<half2>(const half2& a,
return __halves2half2(result1, result2);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 padd<half2>(const half2& a, const half2& b) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 padd(const half2& a,
+ const half2& b) {
#if defined(EIGEN_HIP_DEVICE_COMPILE)
return __hadd2(a, b);
@@ -750,7 +768,8 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 padd<half2>(const half2&
#endif
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 psub<half2>(const half2& a, const half2& b) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 psub(const half2& a,
+ const half2& b) {
#if defined(EIGEN_HIP_DEVICE_COMPILE)
return __hsub2(a, b);
@@ -772,7 +791,7 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 psub<half2>(const half2&
#endif
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pnegate(const half2& a) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pnegate(const half2& a) {
#if defined(EIGEN_HIP_DEVICE_COMPILE)
return __hneg2(a);
@@ -790,9 +809,10 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pnegate(const half2& a) {
#endif
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pconj(const half2& a) { return a; }
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pconj(const half2& a) { return a; }
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmul<half2>(const half2& a, const half2& b) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmul(const half2& a,
+ const half2& b) {
#if defined(EIGEN_HIP_DEVICE_COMPILE)
return __hmul2(a, b);
@@ -814,7 +834,9 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmul<half2>(const half2&
#endif
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmadd<half2>(const half2& a, const half2& b, const half2& c) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmadd(const half2& a,
+ const half2& b,
+ const half2& c) {
#if defined(EIGEN_HIP_DEVICE_COMPILE)
return __hfma2(a, b, c);
@@ -838,7 +860,8 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmadd<half2>(const half2&
#endif
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pdiv<half2>(const half2& a, const half2& b) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pdiv(const half2& a,
+ const half2& b) {
#if defined(EIGEN_HIP_DEVICE_COMPILE)
return __h2div(a, b);
@@ -856,7 +879,8 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pdiv<half2>(const half2&
#endif
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmin<half2>(const half2& a, const half2& b) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmin(const half2& a,
+ const half2& b) {
float a1 = __low2float(a);
float a2 = __high2float(a);
float b1 = __low2float(b);
@@ -866,7 +890,8 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmin<half2>(const half2&
return __halves2half2(r1, r2);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmax<half2>(const half2& a, const half2& b) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmax(const half2& a,
+ const half2& b) {
float a1 = __low2float(a);
float a2 = __high2float(a);
float b1 = __low2float(b);
@@ -876,7 +901,7 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmax<half2>(const half2&
return __halves2half2(r1, r2);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux<half2>(const half2& a) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux(const half2& a) {
#if defined(EIGEN_HIP_DEVICE_COMPILE)
return __hadd(__low2half(a), __high2half(a));
@@ -894,7 +919,7 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux<half2>(const
#endif
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_max<half2>(const half2& a) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_max(const half2& a) {
#if defined(EIGEN_HIP_DEVICE_COMPILE)
__half first = __low2half(a);
@@ -916,7 +941,7 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_max<half2>(c
#endif
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_min<half2>(const half2& a) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_min(const half2& a) {
#if defined(EIGEN_HIP_DEVICE_COMPILE)
__half first = __low2half(a);
@@ -938,7 +963,7 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_min<half2>(c
#endif
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_mul<half2>(const half2& a) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_mul(const half2& a) {
#if defined(EIGEN_HIP_DEVICE_COMPILE)
return __hmul(__low2half(a), __high2half(a));
@@ -956,7 +981,7 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_mul<half2>(c
#endif
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 plog1p<half2>(const half2& a) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 plog1p(const half2& a) {
float a1 = __low2float(a);
float a2 = __high2float(a);
float r1 = log1pf(a1);
@@ -964,7 +989,7 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 plog1p<half2>(const half2
return __floats2half2_rn(r1, r2);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pexpm1<half2>(const half2& a) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pexpm1(const half2& a) {
float a1 = __low2float(a);
float a2 = __high2float(a);
float r1 = expm1f(a1);
@@ -975,29 +1000,29 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pexpm1<half2>(const half2
#if (EIGEN_CUDA_SDK_VER >= 80000 && defined EIGEN_CUDA_ARCH && EIGEN_CUDA_ARCH >= 530) || \
defined(EIGEN_HIP_DEVICE_COMPILE)
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
-half2 plog<half2>(const half2& a) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
+half2 plog(const half2& a) {
return h2log(a);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
-half2 pexp<half2>(const half2& a) {
+ EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
+half2 pexp(const half2& a) {
return h2exp(a);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
-half2 psqrt<half2>(const half2& a) {
+ EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
+half2 psqrt(const half2& a) {
return h2sqrt(a);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
-half2 prsqrt<half2>(const half2& a) {
+ EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE
+half2 prsqrt(const half2& a) {
return h2rsqrt(a);
}
#else
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 plog<half2>(const half2& a) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 plog(const half2& a) {
float a1 = __low2float(a);
float a2 = __high2float(a);
float r1 = logf(a1);
@@ -1005,7 +1030,7 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 plog<half2>(const half2&
return __floats2half2_rn(r1, r2);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pexp<half2>(const half2& a) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pexp(const half2& a) {
float a1 = __low2float(a);
float a2 = __high2float(a);
float r1 = expf(a1);
@@ -1013,7 +1038,7 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pexp<half2>(const half2&
return __floats2half2_rn(r1, r2);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 psqrt<half2>(const half2& a) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 psqrt(const half2& a) {
float a1 = __low2float(a);
float a2 = __high2float(a);
float r1 = sqrtf(a1);
@@ -1021,7 +1046,7 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 psqrt<half2>(const half2&
return __floats2half2_rn(r1, r2);
}
-template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 prsqrt<half2>(const half2& a) {
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 prsqrt(const half2& a) {
float a1 = __low2float(a);
float a2 = __high2float(a);
float r1 = rsqrtf(a1);
@@ -1029,8 +1054,728 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 prsqrt<half2>(const half2
return __floats2half2_rn(r1, r2);
}
#endif
+} // namespace
+
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2
+pload<Packet4h2>(const Eigen::half* from) {
+ return *reinterpret_cast<const Packet4h2*>(from);
+}
+
+// unaligned load;
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2
+ploadu<Packet4h2>(const Eigen::half* from) {
+ Packet4h2 r;
+ half2* p_alias = reinterpret_cast<half2*>(&r);
+ p_alias[0] = ploadu(from + 0);
+ p_alias[1] = ploadu(from + 2);
+ p_alias[2] = ploadu(from + 4);
+ p_alias[3] = ploadu(from + 6);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2
+ploaddup<Packet4h2>(const Eigen::half* from) {
+ Packet4h2 r;
+ half2* p_alias = reinterpret_cast<half2*>(&r);
+ p_alias[0] = ploaddup(from + 0);
+ p_alias[1] = ploaddup(from + 1);
+ p_alias[2] = ploaddup(from + 2);
+ p_alias[3] = ploaddup(from + 3);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstore<Eigen::half>(
+ Eigen::half* to, const Packet4h2& from) {
+ *reinterpret_cast<Packet4h2*>(to) = from;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pstoreu<Eigen::half>(
+ Eigen::half* to, const Packet4h2& from) {
+ const half2* from_alias = reinterpret_cast<const half2*>(&from);
+ pstoreu(to + 0,from_alias[0]);
+ pstoreu(to + 2,from_alias[1]);
+ pstoreu(to + 4,from_alias[2]);
+ pstoreu(to + 6,from_alias[3]);
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Packet4h2
+ploadt_ro<Packet4h2, Aligned>(const Eigen::half* from) {
+#if defined(EIGEN_HIP_DEVICE_COMPILE)
+
+ Packet4h2 r;
+ r = __ldg((const Packet4h2*)from);
+ return r;
+#else // EIGEN_CUDA_ARCH
+
+#if EIGEN_CUDA_ARCH >= 350
+ Packet4h2 r;
+ r = __ldg((const Packet4h2*)from);
+ return r;
+#else
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ r_alias[0] = ploadt_ro_aligned(from + 0);
+ r_alias[1] = ploadt_ro_aligned(from + 2);
+ r_alias[2] = ploadt_ro_aligned(from + 4);
+ r_alias[3] = ploadt_ro_aligned(from + 6);
+ return r;
+#endif
+
+#endif
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Packet4h2
+ploadt_ro<Packet4h2, Unaligned>(const Eigen::half* from) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ r_alias[0] = ploadt_ro_unaligned(from + 0);
+ r_alias[1] = ploadt_ro_unaligned(from + 2);
+ r_alias[2] = ploadt_ro_unaligned(from + 4);
+ r_alias[3] = ploadt_ro_unaligned(from + 6);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2
+pgather<Eigen::half, Packet4h2>(const Eigen::half* from, Index stride) {
+ Packet4h2 r;
+ half2* p_alias = reinterpret_cast<half2*>(&r);
+ p_alias[0] = __halves2half2(from[0 * stride], from[1 * stride]);
+ p_alias[1] = __halves2half2(from[2 * stride], from[3 * stride]);
+ p_alias[2] = __halves2half2(from[4 * stride], from[5 * stride]);
+ p_alias[3] = __halves2half2(from[6 * stride], from[7 * stride]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pscatter<Eigen::half, Packet4h2>(
+ Eigen::half* to, const Packet4h2& from, Index stride) {
+ const half2* from_alias = reinterpret_cast<const half2*>(&from);
+ pscatter(to + stride * 0, from_alias[0], stride);
+ pscatter(to + stride * 2, from_alias[1], stride);
+ pscatter(to + stride * 4, from_alias[2], stride);
+ pscatter(to + stride * 6, from_alias[3], stride);
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half pfirst<Packet4h2>(
+ const Packet4h2& a) {
+ return pfirst(*(reinterpret_cast<const half2*>(&a)));
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pabs<Packet4h2>(
+ const Packet4h2& a) {
+ Packet4h2 r;
+ half2* p_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ p_alias[0] = pabs(a_alias[0]);
+ p_alias[1] = pabs(a_alias[1]);
+ p_alias[2] = pabs(a_alias[2]);
+ p_alias[3] = pabs(a_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 ptrue<Packet4h2>(
+ const Packet4h2& a) {
+ half true_half = half_impl::raw_uint16_to_half(0xffffu);
+ return pset1<Packet4h2>(true_half);
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pzero<Packet4h2>(const Packet4h2& a) {
+ half false_half = half_impl::raw_uint16_to_half(0x0000u);
+ return pset1<Packet4h2>(false_half);
+}
+
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void ptranspose_double(
+ double* d_row0, double* d_row1, double* d_row2, double* d_row3,
+ double* d_row4, double* d_row5, double* d_row6, double* d_row7) {
+ double d_tmp;
+ d_tmp = d_row0[1];
+ d_row0[1] = d_row4[0];
+ d_row4[0] = d_tmp;
+
+ d_tmp = d_row1[1];
+ d_row1[1] = d_row5[0];
+ d_row5[0] = d_tmp;
+
+ d_tmp = d_row2[1];
+ d_row2[1] = d_row6[0];
+ d_row6[0] = d_tmp;
+
+ d_tmp = d_row3[1];
+ d_row3[1] = d_row7[0];
+ d_row7[0] = d_tmp;
+}
+
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void ptranspose_half2(
+ half2* f_row0, half2* f_row1, half2* f_row2, half2* f_row3) {
+ half2 f_tmp;
+ f_tmp = f_row0[1];
+ f_row0[1] = f_row2[0];
+ f_row2[0] = f_tmp;
+
+ f_tmp = f_row1[1];
+ f_row1[1] = f_row3[0];
+ f_row3[0] = f_tmp;
+}
+
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void
+ptranspose_half(half2& f0, half2& f1) {
+ __half a1 = __low2half(f0);
+ __half a2 = __high2half(f0);
+ __half b1 = __low2half(f1);
+ __half b2 = __high2half(f1);
+ f0 = __halves2half2(a1, b1);
+ f1 = __halves2half2(a2, b2);
+}
+
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void
+ptranspose(PacketBlock<Packet4h2,8>& kernel) {
+ double* d_row0 = reinterpret_cast<double*>(&kernel.packet[0]);
+ double* d_row1 = reinterpret_cast<double*>(&kernel.packet[1]);
+ double* d_row2 = reinterpret_cast<double*>(&kernel.packet[2]);
+ double* d_row3 = reinterpret_cast<double*>(&kernel.packet[3]);
+ double* d_row4 = reinterpret_cast<double*>(&kernel.packet[4]);
+ double* d_row5 = reinterpret_cast<double*>(&kernel.packet[5]);
+ double* d_row6 = reinterpret_cast<double*>(&kernel.packet[6]);
+ double* d_row7 = reinterpret_cast<double*>(&kernel.packet[7]);
+ ptranspose_double(d_row0, d_row1, d_row2, d_row3,
+ d_row4, d_row5, d_row6, d_row7);
+
+
+ half2* f_row0 = reinterpret_cast<half2*>(d_row0);
+ half2* f_row1 = reinterpret_cast<half2*>(d_row1);
+ half2* f_row2 = reinterpret_cast<half2*>(d_row2);
+ half2* f_row3 = reinterpret_cast<half2*>(d_row3);
+ ptranspose_half2(f_row0, f_row1, f_row2, f_row3);
+ ptranspose_half(f_row0[0], f_row1[0]);
+ ptranspose_half(f_row0[1], f_row1[1]);
+ ptranspose_half(f_row2[0], f_row3[0]);
+ ptranspose_half(f_row2[1], f_row3[1]);
+
+ f_row0 = reinterpret_cast<half2*>(d_row0 + 1);
+ f_row1 = reinterpret_cast<half2*>(d_row1 + 1);
+ f_row2 = reinterpret_cast<half2*>(d_row2 + 1);
+ f_row3 = reinterpret_cast<half2*>(d_row3 + 1);
+ ptranspose_half2(f_row0, f_row1, f_row2, f_row3);
+ ptranspose_half(f_row0[0], f_row1[0]);
+ ptranspose_half(f_row0[1], f_row1[1]);
+ ptranspose_half(f_row2[0], f_row3[0]);
+ ptranspose_half(f_row2[1], f_row3[1]);
+
+ f_row0 = reinterpret_cast<half2*>(d_row4);
+ f_row1 = reinterpret_cast<half2*>(d_row5);
+ f_row2 = reinterpret_cast<half2*>(d_row6);
+ f_row3 = reinterpret_cast<half2*>(d_row7);
+ ptranspose_half2(f_row0, f_row1, f_row2, f_row3);
+ ptranspose_half(f_row0[0], f_row1[0]);
+ ptranspose_half(f_row0[1], f_row1[1]);
+ ptranspose_half(f_row2[0], f_row3[0]);
+ ptranspose_half(f_row2[1], f_row3[1]);
+
+ f_row0 = reinterpret_cast<half2*>(d_row4 + 1);
+ f_row1 = reinterpret_cast<half2*>(d_row5 + 1);
+ f_row2 = reinterpret_cast<half2*>(d_row6 + 1);
+ f_row3 = reinterpret_cast<half2*>(d_row7 + 1);
+ ptranspose_half2(f_row0, f_row1, f_row2, f_row3);
+ ptranspose_half(f_row0[0], f_row1[0]);
+ ptranspose_half(f_row0[1], f_row1[1]);
+ ptranspose_half(f_row2[0], f_row3[0]);
+ ptranspose_half(f_row2[1], f_row3[1]);
+
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2
+plset<Packet4h2>(const Eigen::half& a) {
+#if defined(EIGEN_HIP_DEVICE_COMPILE)
+
+ Packet4h2 r;
+ half2* p_alias = reinterpret_cast<half2*>(&r);
+ p_alias[0] = __halves2half2(a, __hadd(a, __float2half(1.0f)));
+ p_alias[1] = __halves2half2(__hadd(a, __float2half(2.0f)),
+ __hadd(a, __float2half(3.0f)));
+ p_alias[2] = __halves2half2(__hadd(a, __float2half(4.0f)),
+ __hadd(a, __float2half(5.0f)));
+ p_alias[3] = __halves2half2(__hadd(a, __float2half(6.0f)),
+ __hadd(a, __float2half(7.0f)));
+ return r;
+#else // EIGEN_CUDA_ARCH
+
+#if EIGEN_CUDA_ARCH >= 530
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+
+ half2 b = pset1<half2>(a);
+ half2 c;
+ half2 half_offset0 = __halves2half2(__float2half(0.0f),__float2half(2.0f));
+ half2 half_offset1 = __halves2half2(__float2half(4.0f),__float2half(6.0f));
+
+ c = __hadd2(b, half_offset0);
+ r_alias[0] = plset(__low2half(c));
+ r_alias[1] = plset(__high2half(c));
+
+ c = __hadd2(b, half_offset1);
+ r_alias[2] = plset(__low2half(c));
+ r_alias[3] = plset(__high2half(c));
+
+ return r;
+
+#else
+ float f = __half2float(a);
+ Packet4h2 r;
+ half2* p_alias = reinterpret_cast<half2*>(&r);
+ p_alias[0] = __halves2half2(a, __float2half(f + 1.0f));
+ p_alias[1] = __halves2half2(__float2half(f + 2.0f), __float2half(f + 3.0f));
+ p_alias[2] = __halves2half2(__float2half(f + 4.0f), __float2half(f + 5.0f));
+ p_alias[3] = __halves2half2(__float2half(f + 6.0f), __float2half(f + 7.0f));
+ return r;
+#endif
#endif
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2
+pselect<Packet4h2>(const Packet4h2& mask, const Packet4h2& a,
+ const Packet4h2& b) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* mask_alias = reinterpret_cast<const half2*>(&mask);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ const half2* b_alias = reinterpret_cast<const half2*>(&b);
+ r_alias[0] = pselect(mask_alias[0], a_alias[0], b_alias[0]);
+ r_alias[1] = pselect(mask_alias[1], a_alias[1], b_alias[1]);
+ r_alias[2] = pselect(mask_alias[2], a_alias[2], b_alias[2]);
+ r_alias[3] = pselect(mask_alias[3], a_alias[3], b_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2
+pcmp_eq<Packet4h2>(const Packet4h2& a, const Packet4h2& b) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ const half2* b_alias = reinterpret_cast<const half2*>(&b);
+ r_alias[0] = pcmp_eq(a_alias[0], b_alias[0]);
+ r_alias[1] = pcmp_eq(a_alias[1], b_alias[1]);
+ r_alias[2] = pcmp_eq(a_alias[2], b_alias[2]);
+ r_alias[3] = pcmp_eq(a_alias[3], b_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pand<Packet4h2>(
+ const Packet4h2& a, const Packet4h2& b) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ const half2* b_alias = reinterpret_cast<const half2*>(&b);
+ r_alias[0] = pand(a_alias[0], b_alias[0]);
+ r_alias[1] = pand(a_alias[1], b_alias[1]);
+ r_alias[2] = pand(a_alias[2], b_alias[2]);
+ r_alias[3] = pand(a_alias[3], b_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 por<Packet4h2>(
+ const Packet4h2& a, const Packet4h2& b) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ const half2* b_alias = reinterpret_cast<const half2*>(&b);
+ r_alias[0] = por(a_alias[0], b_alias[0]);
+ r_alias[1] = por(a_alias[1], b_alias[1]);
+ r_alias[2] = por(a_alias[2], b_alias[2]);
+ r_alias[3] = por(a_alias[3], b_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pxor<Packet4h2>(
+ const Packet4h2& a, const Packet4h2& b) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ const half2* b_alias = reinterpret_cast<const half2*>(&b);
+ r_alias[0] = pxor(a_alias[0], b_alias[0]);
+ r_alias[1] = pxor(a_alias[1], b_alias[1]);
+ r_alias[2] = pxor(a_alias[2], b_alias[2]);
+ r_alias[3] = pxor(a_alias[3], b_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2
+pandnot<Packet4h2>(const Packet4h2& a, const Packet4h2& b) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ const half2* b_alias = reinterpret_cast<const half2*>(&b);
+ r_alias[0] = pandnot(a_alias[0], b_alias[0]);
+ r_alias[1] = pandnot(a_alias[1], b_alias[1]);
+ r_alias[2] = pandnot(a_alias[2], b_alias[2]);
+ r_alias[3] = pandnot(a_alias[3], b_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 padd<Packet4h2>(
+ const Packet4h2& a, const Packet4h2& b) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ const half2* b_alias = reinterpret_cast<const half2*>(&b);
+ r_alias[0] = padd(a_alias[0], b_alias[0]);
+ r_alias[1] = padd(a_alias[1], b_alias[1]);
+ r_alias[2] = padd(a_alias[2], b_alias[2]);
+ r_alias[3] = padd(a_alias[3], b_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 psub<Packet4h2>(
+ const Packet4h2& a, const Packet4h2& b) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ const half2* b_alias = reinterpret_cast<const half2*>(&b);
+ r_alias[0] = psub(a_alias[0], b_alias[0]);
+ r_alias[1] = psub(a_alias[1], b_alias[1]);
+ r_alias[2] = psub(a_alias[2], b_alias[2]);
+ r_alias[3] = psub(a_alias[3], b_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pnegate(const Packet4h2& a) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ r_alias[0] = pnegate(a_alias[0]);
+ r_alias[1] = pnegate(a_alias[1]);
+ r_alias[2] = pnegate(a_alias[2]);
+ r_alias[3] = pnegate(a_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pconj(const Packet4h2& a) {
+ return a;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pmul<Packet4h2>(
+ const Packet4h2& a, const Packet4h2& b) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ const half2* b_alias = reinterpret_cast<const half2*>(&b);
+ r_alias[0] = pmul(a_alias[0], b_alias[0]);
+ r_alias[1] = pmul(a_alias[1], b_alias[1]);
+ r_alias[2] = pmul(a_alias[2], b_alias[2]);
+ r_alias[3] = pmul(a_alias[3], b_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pmadd<Packet4h2>(
+ const Packet4h2& a, const Packet4h2& b, const Packet4h2& c) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ const half2* b_alias = reinterpret_cast<const half2*>(&b);
+ const half2* c_alias = reinterpret_cast<const half2*>(&c);
+ r_alias[0] = pmadd(a_alias[0], b_alias[0], c_alias[0]);
+ r_alias[1] = pmadd(a_alias[1], b_alias[1], c_alias[1]);
+ r_alias[2] = pmadd(a_alias[2], b_alias[2], c_alias[2]);
+ r_alias[3] = pmadd(a_alias[3], b_alias[3], c_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pdiv<Packet4h2>(
+ const Packet4h2& a, const Packet4h2& b) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ const half2* b_alias = reinterpret_cast<const half2*>(&b);
+ r_alias[0] = pdiv(a_alias[0], b_alias[0]);
+ r_alias[1] = pdiv(a_alias[1], b_alias[1]);
+ r_alias[2] = pdiv(a_alias[2], b_alias[2]);
+ r_alias[3] = pdiv(a_alias[3], b_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pmin<Packet4h2>(
+ const Packet4h2& a, const Packet4h2& b) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ const half2* b_alias = reinterpret_cast<const half2*>(&b);
+ r_alias[0] = pmin(a_alias[0], b_alias[0]);
+ r_alias[1] = pmin(a_alias[1], b_alias[1]);
+ r_alias[2] = pmin(a_alias[2], b_alias[2]);
+ r_alias[3] = pmin(a_alias[3], b_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pmax<Packet4h2>(
+ const Packet4h2& a, const Packet4h2& b) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ const half2* b_alias = reinterpret_cast<const half2*>(&b);
+ r_alias[0] = pmax(a_alias[0], b_alias[0]);
+ r_alias[1] = pmax(a_alias[1], b_alias[1]);
+ r_alias[2] = pmax(a_alias[2], b_alias[2]);
+ r_alias[3] = pmax(a_alias[3], b_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux<Packet4h2>(
+ const Packet4h2& a) {
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+
+ return predux(a_alias[0]) + predux(a_alias[1]) +
+ predux(a_alias[2]) + predux(a_alias[3]);
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_max<Packet4h2>(
+ const Packet4h2& a) {
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ half2 m0 = __halves2half2(predux_max(a_alias[0]),
+ predux_max(a_alias[1]));
+ half2 m1 = __halves2half2(predux_max(a_alias[2]),
+ predux_max(a_alias[3]));
+ __half first = predux_max(m0);
+ __half second = predux_max(m1);
+#if EIGEN_CUDA_ARCH >= 530
+ return (__hgt(first, second) ? first : second);
+#else
+ float ffirst = __half2float(first);
+ float fsecond = __half2float(second);
+ return (ffirst > fsecond)? first: second;
+#endif
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_min<Packet4h2>(
+ const Packet4h2& a) {
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ half2 m0 = __halves2half2(predux_min(a_alias[0]),
+ predux_min(a_alias[1]));
+ half2 m1 = __halves2half2(predux_min(a_alias[2]),
+ predux_min(a_alias[3]));
+ __half first = predux_min(m0);
+ __half second = predux_min(m1);
+#if EIGEN_CUDA_ARCH >= 530
+ return (__hlt(first, second) ? first : second);
+#else
+ float ffirst = __half2float(first);
+ float fsecond = __half2float(second);
+ return (ffirst < fsecond)? first: second;
+#endif
+}
+
+// likely overflow/underflow
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Eigen::half predux_mul<Packet4h2>(
+ const Packet4h2& a) {
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ return predux_mul(pmul(pmul(a_alias[0], a_alias[1]),
+ pmul(a_alias[2], a_alias[3])));
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2
+plog1p<Packet4h2>(const Packet4h2& a) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ r_alias[0] = plog1p(a_alias[0]);
+ r_alias[1] = plog1p(a_alias[1]);
+ r_alias[2] = plog1p(a_alias[2]);
+ r_alias[3] = plog1p(a_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2
+pexpm1<Packet4h2>(const Packet4h2& a) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ r_alias[0] = pexpm1(a_alias[0]);
+ r_alias[1] = pexpm1(a_alias[1]);
+ r_alias[2] = pexpm1(a_alias[2]);
+ r_alias[3] = pexpm1(a_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 plog<Packet4h2>(const Packet4h2& a) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ r_alias[0] = plog(a_alias[0]);
+ r_alias[1] = plog(a_alias[1]);
+ r_alias[2] = plog(a_alias[2]);
+ r_alias[3] = plog(a_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pexp<Packet4h2>(const Packet4h2& a) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ r_alias[0] = pexp(a_alias[0]);
+ r_alias[1] = pexp(a_alias[1]);
+ r_alias[2] = pexp(a_alias[2]);
+ r_alias[3] = pexp(a_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 psqrt<Packet4h2>(const Packet4h2& a) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ r_alias[0] = psqrt(a_alias[0]);
+ r_alias[1] = psqrt(a_alias[1]);
+ r_alias[2] = psqrt(a_alias[2]);
+ r_alias[3] = psqrt(a_alias[3]);
+ return r;
+}
+
+template <>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2
+prsqrt<Packet4h2>(const Packet4h2& a) {
+ Packet4h2 r;
+ half2* r_alias = reinterpret_cast<half2*>(&r);
+ const half2* a_alias = reinterpret_cast<const half2*>(&a);
+ r_alias[0] = prsqrt(a_alias[0]);
+ r_alias[1] = prsqrt(a_alias[1]);
+ r_alias[2] = prsqrt(a_alias[2]);
+ r_alias[3] = prsqrt(a_alias[3]);
+ return r;
+}
+
+// The following specialized padd, pmul, pdiv, pmin, pmax, pset1 are needed for
+// the implementation of GPU half reduction.
+template<>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 padd<half2>(const half2& a,
+ const half2& b) {
+#if defined(EIGEN_HIP_DEVICE_COMPILE)
+
+ return __hadd2(a, b);
+
+#else // EIGEN_CUDA_ARCH
+
+#if EIGEN_CUDA_ARCH >= 530
+ return __hadd2(a, b);
+#else
+ float a1 = __low2float(a);
+ float a2 = __high2float(a);
+ float b1 = __low2float(b);
+ float b2 = __high2float(b);
+ float r1 = a1 + b1;
+ float r2 = a2 + b2;
+ return __floats2half2_rn(r1, r2);
+#endif
+
+#endif
+}
+
+template<>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmul<half2>(const half2& a,
+ const half2& b) {
+#if defined(EIGEN_HIP_DEVICE_COMPILE)
+
+ return __hmul2(a, b);
+
+#else // EIGEN_CUDA_ARCH
+
+#if EIGEN_CUDA_ARCH >= 530
+ return __hmul2(a, b);
+#else
+ float a1 = __low2float(a);
+ float a2 = __high2float(a);
+ float b1 = __low2float(b);
+ float b2 = __high2float(b);
+ float r1 = a1 * b1;
+ float r2 = a2 * b2;
+ return __floats2half2_rn(r1, r2);
+#endif
+
+#endif
+}
+
+template<>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pdiv<half2>(const half2& a,
+ const half2& b) {
+#if defined(EIGEN_HIP_DEVICE_COMPILE)
+
+ return __h2div(a, b);
+
+#else // EIGEN_CUDA_ARCH
+
+ float a1 = __low2float(a);
+ float a2 = __high2float(a);
+ float b1 = __low2float(b);
+ float b2 = __high2float(b);
+ float r1 = a1 / b1;
+ float r2 = a2 / b2;
+ return __floats2half2_rn(r1, r2);
+
+#endif
+}
+
+template<>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmin<half2>(const half2& a,
+ const half2& b) {
+ float a1 = __low2float(a);
+ float a2 = __high2float(a);
+ float b1 = __low2float(b);
+ float b2 = __high2float(b);
+ __half r1 = a1 < b1 ? __low2half(a) : __low2half(b);
+ __half r2 = a2 < b2 ? __high2half(a) : __high2half(b);
+ return __halves2half2(r1, r2);
+}
+
+template<>
+EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pmax<half2>(const half2& a,
+ const half2& b) {
+ float a1 = __low2float(a);
+ float a2 = __high2float(a);
+ float b1 = __low2float(b);
+ float b2 = __high2float(b);
+ __half r1 = a1 > b1 ? __low2half(a) : __low2half(b);
+ __half r2 = a2 > b2 ? __high2half(a) : __high2half(b);
+ return __halves2half2(r1, r2);
+}
+
+#endif // defined(EIGEN_CUDA_ARCH)
+
+#endif // defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDACC)
} // end namespace internal
diff --git a/Eigen/src/Core/arch/GPU/TypeCasting.h b/Eigen/src/Core/arch/GPU/TypeCasting.h
index c278f3fe8..754546225 100644
--- a/Eigen/src/Core/arch/GPU/TypeCasting.h
+++ b/Eigen/src/Core/arch/GPU/TypeCasting.h
@@ -17,12 +17,13 @@ namespace internal {
#if (defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 300) || \
(defined(EIGEN_HAS_HIP_FP16) && defined(EIGEN_HIP_DEVICE_COMPILE))
+
template <>
struct type_casting_traits<Eigen::half, float> {
enum {
VectorizedCast = 1,
- SrcCoeffRatio = 2,
- TgtCoeffRatio = 1
+ SrcCoeffRatio = 1,
+ TgtCoeffRatio = 2
};
};
@@ -32,15 +33,39 @@ template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pcast<half2, float4>(con
return make_float4(r1.x, r1.y, r2.x, r2.y);
}
+
+template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Packet4h2 pcast<float4, Packet4h2>(const float4& a, const float4& b) {
+ Packet4h2 r;
+ half2* r_alias=reinterpret_cast<half2*>(&r);
+ r_alias[0]=__floats2half2_rn(a.x,a.y);
+ r_alias[1]=__floats2half2_rn(a.z,a.w);
+ r_alias[2]=__floats2half2_rn(b.x,b.y);
+ r_alias[3]=__floats2half2_rn(b.z,b.w);
+ return r;
+}
+
template <>
struct type_casting_traits<float, Eigen::half> {
enum {
VectorizedCast = 1,
- SrcCoeffRatio = 1,
- TgtCoeffRatio = 2
+ SrcCoeffRatio = 2,
+ TgtCoeffRatio = 1
};
};
+template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float4 pcast<Packet4h2, float4>(const Packet4h2& a) {
+ // Simply discard the second half of the input
+ float4 r;
+ const half2* a_alias=reinterpret_cast<const half2*>(&a);
+ float2 r1 = __half22float2(a_alias[0]);
+ float2 r2 = __half22float2(a_alias[1]);
+ r.x=static_cast<float>(r1.x);
+ r.y=static_cast<float>(r1.y);
+ r.z=static_cast<float>(r2.x);
+ r.w=static_cast<float>(r2.y);
+ return r;
+}
+
template<> EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE half2 pcast<float4, half2>(const float4& a) {
// Simply discard the second half of the input
return __floats2half2_rn(a.x, a.y);
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorMeta.h b/unsupported/Eigen/CXX11/src/Tensor/TensorMeta.h
index 6afc98877..a3a750f21 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorMeta.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorMeta.h
@@ -53,10 +53,12 @@ struct PacketType : internal::packet_traits<Scalar> {
// For CUDA packet types when using a GpuDevice
#if defined(EIGEN_USE_GPU) && defined(EIGEN_HAS_GPU_FP16)
-template <>
+
+typedef ulonglong2 Packet4h2;
+template<>
struct PacketType<half, GpuDevice> {
- typedef half2 type;
- static const int size = 2;
+ typedef Packet4h2 type;
+ static const int size = 8;
enum {
HasAdd = 1,
HasSub = 1,
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorReduction.h b/unsupported/Eigen/CXX11/src/Tensor/TensorReduction.h
index 5ca694062..8332a9ae0 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorReduction.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorReduction.h
@@ -420,9 +420,9 @@ __global__ void FullReductionKernel(R, const S, I_, typename S::CoeffReturnType*
#if defined(EIGEN_HAS_GPU_FP16)
template <typename S, typename R, typename I_>
-__global__ void ReductionInitFullReduxKernelHalfFloat(R, const S, I_, half2*);
+__global__ void ReductionInitFullReduxKernelHalfFloat(R, const S, I_, internal::packet_traits<half>::type*);
template <int B, int N, typename S, typename R, typename I_>
-__global__ void FullReductionKernelHalfFloat(R, const S, I_, half*, half2*);
+__global__ void FullReductionKernelHalfFloat(R, const S, I_, half*, internal::packet_traits<half>::type*);
template <int NPT, typename S, typename R, typename I_>
__global__ void InnerReductionKernelHalfFloat(R, const S, I_, I_, half*);
@@ -863,8 +863,8 @@ struct TensorReductionEvaluatorBase<const TensorReductionOp<Op, Dims, ArgType, M
#if defined(EIGEN_USE_GPU) && (defined(EIGEN_GPUCC))
template <int B, int N, typename S, typename R, typename I_> KERNEL_FRIEND void internal::FullReductionKernel(R, const S, I_, typename S::CoeffReturnType*, unsigned int*);
#if defined(EIGEN_HAS_GPU_FP16)
- template <typename S, typename R, typename I_> KERNEL_FRIEND void internal::ReductionInitFullReduxKernelHalfFloat(R, const S, I_, half2*);
- template <int B, int N, typename S, typename R, typename I_> KERNEL_FRIEND void internal::FullReductionKernelHalfFloat(R, const S, I_, half*, half2*);
+ template <typename S, typename R, typename I_> KERNEL_FRIEND void internal::ReductionInitFullReduxKernelHalfFloat(R, const S, I_, internal::packet_traits<Eigen::half>::type*);
+ template <int B, int N, typename S, typename R, typename I_> KERNEL_FRIEND void internal::FullReductionKernelHalfFloat(R, const S, I_, half*, internal::packet_traits<Eigen::half>::type*);
template <int NPT, typename S, typename R, typename I_> KERNEL_FRIEND void internal::InnerReductionKernelHalfFloat(R, const S, I_, I_, half*);
#endif
template <int NPT, typename S, typename R, typename I_> KERNEL_FRIEND void internal::InnerReductionKernel(R, const S, I_, I_, typename S::CoeffReturnType*);
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorReductionGpu.h b/unsupported/Eigen/CXX11/src/Tensor/TensorReductionGpu.h
index 095bb54cc..9d3305cfd 100644
--- a/unsupported/Eigen/CXX11/src/Tensor/TensorReductionGpu.h
+++ b/unsupported/Eigen/CXX11/src/Tensor/TensorReductionGpu.h
@@ -98,7 +98,17 @@ __device__ inline void atomicReduce(half2* output, half2 accum, R<half>& reducer
}
}
}
-#endif // EIGEN_HAS_GPU_FP16
+// reduction should be associative since reduction is not atomic in wide vector but atomic in half2 operations
+template <template <typename T> class R>
+__device__ inline void atomicReduce(Packet4h2* output, Packet4h2 accum,
+ R<half>& reducer) {
+ half2* houtput=reinterpret_cast<half2*>(output);
+ half2* haccum=reinterpret_cast<half2*>(&accum);
+ for(int i=0;i<4;++i){
+ atomicReduce(houtput+i,*(haccum+i),reducer);
+ }
+}
+#endif // EIGEN_HAS_GPU_FP16
template <>
__device__ inline void atomicReduce(float* output, float accum, SumReducer<float>&) {
@@ -204,14 +214,26 @@ __global__ void FullReductionKernel(Reducer reducer, const Self input, Index num
#ifdef EIGEN_HAS_GPU_FP16
template <typename Self,
typename Reducer, typename Index>
-__global__ void ReductionInitFullReduxKernelHalfFloat(Reducer reducer, const Self input, Index num_coeffs, half2* scratch) {
+__global__ void ReductionInitFullReduxKernelHalfFloat(Reducer reducer, const Self input, Index num_coeffs,
+ packet_traits<Eigen::half>::type* scratch) {
eigen_assert(blockDim.x == 1);
eigen_assert(gridDim.x == 1);
- if (num_coeffs % 2 != 0) {
- half lastCoeff = input.m_impl.coeff(num_coeffs-1);
- *scratch = __halves2half2(lastCoeff, reducer.initialize());
+ typedef packet_traits<Eigen::half>::type packet_type;
+ Index packet_remainder =
+ num_coeffs % Index(unpacket_traits<packet_type>::size);
+ if (packet_remainder != 0) {
+ half2* h2scratch = reinterpret_cast<half2*>(scratch);
+ for (Index i = num_coeffs - packet_remainder; i + 2 <= num_coeffs; i += 2) {
+ *h2scratch =
+ __halves2half2(input.m_impl.coeff(i), input.m_impl.coeff(i + 1));
+ h2scratch++;
+ }
+ if ((num_coeffs & 1) != 0) {
+ half lastCoeff = input.m_impl.coeff(num_coeffs - 1);
+ *h2scratch = __halves2half2(lastCoeff, reducer.initialize());
+ }
} else {
- *scratch = reducer.template initializePacket<half2>();
+ *scratch = reducer.template initializePacket<packet_type>();
}
}
@@ -220,44 +242,64 @@ template <typename Self,
__global__ void ReductionInitKernelHalfFloat(Reducer reducer, const Self input, Index num_coeffs, half* output) {
const Index thread_id = blockIdx.x * blockDim.x + threadIdx.x;
const Index num_threads = blockDim.x * gridDim.x;
- const Index num_packets = num_coeffs / 2;
+ typedef typename packet_traits<Eigen::half>::type PacketType;
+
+ const Index num_packets =
+ num_coeffs / Index(unpacket_traits<PacketType>::size);
+ PacketType* p_output = reinterpret_cast<PacketType*>(output);
for (Index i = thread_id; i < num_packets; i += num_threads) {
- ((half2*)output)[i] = reducer.template initializePacket<half2>();
+ p_output[i] = reducer.template initializePacket<PacketType>();
}
-
- if (thread_id == 0 && num_coeffs % 2 != 0) {
- output[num_coeffs-1] = reducer.initialize();
+ Index packet_remainder =
+ num_coeffs % Index(unpacket_traits<PacketType>::size);
+ if (thread_id < packet_remainder) {
+ output[num_coeffs - packet_remainder + thread_id] = reducer.initialize();
}
}
template <int BlockSize, int NumPerThread, typename Self,
typename Reducer, typename Index>
__global__ void FullReductionKernelHalfFloat(Reducer reducer, const Self input, Index num_coeffs,
- half* output, half2* scratch) {
- eigen_assert(NumPerThread % 2 == 0);
-
- const Index first_index = blockIdx.x * BlockSize * NumPerThread + 2*threadIdx.x;
+ half* output, packet_traits<Eigen::half>::type* scratch) {
+ typedef typename packet_traits<Eigen::half>::type PacketType;
+ const int packet_width = unpacket_traits<PacketType>::size;
+ eigen_assert(NumPerThread % packet_width == 0);
+ const Index first_index =
+ blockIdx.x * BlockSize * NumPerThread + packet_width * threadIdx.x;
// Initialize the output value if it wasn't initialized by the ReductionInitKernel
if (gridDim.x == 1) {
if (first_index == 0) {
- if (num_coeffs % 2 != 0) {
- half last = input.m_impl.coeff(num_coeffs-1);
- *scratch = __halves2half2(last, reducer.initialize());
+ int rem = num_coeffs % packet_width;
+ if (rem != 0) {
+ half2* p_scratch = reinterpret_cast<half2*>(scratch);
+ *scratch = reducer.template initializePacket<PacketType>();
+ for (int i = 0; i < rem / 2; i++) {
+ *p_scratch = __halves2half2(
+ input.m_impl.coeff(num_coeffs - packet_width + 2 * i),
+ input.m_impl.coeff(num_coeffs - packet_width + 2 * i + 1));
+ p_scratch++;
+ }
+ if ((num_coeffs & 1) != 0) {
+ half last = input.m_impl.coeff(num_coeffs - 1);
+ *p_scratch = __halves2half2(last, reducer.initialize());
+ }
} else {
- *scratch = reducer.template initializePacket<half2>();
+ *scratch = reducer.template initializePacket<PacketType>();
}
}
__syncthreads();
}
-
- half2 accum = reducer.template initializePacket<half2>();
- const Index max_iter = numext::mini<Index>((num_coeffs - first_index) / 2, NumPerThread*BlockSize / 2);
+
+ PacketType accum = reducer.template initializePacket<PacketType>();
+ const Index max_iter =
+ numext::mini<Index>((num_coeffs - first_index) / packet_width,
+ NumPerThread * BlockSize / packet_width);
for (Index i = 0; i < max_iter; i += BlockSize) {
- const Index index = first_index + 2*i;
- eigen_assert(index + 1 < num_coeffs);
- half2 val = input.m_impl.template packet<Unaligned>(index);
+ const Index index = first_index + packet_width * i;
+ eigen_assert(index + packet_width < num_coeffs);
+ PacketType val = input.m_impl.template packet<Unaligned>(index);
reducer.reducePacket(val, &accum);
}
@@ -270,10 +312,22 @@ __global__ void FullReductionKernelHalfFloat(Reducer reducer, const Self input,
wka_out.i = __shfl_down(wka_in.i, offset, warpSize);
reducer.reducePacket(wka_out.h, &accum);
#elif defined(EIGEN_CUDA_SDK_VER) && EIGEN_CUDA_SDK_VER < 90000
- reducer.reducePacket(__shfl_down(accum, offset, warpSize), &accum);
+ PacketType r1;
+ half2* hr = reinterpret_cast<half2*>(&r1);
+ half2* hacc = reinterpret_cast<half2*>(&accum);
+ for (int i = 0; i < packet_width / 2; i++) {
+ hr[i] = __shfl_down(hacc[i], offset, warpSize);
+ }
+ reducer.reducePacket(r1, &accum);
#else
- int temp = __shfl_down_sync(0xFFFFFFFF, *(int*)(&accum), (unsigned)offset, warpSize);
- reducer.reducePacket(*(half2*)(&temp), &accum);
+ PacketType r1;
+ half2* hr = reinterpret_cast<half2*>(&r1);
+ half2* hacc = reinterpret_cast<half2*>(&accum);
+ for (int i = 0; i < packet_width / 2; i++) {
+ hr[i] = __shfl_down_sync(0xFFFFFFFF, hacc[i], (unsigned)offset, warpSize);
+ }
+ reducer.reducePacket(r1, &accum);
+
#endif
}
@@ -281,21 +335,33 @@ __global__ void FullReductionKernelHalfFloat(Reducer reducer, const Self input,
atomicReduce(scratch, accum, reducer);
}
+ __syncthreads();
+ half2* rv1 = reinterpret_cast<half2*>(scratch);
+ if (packet_width > 2) {
+ reducer.reducePacket(rv1[2], rv1);
+ reducer.reducePacket(rv1[3], rv1 + 1);
+ reducer.reducePacket(rv1[1], rv1);
+ }
if (gridDim.x == 1) {
- __syncthreads();
if (first_index == 0) {
- half tmp = __low2half(*scratch);
- reducer.reduce(__high2half(*scratch), &tmp);
+ half tmp = __low2half(*rv1);
+ reducer.reduce(__high2half(*rv1), &tmp);
*output = tmp;
}
}
}
template <typename Op>
-__global__ void ReductionCleanupKernelHalfFloat(Op reducer, half* output, half2* scratch) {
+__global__ void ReductionCleanupKernelHalfFloat(Op reducer, half* output, packet_traits<Eigen::half>::type* scratch) {
eigen_assert(threadIdx.x == 1);
- half tmp = __low2half(*scratch);
- reducer.reduce(__high2half(*scratch), &tmp);
+ half2* pscratch = reinterpret_cast<half2*>(scratch);
+ half tmp = __float2half(0.f);
+ typedef packet_traits<Eigen::half>::type packet_type;
+ for (int i = 0; i < unpacket_traits<packet_type>::size; i += 2) {
+ reducer.reduce(__low2half(*pscratch), &tmp);
+ reducer.reduce(__high2half(*pscratch), &tmp);
+ pscratch++;
+ }
*output = tmp;
}
@@ -345,11 +411,13 @@ template <typename Self, typename Op>
struct FullReductionLauncher<Self, Op, Eigen::half, true> {
static void run(const Self& self, Op& reducer, const GpuDevice& device, half* output, typename Self::Index num_coeffs) {
typedef typename Self::Index Index;
+ typedef typename packet_traits<Eigen::half>::type PacketType;
const int block_size = 256;
const int num_per_thread = 128;
const int num_blocks = divup<int>(num_coeffs, block_size * num_per_thread);
- half2* scratch = static_cast<half2*>(device.scratchpad());
+ PacketType* scratch = static_cast<PacketType*>(device.scratchpad());
+ // half2* scratch = static_cast<half2*>(device.scratchpad());
if (num_blocks > 1) {
// We initialize the output and the scrathpad outside the reduction kernel when we can't be sure that there
@@ -459,8 +527,8 @@ __global__ void InnerReductionKernel(Reducer reducer, const Self input, Index nu
for (int offset = warpSize/2; offset > 0; offset /= 2) {
#if defined(EIGEN_HIPCC)
// use std::is_floating_point to determine the type of reduced_val
- // This is needed because when Type == double, hipcc will give a "call to __shfl_down is ambguous" error
- // and list the float and int versions of __shfl_down as the candidate functions.
+ // This is needed because when Type == double, hipcc will give a "call to __shfl_down is ambguous" error
+ // and list the float and int versions of __shfl_down as the candidate functions.
if (std::is_floating_point<Type>::value) {
reducer.reduce(__shfl_down(static_cast<float>(reduced_val), offset), &reduced_val);
} else {
@@ -494,7 +562,9 @@ __global__ void InnerReductionKernelHalfFloat(Reducer reducer, const Self input,
eigen_assert(gridDim.y == 1);
eigen_assert(gridDim.z == 1);
- const int unroll_times = 16;
+ typedef typename packet_traits<Eigen::half>::type PacketType;
+ const int packet_width = unpacket_traits<PacketType>::size;
+ const int unroll_times = 16 / packet_width;
eigen_assert(NumPerThread % unroll_times == 0);
eigen_assert(unroll_times % 2 == 0);
@@ -506,10 +576,11 @@ __global__ void InnerReductionKernelHalfFloat(Reducer reducer, const Self input,
// Initialize the output values if they weren't initialized by the ReductionInitKernel
if (gridDim.x == 1) {
- Index i = 2*thread_id;
- for (; i + 1 < num_preserved_coeffs; i += 2*num_threads) {
- half* loc = output + i;
- *((half2*)loc) = reducer.template initializePacket<half2>();
+ Index i = packet_width * thread_id;
+ for (; i + packet_width <= num_preserved_coeffs;
+ i += packet_width * num_threads) {
+ PacketType* poutput = reinterpret_cast<PacketType*>(output + i);
+ *poutput = reducer.template initializePacket<PacketType>();
}
if (i < num_preserved_coeffs) {
output[i] = reducer.initialize();
@@ -518,42 +589,71 @@ __global__ void InnerReductionKernelHalfFloat(Reducer reducer, const Self input,
}
for (Index i = blockIdx.x; i < num_input_blocks; i += gridDim.x) {
- const Index row = 2 * (i / input_col_blocks);
+ const Index row = 2 * (i / input_col_blocks); // everybody takes 2 rows
if (row + 1 < num_preserved_coeffs) {
const Index col_block = i % input_col_blocks;
- const Index col_begin = 2 * (col_block * blockDim.x * NumPerThread + threadIdx.x);
+ const Index col_begin =
+ packet_width * (col_block * blockDim.x * NumPerThread + threadIdx.x);
- half2 reduced_val1 = reducer.template initializePacket<half2>();
- half2 reduced_val2 = reducer.template initializePacket<half2>();
+ PacketType reduced_val1 = reducer.template initializePacket<PacketType>();
+ PacketType reduced_val2 = reducer.template initializePacket<PacketType>();
for (Index j = 0; j < NumPerThread; j += unroll_times) {
- const Index last_col = col_begin + blockDim.x * (j + unroll_times - 1) * 2;
+ const Index last_col =
+ col_begin + blockDim.x * (j + unroll_times - 1) * packet_width;
if (last_col >= num_coeffs_to_reduce) {
Index col = col_begin + blockDim.x * j;
- for (; col + 1 < num_coeffs_to_reduce; col += blockDim.x) {
- const half2 val1 = input.m_impl.template packet<Unaligned>(row * num_coeffs_to_reduce + col);
+ for (; col + packet_width <= num_coeffs_to_reduce;
+ col += blockDim.x) {
+ const PacketType val1 = input.m_impl.template packet<Unaligned>(
+ row * num_coeffs_to_reduce + col);
reducer.reducePacket(val1, &reduced_val1);
- const half2 val2 = input.m_impl.template packet<Unaligned>((row+1) * num_coeffs_to_reduce + col);
+ const PacketType val2 = input.m_impl.template packet<Unaligned>(
+ (row + 1) * num_coeffs_to_reduce + col);
reducer.reducePacket(val2, &reduced_val2);
}
if (col < num_coeffs_to_reduce) {
- // Peel;
- const half last1 = input.m_impl.coeff(row * num_coeffs_to_reduce + col);
- const half2 val1 = __halves2half2(last1, reducer.initialize());
- reducer.reducePacket(val1, &reduced_val1);
- const half last2 = input.m_impl.coeff((row+1) * num_coeffs_to_reduce + col);
- const half2 val2 = __halves2half2(last2, reducer.initialize());
- reducer.reducePacket(val2, &reduced_val2);
+ PacketType r1 = reducer.template initializePacket<PacketType>();
+ PacketType r2 = reducer.template initializePacket<PacketType>();
+ half2* hr1 = reinterpret_cast<half2*>(&r1);
+ half2* hr2 = reinterpret_cast<half2*>(&r2);
+ while (col + 1 < num_coeffs_to_reduce) {
+ *hr1 = __halves2half2(
+ input.m_impl.coeff(row * num_coeffs_to_reduce + col),
+ input.m_impl.coeff(row * num_coeffs_to_reduce + col + 1));
+ *hr2 = __halves2half2(
+ input.m_impl.coeff((row + 1) * num_coeffs_to_reduce + col),
+ input.m_impl.coeff((row + 1) * num_coeffs_to_reduce + col +
+ 1));
+ hr1++;
+ hr2++;
+ col += 2;
+ }
+ if (col < num_coeffs_to_reduce) {
+ // Peel;
+ const half last1 =
+ input.m_impl.coeff(row * num_coeffs_to_reduce + col);
+ *hr1 = __halves2half2(last1, reducer.initialize());
+ const half last2 =
+ input.m_impl.coeff((row + 1) * num_coeffs_to_reduce + col);
+ *hr2 = __halves2half2(last2, reducer.initialize());
+ }
+ reducer.reducePacket(r1, &reduced_val1);
+ reducer.reducePacket(r2, &reduced_val2);
}
break;
} else {
// Faster version of the loop with no branches after unrolling.
#pragma unroll
for (int k = 0; k < unroll_times; ++k) {
- const Index col = col_begin + blockDim.x * (j + k) * 2;
- reducer.reducePacket(input.m_impl.template packet<Unaligned>(row * num_coeffs_to_reduce + col), &reduced_val1);
- reducer.reducePacket(input.m_impl.template packet<Unaligned>((row + 1)* num_coeffs_to_reduce + col), &reduced_val2);
+ const Index col = col_begin + blockDim.x * (j + k) * packet_width;
+ reducer.reducePacket(input.m_impl.template packet<Unaligned>(
+ row * num_coeffs_to_reduce + col),
+ &reduced_val1);
+ reducer.reducePacket(input.m_impl.template packet<Unaligned>(
+ (row + 1) * num_coeffs_to_reduce + col),
+ &reduced_val2);
}
}
}
@@ -561,33 +661,63 @@ __global__ void InnerReductionKernelHalfFloat(Reducer reducer, const Self input,
#pragma unroll
for (int offset = warpSize/2; offset > 0; offset /= 2) {
#if defined(EIGEN_HIPCC)
- // FIXME : remove this workaround once we have native half/half2 support for __shfl_down
- union { int i; half2 h; } wka_in, wka_out;
+ // FIXME : remove this workaround once we have native half/half2 support for __shfl_down
+ union { int i; half2 h; } wka_in, wka_out;
- wka_in.h = reduced_val1;
- wka_out.i = __shfl_down(wka_in.i, offset, warpSize);
+ wka_in.h = reduced_val1;
+ wka_out.i = __shfl_down(wka_in.i, offset, warpSize);
reducer.reducePacket(wka_out.h, &reduced_val1);
-
- wka_in.h = reduced_val2;
- wka_out.i = __shfl_down(wka_in.i, offset, warpSize);
+
+ wka_in.h = reduced_val2;
+ wka_out.i = __shfl_down(wka_in.i, offset, warpSize);
reducer.reducePacket(wka_out.h, &reduced_val2);
#elif defined(EIGEN_CUDA_SDK_VER) && EIGEN_CUDA_SDK_VER < 90000
- reducer.reducePacket(__shfl_down(reduced_val1, offset, warpSize), &reduced_val1);
- reducer.reducePacket(__shfl_down(reduced_val2, offset, warpSize), &reduced_val2);
+ PacketType r1;
+ PacketType r2;
+ half2* hr1 = reinterpret_cast<half2*>(&r1);
+ half2* hr2 = reinterpret_cast<half2*>(&r2);
+ half2* rv1 = reinterpret_cast<half2*>(&reduced_val1);
+ half2* rv2 = reinterpret_cast<half2*>(&reduced_val2);
+ for (int i = 0; i < packet_width / 2; i++) {
+ hr1[i] = __shfl_down(rv1[i], offset, warpSize);
+ hr2[i] = __shfl_down(rv2[i], offset, warpSize);
+ }
+ reducer.reducePacket(r1, &reduced_val1);
+ reducer.reducePacket(r2, &reduced_val2);
#else
- int temp1 = __shfl_down_sync(0xFFFFFFFF, *(int*)(&reduced_val1), (unsigned)offset, warpSize);
- int temp2 = __shfl_down_sync(0xFFFFFFFF, *(int*)(&reduced_val2), (unsigned)offset, warpSize);
- reducer.reducePacket(*(half2*)(&temp1), &reduced_val1);
- reducer.reducePacket(*(half2*)(&temp2), &reduced_val2);
+ PacketType r1;
+ PacketType r2;
+ half2* hr1 = reinterpret_cast<half2*>(&r1);
+ half2* hr2 = reinterpret_cast<half2*>(&r2);
+ half2* rr1 = reinterpret_cast<half2*>(&reduced_val1);
+ half2* rr2 = reinterpret_cast<half2*>(&reduced_val2);
+ for (int i = 0; i < packet_width / 2; i++) {
+ hr1[i] =
+ __shfl_down_sync(0xFFFFFFFF, rr1[i], (unsigned)offset, warpSize);
+ hr2[i] =
+ __shfl_down_sync(0xFFFFFFFF, rr2[i], (unsigned)offset, warpSize);
+ }
+ reducer.reducePacket(r1, &reduced_val1);
+ reducer.reducePacket(r2, &reduced_val2);
+
#endif
}
-
- half val1 = __low2half(reduced_val1);
- reducer.reduce(__high2half(reduced_val1), &val1);
- half val2 = __low2half(reduced_val2);
- reducer.reduce(__high2half(reduced_val2), &val2);
- half2 val = __halves2half2(val1, val2);
-
+ half2* rv1 = reinterpret_cast<half2*>(&reduced_val1);
+ half2* rv2 = reinterpret_cast<half2*>(&reduced_val2);
+ half2 val;
+ if (packet_width > 2) {
+ reducer.reducePacket(rv1[2], rv1);
+ reducer.reducePacket(rv1[3], rv1 + 1);
+ reducer.reducePacket(rv1[1], rv1);
+ reducer.reducePacket(rv2[2], rv2);
+ reducer.reducePacket(rv2[3], rv2 + 1);
+ reducer.reducePacket(rv2[1], rv2);
+ }
+ half val1 = __low2half(*rv1);
+ reducer.reduce(__high2half(*rv1), &val1);
+ half val2 = __low2half(*rv2);
+ reducer.reduce(__high2half(*rv2), &val2);
+ val = __halves2half2(val1, val2);
if ((threadIdx.x & (warpSize - 1)) == 0) {
half* loc = output + row;
atomicReduce((half2*)loc, val, reducer);
generated by cgit v1.2.3 (git 2.39.1) at 2024-05-29 21:09:44 +0000