Update Eigen to commit:e8515f78ac098329ab9f8cab21c87caede090a3f CHANGELOG ========= e8515f78a - Fix sparse triangular view iterator 6d829e766 - Fix extra semicolon in XprHelper ba47341a1 - [SYCL-2020] Enabling half precision support for SYCL. 92a77a596 - Fix call to static functions from device by adding EIGEN_DEVICE_FUNC attribute to run methods 8f858a4ea - Export ThreadPool symbols from legacy header. 4e598ad25 - New panel modes for GEMM MMA (real & complex). 2c64a655f - Stage will not be ok if pardiso returned error PiperOrigin-RevId: 571952024 Change-Id: Ic4ba7448c04e842e42dfb20a19f7cc146a00f66f
diff --git a/Eigen/src/Core/arch/AVX512/PacketMathFP16.h b/Eigen/src/Core/arch/AVX512/PacketMathFP16.h index fc11174..faa3853 100644 --- a/Eigen/src/Core/arch/AVX512/PacketMathFP16.h +++ b/Eigen/src/Core/arch/AVX512/PacketMathFP16.h
@@ -1,870 +1,870 @@ -// This file is part of Eigen, a lightweight C++ template library -// for linear algebra. -// -// -// -// This Source Code Form is subject to the terms of the Mozilla -// Public License v. 2.0. If a copy of the MPL was not distributed -// with this file, You can obtain one at http://mozilla.org/MPL/2.0/. - -#ifndef EIGEN_PACKET_MATH_FP16_AVX512_H -#define EIGEN_PACKET_MATH_FP16_AVX512_H - +// This file is part of Eigen, a lightweight C++ template library +// for linear algebra. +// +// +// +// This Source Code Form is subject to the terms of the Mozilla +// Public License v. 2.0. If a copy of the MPL was not distributed +// with this file, You can obtain one at http://mozilla.org/MPL/2.0/. + +#ifndef EIGEN_PACKET_MATH_FP16_AVX512_H +#define EIGEN_PACKET_MATH_FP16_AVX512_H + // IWYU pragma: private -#include "../../InternalHeaderCheck.h" - -namespace Eigen { - -namespace internal { - -typedef __m512h Packet32h; -typedef eigen_packet_wrapper<__m256i, 1> Packet16h; -typedef eigen_packet_wrapper<__m128i, 2> Packet8h; - -template <> -struct is_arithmetic<Packet8h> { - enum { value = true }; -}; - -template <> -struct packet_traits<half> : default_packet_traits { - typedef Packet32h type; - typedef Packet16h half; - enum { - Vectorizable = 1, - AlignedOnScalar = 1, - size = 32, - - HasCmp = 1, - HasAdd = 1, - HasSub = 1, - HasMul = 1, - HasDiv = 1, - HasNegate = 1, - HasAbs = 1, - HasAbs2 = 0, - HasMin = 1, - HasMax = 1, - HasConj = 1, - HasSetLinear = 0, - HasLog = 1, - HasLog1p = 1, - HasExp = 1, - HasExpm1 = 1, - HasSqrt = 1, - HasRsqrt = 1, - // These ones should be implemented in future - HasBessel = 0, - HasNdtri = 0, - HasSin = EIGEN_FAST_MATH, - HasCos = EIGEN_FAST_MATH, - HasTanh = EIGEN_FAST_MATH, - HasErf = 0, // EIGEN_FAST_MATH, - HasBlend = 0, - HasRound = 1, - HasFloor = 1, - HasCeil = 1, - HasRint = 1 - }; -}; - -template <> -struct unpacket_traits<Packet32h> { - typedef Eigen::half type; - typedef Packet16h half; - enum { - size = 32, - alignment = Aligned64, - vectorizable = true, - masked_load_available = false, - masked_store_available = false - }; -}; - -template <> -struct unpacket_traits<Packet16h> { - typedef Eigen::half type; - typedef Packet8h half; - enum { - size = 16, - alignment = Aligned32, - vectorizable = true, - masked_load_available = false, - masked_store_available = false - }; -}; - -template <> -struct unpacket_traits<Packet8h> { - typedef Eigen::half type; - typedef Packet8h half; - enum { - size = 8, - alignment = Aligned16, - vectorizable = true, - masked_load_available = false, - masked_store_available = false - }; -}; - -// Memory functions - -// pset1 - -template <> -EIGEN_STRONG_INLINE Packet32h pset1<Packet32h>(const Eigen::half& from) { - return _mm512_set1_ph(static_cast<_Float16>(from)); -} - -// pset1frombits -template <> -EIGEN_STRONG_INLINE Packet32h pset1frombits<Packet32h>(unsigned short from) { - return _mm512_castsi512_ph(_mm512_set1_epi16(from)); -} - -// pfirst - -template <> -EIGEN_STRONG_INLINE Eigen::half pfirst<Packet32h>(const Packet32h& from) { -#ifdef EIGEN_VECTORIZE_AVX512DQ - return half_impl::raw_uint16_to_half( - static_cast<unsigned short>(_mm256_extract_epi16(_mm512_extracti32x8_epi32(_mm512_castph_si512(from), 0), 0))); -#else - Eigen::half dest[32]; - _mm512_storeu_ph(dest, from); - return dest[0]; -#endif -} - -// pload - -template <> -EIGEN_STRONG_INLINE Packet32h pload<Packet32h>(const Eigen::half* from) { - EIGEN_DEBUG_ALIGNED_LOAD return _mm512_load_ph(from); -} - -// ploadu - -template <> -EIGEN_STRONG_INLINE Packet32h ploadu<Packet32h>(const Eigen::half* from) { - EIGEN_DEBUG_UNALIGNED_LOAD return _mm512_loadu_ph(from); -} - -// pstore - -template <> -EIGEN_STRONG_INLINE void pstore<half>(Eigen::half* to, const Packet32h& from) { - EIGEN_DEBUG_ALIGNED_STORE _mm512_store_ph(to, from); -} - -// pstoreu - -template <> -EIGEN_STRONG_INLINE void pstoreu<half>(Eigen::half* to, const Packet32h& from) { - EIGEN_DEBUG_UNALIGNED_STORE _mm512_storeu_ph(to, from); -} - -// ploaddup -template <> -EIGEN_STRONG_INLINE Packet32h ploaddup<Packet32h>(const Eigen::half* from) { - __m512h a = _mm512_castph256_ph512(_mm256_loadu_ph(from)); - return _mm512_permutexvar_ph(_mm512_set_epi16(15, 15, 14, 14, 13, 13, 12, 12, 11, 11, 10, 10, 9, 9, 8, 8, 7, 7, 6, 6, - 5, 5, 4, 4, 3, 3, 2, 2, 1, 1, 0, 0), - a); -} - -// ploadquad -template <> -EIGEN_STRONG_INLINE Packet32h ploadquad<Packet32h>(const Eigen::half* from) { - __m512h a = _mm512_castph128_ph512(_mm_loadu_ph(from)); - return _mm512_permutexvar_ph( - _mm512_set_epi16(7, 7, 7, 7, 6, 6, 6, 6, 5, 5, 5, 5, 4, 4, 4, 4, 3, 3, 3, 3, 2, 2, 2, 2, 1, 1, 1, 1, 0, 0, 0, 0), - a); -} - -// pabs - -template <> -EIGEN_STRONG_INLINE Packet32h pabs<Packet32h>(const Packet32h& a) { - return _mm512_abs_ph(a); -} - -// psignbit - -template <> -EIGEN_STRONG_INLINE Packet32h psignbit<Packet32h>(const Packet32h& a) { - return _mm512_castsi512_ph(_mm512_srai_epi16(_mm512_castph_si512(a), 15)); -} - -// pmin - -template <> -EIGEN_STRONG_INLINE Packet32h pmin<Packet32h>(const Packet32h& a, const Packet32h& b) { - return _mm512_min_ph(a, b); -} - -// pmax - -template <> -EIGEN_STRONG_INLINE Packet32h pmax<Packet32h>(const Packet32h& a, const Packet32h& b) { - return _mm512_max_ph(a, b); -} - -// plset -template <> -EIGEN_STRONG_INLINE Packet32h plset<Packet32h>(const half& a) { - return _mm512_add_ph(_mm512_set1_ph(a), - _mm512_set_ph(31.0f, 30.0f, 29.0f, 28.0f, 27.0f, 26.0f, 25.0f, 24.0f, 23.0f, 22.0f, 21.0f, 20.0f, - 19.0f, 18.0f, 17.0f, 16.0f, 15.0f, 14.0f, 13.0f, 12.0f, 11.0f, 10.0f, 9.0f, 8.0f, - 7.0f, 6.0f, 5.0f, 4.0f, 3.0f, 2.0f, 1.0f, 0.0f)); -} - -// por - -template <> -EIGEN_STRONG_INLINE Packet32h por(const Packet32h& a, const Packet32h& b) { - return _mm512_castsi512_ph(_mm512_or_si512(_mm512_castph_si512(a), _mm512_castph_si512(b))); -} - -// pxor - -template <> -EIGEN_STRONG_INLINE Packet32h pxor(const Packet32h& a, const Packet32h& b) { - return _mm512_castsi512_ph(_mm512_xor_si512(_mm512_castph_si512(a), _mm512_castph_si512(b))); -} - -// pand - -template <> -EIGEN_STRONG_INLINE Packet32h pand(const Packet32h& a, const Packet32h& b) { - return _mm512_castsi512_ph(_mm512_and_si512(_mm512_castph_si512(a), _mm512_castph_si512(b))); -} - -// pandnot - -template <> -EIGEN_STRONG_INLINE Packet32h pandnot(const Packet32h& a, const Packet32h& b) { - return _mm512_castsi512_ph(_mm512_andnot_si512(_mm512_castph_si512(b), _mm512_castph_si512(a))); -} - -// pselect - -template <> -EIGEN_DEVICE_FUNC inline Packet32h pselect(const Packet32h& mask, const Packet32h& a, const Packet32h& b) { - __mmask32 mask32 = _mm512_cmp_epi16_mask(_mm512_castph_si512(mask), _mm512_setzero_epi32(), _MM_CMPINT_EQ); - return _mm512_mask_blend_ph(mask32, a, b); -} - -// pcmp_eq - -template <> -EIGEN_STRONG_INLINE Packet32h pcmp_eq(const Packet32h& a, const Packet32h& b) { - __mmask32 mask = _mm512_cmp_ph_mask(a, b, _CMP_EQ_OQ); - return _mm512_castsi512_ph(_mm512_mask_set1_epi16(_mm512_set1_epi32(0), mask, 0xffffu)); -} - -// pcmp_le - -template <> -EIGEN_STRONG_INLINE Packet32h pcmp_le(const Packet32h& a, const Packet32h& b) { - __mmask32 mask = _mm512_cmp_ph_mask(a, b, _CMP_LE_OQ); - return _mm512_castsi512_ph(_mm512_mask_set1_epi16(_mm512_set1_epi32(0), mask, 0xffffu)); -} - -// pcmp_lt - -template <> -EIGEN_STRONG_INLINE Packet32h pcmp_lt(const Packet32h& a, const Packet32h& b) { - __mmask32 mask = _mm512_cmp_ph_mask(a, b, _CMP_LT_OQ); - return _mm512_castsi512_ph(_mm512_mask_set1_epi16(_mm512_set1_epi32(0), mask, 0xffffu)); -} - -// pcmp_lt_or_nan - -template <> -EIGEN_STRONG_INLINE Packet32h pcmp_lt_or_nan(const Packet32h& a, const Packet32h& b) { - __mmask32 mask = _mm512_cmp_ph_mask(a, b, _CMP_NGE_UQ); - return _mm512_castsi512_ph(_mm512_mask_set1_epi16(_mm512_set1_epi16(0), mask, 0xffffu)); -} - -// padd - -template <> -EIGEN_STRONG_INLINE Packet32h padd<Packet32h>(const Packet32h& a, const Packet32h& b) { - return _mm512_add_ph(a, b); -} - -template <> -EIGEN_STRONG_INLINE Packet16h padd<Packet16h>(const Packet16h& a, const Packet16h& b) { - return _mm256_castph_si256(_mm256_add_ph(_mm256_castsi256_ph(a), _mm256_castsi256_ph(b))); -} - -template <> -EIGEN_STRONG_INLINE Packet8h padd<Packet8h>(const Packet8h& a, const Packet8h& b) { - return _mm_castph_si128(_mm_add_ph(_mm_castsi128_ph(a), _mm_castsi128_ph(b))); -} - -// psub - -template <> -EIGEN_STRONG_INLINE Packet32h psub<Packet32h>(const Packet32h& a, const Packet32h& b) { - return _mm512_sub_ph(a, b); -} - -template <> -EIGEN_STRONG_INLINE Packet16h psub<Packet16h>(const Packet16h& a, const Packet16h& b) { - return _mm256_castph_si256(_mm256_sub_ph(_mm256_castsi256_ph(a), _mm256_castsi256_ph(b))); -} - -template <> -EIGEN_STRONG_INLINE Packet8h psub<Packet8h>(const Packet8h& a, const Packet8h& b) { - return _mm_castph_si128(_mm_sub_ph(_mm_castsi128_ph(a), _mm_castsi128_ph(b))); -} - -// pmul - -template <> -EIGEN_STRONG_INLINE Packet32h pmul<Packet32h>(const Packet32h& a, const Packet32h& b) { - return _mm512_mul_ph(a, b); -} - -template <> -EIGEN_STRONG_INLINE Packet16h pmul<Packet16h>(const Packet16h& a, const Packet16h& b) { - return _mm256_castph_si256(_mm256_mul_ph(_mm256_castsi256_ph(a), _mm256_castsi256_ph(b))); -} - -template <> -EIGEN_STRONG_INLINE Packet8h pmul<Packet8h>(const Packet8h& a, const Packet8h& b) { - return _mm_castph_si128(_mm_mul_ph(_mm_castsi128_ph(a), _mm_castsi128_ph(b))); -} - -// pdiv - -template <> -EIGEN_STRONG_INLINE Packet32h pdiv<Packet32h>(const Packet32h& a, const Packet32h& b) { - return _mm512_div_ph(a, b); -} - -template <> -EIGEN_STRONG_INLINE Packet16h pdiv<Packet16h>(const Packet16h& a, const Packet16h& b) { - return _mm256_castph_si256(_mm256_div_ph(_mm256_castsi256_ph(a), _mm256_castsi256_ph(b))); -} - -template <> -EIGEN_STRONG_INLINE Packet8h pdiv<Packet8h>(const Packet8h& a, const Packet8h& b) { - return _mm_castph_si128(_mm_div_ph(_mm_castsi128_ph(a), _mm_castsi128_ph(b))); -} - -// pround - -template <> -EIGEN_STRONG_INLINE Packet32h pround<Packet32h>(const Packet32h& a) { - // Work-around for default std::round rounding mode. - - // Mask for the sign bit - const Packet32h signMask = pset1frombits<Packet32h>(static_cast<numext::uint16_t>(0x8000u)); - // The largest half-preicision float less than 0.5 - const Packet32h prev0dot5 = pset1frombits<Packet32h>(static_cast<numext::uint16_t>(0x37FFu)); - - return _mm512_roundscale_ph(padd(por(pand(a, signMask), prev0dot5), a), _MM_FROUND_TO_ZERO); -} - -// print - -template <> -EIGEN_STRONG_INLINE Packet32h print<Packet32h>(const Packet32h& a) { - return _mm512_roundscale_ph(a, _MM_FROUND_CUR_DIRECTION); -} - -// pceil - -template <> -EIGEN_STRONG_INLINE Packet32h pceil<Packet32h>(const Packet32h& a) { - return _mm512_roundscale_ph(a, _MM_FROUND_TO_POS_INF); -} - -// pfloor - -template <> -EIGEN_STRONG_INLINE Packet32h pfloor<Packet32h>(const Packet32h& a) { - return _mm512_roundscale_ph(a, _MM_FROUND_TO_NEG_INF); -} - -// predux -template <> -EIGEN_STRONG_INLINE half predux<Packet32h>(const Packet32h& a) { - return (half)_mm512_reduce_add_ph(a); -} - -template <> -EIGEN_STRONG_INLINE half predux<Packet16h>(const Packet16h& a) { - return (half)_mm256_reduce_add_ph(_mm256_castsi256_ph(a)); -} - -template <> -EIGEN_STRONG_INLINE half predux<Packet8h>(const Packet8h& a) { - return (half)_mm_reduce_add_ph(_mm_castsi128_ph(a)); -} - -// predux_half_dowto4 -template <> -EIGEN_STRONG_INLINE Packet16h predux_half_dowto4<Packet32h>(const Packet32h& a) { -#ifdef EIGEN_VECTORIZE_AVX512DQ - __m256i lowHalf = _mm256_castps_si256(_mm512_extractf32x8_ps(_mm512_castph_ps(a), 0)); - __m256i highHalf = _mm256_castps_si256(_mm512_extractf32x8_ps(_mm512_castph_ps(a), 1)); - - return Packet16h(padd<Packet16h>(lowHalf, highHalf)); -#else - Eigen::half data[32]; - _mm512_storeu_ph(data, a); - - __m256i lowHalf = _mm256_castph_si256(_mm256_loadu_ph(data)); - __m256i highHalf = _mm256_castph_si256(_mm256_loadu_ph(data + 16)); - - return Packet16h(padd<Packet16h>(lowHalf, highHalf)); -#endif -} - -// predux_max - -// predux_min - -// predux_mul - -#ifdef EIGEN_VECTORIZE_FMA - -// pmadd - -template <> -EIGEN_STRONG_INLINE Packet32h pmadd(const Packet32h& a, const Packet32h& b, const Packet32h& c) { - return _mm512_fmadd_ph(a, b, c); -} - -template <> -EIGEN_STRONG_INLINE Packet16h pmadd(const Packet16h& a, const Packet16h& b, const Packet16h& c) { - return _mm256_castph_si256(_mm256_fmadd_ph(_mm256_castsi256_ph(a), _mm256_castsi256_ph(b), _mm256_castsi256_ph(c))); -} - -template <> -EIGEN_STRONG_INLINE Packet8h pmadd(const Packet8h& a, const Packet8h& b, const Packet8h& c) { - return _mm_castph_si128(_mm_fmadd_ph(_mm_castsi128_ph(a), _mm_castsi128_ph(b), _mm_castsi128_ph(c))); -} - -// pmsub - -template <> -EIGEN_STRONG_INLINE Packet32h pmsub(const Packet32h& a, const Packet32h& b, const Packet32h& c) { - return _mm512_fmsub_ph(a, b, c); -} - -template <> -EIGEN_STRONG_INLINE Packet16h pmsub(const Packet16h& a, const Packet16h& b, const Packet16h& c) { - return _mm256_castph_si256(_mm256_fmsub_ph(_mm256_castsi256_ph(a), _mm256_castsi256_ph(b), _mm256_castsi256_ph(c))); -} - -template <> -EIGEN_STRONG_INLINE Packet8h pmsub(const Packet8h& a, const Packet8h& b, const Packet8h& c) { - return _mm_castph_si128(_mm_fmsub_ph(_mm_castsi128_ph(a), _mm_castsi128_ph(b), _mm_castsi128_ph(c))); -} - -// pnmadd - -template <> -EIGEN_STRONG_INLINE Packet32h pnmadd(const Packet32h& a, const Packet32h& b, const Packet32h& c) { - return _mm512_fnmadd_ph(a, b, c); -} - -template <> -EIGEN_STRONG_INLINE Packet16h pnmadd(const Packet16h& a, const Packet16h& b, const Packet16h& c) { - return _mm256_castph_si256(_mm256_fnmadd_ph(_mm256_castsi256_ph(a), _mm256_castsi256_ph(b), _mm256_castsi256_ph(c))); -} - -template <> -EIGEN_STRONG_INLINE Packet8h pnmadd(const Packet8h& a, const Packet8h& b, const Packet8h& c) { - return _mm_castph_si128(_mm_fnmadd_ph(_mm_castsi128_ph(a), _mm_castsi128_ph(b), _mm_castsi128_ph(c))); -} - -// pnmsub - -template <> -EIGEN_STRONG_INLINE Packet32h pnmsub(const Packet32h& a, const Packet32h& b, const Packet32h& c) { - return _mm512_fnmsub_ph(a, b, c); -} - -template <> -EIGEN_STRONG_INLINE Packet16h pnmsub(const Packet16h& a, const Packet16h& b, const Packet16h& c) { - return _mm256_castph_si256(_mm256_fnmsub_ph(_mm256_castsi256_ph(a), _mm256_castsi256_ph(b), _mm256_castsi256_ph(c))); -} - -template <> -EIGEN_STRONG_INLINE Packet8h pnmsub(const Packet8h& a, const Packet8h& b, const Packet8h& c) { - return _mm_castph_si128(_mm_fnmsub_ph(_mm_castsi128_ph(a), _mm_castsi128_ph(b), _mm_castsi128_ph(c))); -} - -#endif - -// pnegate - -template <> -EIGEN_STRONG_INLINE Packet32h pnegate<Packet32h>(const Packet32h& a) { - return _mm512_sub_ph(_mm512_set1_ph(0.0), a); -} - -// pconj - -template <> -EIGEN_STRONG_INLINE Packet32h pconj<Packet32h>(const Packet32h& a) { - return a; -} - -// psqrt - -template <> -EIGEN_STRONG_INLINE Packet32h psqrt<Packet32h>(const Packet32h& a) { - return _mm512_sqrt_ph(a); -} - -// prsqrt - -template <> -EIGEN_STRONG_INLINE Packet32h prsqrt<Packet32h>(const Packet32h& a) { - return _mm512_rsqrt_ph(a); -} - -// preciprocal - -template <> -EIGEN_STRONG_INLINE Packet32h preciprocal<Packet32h>(const Packet32h& a) { - return _mm512_rcp_ph(a); -} - -// ptranspose - -EIGEN_DEVICE_FUNC inline void ptranspose(PacketBlock<Packet32h, 32>& a) { - __m512i t[32]; - - EIGEN_UNROLL_LOOP - for (int i = 0; i < 16; i++) { - t[2 * i] = _mm512_unpacklo_epi16(_mm512_castph_si512(a.packet[2 * i]), _mm512_castph_si512(a.packet[2 * i + 1])); - t[2 * i + 1] = - _mm512_unpackhi_epi16(_mm512_castph_si512(a.packet[2 * i]), _mm512_castph_si512(a.packet[2 * i + 1])); - } - - __m512i p[32]; - - EIGEN_UNROLL_LOOP - for (int i = 0; i < 8; i++) { - p[4 * i] = _mm512_unpacklo_epi32(t[4 * i], t[4 * i + 2]); - p[4 * i + 1] = _mm512_unpackhi_epi32(t[4 * i], t[4 * i + 2]); - p[4 * i + 2] = _mm512_unpacklo_epi32(t[4 * i + 1], t[4 * i + 3]); - p[4 * i + 3] = _mm512_unpackhi_epi32(t[4 * i + 1], t[4 * i + 3]); - } - - __m512i q[32]; - - EIGEN_UNROLL_LOOP - for (int i = 0; i < 4; i++) { - q[8 * i] = _mm512_unpacklo_epi64(p[8 * i], p[8 * i + 4]); - q[8 * i + 1] = _mm512_unpackhi_epi64(p[8 * i], p[8 * i + 4]); - q[8 * i + 2] = _mm512_unpacklo_epi64(p[8 * i + 1], p[8 * i + 5]); - q[8 * i + 3] = _mm512_unpackhi_epi64(p[8 * i + 1], p[8 * i + 5]); - q[8 * i + 4] = _mm512_unpacklo_epi64(p[8 * i + 2], p[8 * i + 6]); - q[8 * i + 5] = _mm512_unpackhi_epi64(p[8 * i + 2], p[8 * i + 6]); - q[8 * i + 6] = _mm512_unpacklo_epi64(p[8 * i + 3], p[8 * i + 7]); - q[8 * i + 7] = _mm512_unpackhi_epi64(p[8 * i + 3], p[8 * i + 7]); - } - - __m512i f[32]; - -#define PACKET32H_TRANSPOSE_HELPER(X, Y) \ - do { \ - f[Y * 8] = _mm512_inserti32x4(f[Y * 8], _mm512_extracti32x4_epi32(q[X * 8], Y), X); \ - f[Y * 8 + 1] = _mm512_inserti32x4(f[Y * 8 + 1], _mm512_extracti32x4_epi32(q[X * 8 + 1], Y), X); \ - f[Y * 8 + 2] = _mm512_inserti32x4(f[Y * 8 + 2], _mm512_extracti32x4_epi32(q[X * 8 + 2], Y), X); \ - f[Y * 8 + 3] = _mm512_inserti32x4(f[Y * 8 + 3], _mm512_extracti32x4_epi32(q[X * 8 + 3], Y), X); \ - f[Y * 8 + 4] = _mm512_inserti32x4(f[Y * 8 + 4], _mm512_extracti32x4_epi32(q[X * 8 + 4], Y), X); \ - f[Y * 8 + 5] = _mm512_inserti32x4(f[Y * 8 + 5], _mm512_extracti32x4_epi32(q[X * 8 + 5], Y), X); \ - f[Y * 8 + 6] = _mm512_inserti32x4(f[Y * 8 + 6], _mm512_extracti32x4_epi32(q[X * 8 + 6], Y), X); \ - f[Y * 8 + 7] = _mm512_inserti32x4(f[Y * 8 + 7], _mm512_extracti32x4_epi32(q[X * 8 + 7], Y), X); \ - } while (false); - - PACKET32H_TRANSPOSE_HELPER(0, 0); - PACKET32H_TRANSPOSE_HELPER(1, 1); - PACKET32H_TRANSPOSE_HELPER(2, 2); - PACKET32H_TRANSPOSE_HELPER(3, 3); - - PACKET32H_TRANSPOSE_HELPER(1, 0); - PACKET32H_TRANSPOSE_HELPER(2, 0); - PACKET32H_TRANSPOSE_HELPER(3, 0); - PACKET32H_TRANSPOSE_HELPER(2, 1); - PACKET32H_TRANSPOSE_HELPER(3, 1); - PACKET32H_TRANSPOSE_HELPER(3, 2); - - PACKET32H_TRANSPOSE_HELPER(0, 1); - PACKET32H_TRANSPOSE_HELPER(0, 2); - PACKET32H_TRANSPOSE_HELPER(0, 3); - PACKET32H_TRANSPOSE_HELPER(1, 2); - PACKET32H_TRANSPOSE_HELPER(1, 3); - PACKET32H_TRANSPOSE_HELPER(2, 3); - -#undef PACKET32H_TRANSPOSE_HELPER - - EIGEN_UNROLL_LOOP - for (int i = 0; i < 32; i++) { - a.packet[i] = _mm512_castsi512_ph(f[i]); - } -} - -EIGEN_DEVICE_FUNC inline void ptranspose(PacketBlock<Packet32h, 4>& a) { - __m512i p0, p1, p2, p3, t0, t1, t2, t3, a0, a1, a2, a3; - t0 = _mm512_unpacklo_epi16(_mm512_castph_si512(a.packet[0]), _mm512_castph_si512(a.packet[1])); - t1 = _mm512_unpackhi_epi16(_mm512_castph_si512(a.packet[0]), _mm512_castph_si512(a.packet[1])); - t2 = _mm512_unpacklo_epi16(_mm512_castph_si512(a.packet[2]), _mm512_castph_si512(a.packet[3])); - t3 = _mm512_unpackhi_epi16(_mm512_castph_si512(a.packet[2]), _mm512_castph_si512(a.packet[3])); - - p0 = _mm512_unpacklo_epi32(t0, t2); - p1 = _mm512_unpackhi_epi32(t0, t2); - p2 = _mm512_unpacklo_epi32(t1, t3); - p3 = _mm512_unpackhi_epi32(t1, t3); - - a0 = p0; - a1 = p1; - a2 = p2; - a3 = p3; - - a0 = _mm512_inserti32x4(a0, _mm512_extracti32x4_epi32(p1, 0), 1); - a1 = _mm512_inserti32x4(a1, _mm512_extracti32x4_epi32(p0, 1), 0); - - a0 = _mm512_inserti32x4(a0, _mm512_extracti32x4_epi32(p2, 0), 2); - a2 = _mm512_inserti32x4(a2, _mm512_extracti32x4_epi32(p0, 2), 0); - - a0 = _mm512_inserti32x4(a0, _mm512_extracti32x4_epi32(p3, 0), 3); - a3 = _mm512_inserti32x4(a3, _mm512_extracti32x4_epi32(p0, 3), 0); - - a1 = _mm512_inserti32x4(a1, _mm512_extracti32x4_epi32(p2, 1), 2); - a2 = _mm512_inserti32x4(a2, _mm512_extracti32x4_epi32(p1, 2), 1); - - a2 = _mm512_inserti32x4(a2, _mm512_extracti32x4_epi32(p3, 2), 3); - a3 = _mm512_inserti32x4(a3, _mm512_extracti32x4_epi32(p2, 3), 2); - - a1 = _mm512_inserti32x4(a1, _mm512_extracti32x4_epi32(p3, 1), 3); - a3 = _mm512_inserti32x4(a3, _mm512_extracti32x4_epi32(p1, 3), 1); - - a.packet[0] = _mm512_castsi512_ph(a0); - a.packet[1] = _mm512_castsi512_ph(a1); - a.packet[2] = _mm512_castsi512_ph(a2); - a.packet[3] = _mm512_castsi512_ph(a3); -} - -// preverse - -template <> -EIGEN_STRONG_INLINE Packet32h preverse(const Packet32h& a) { - return _mm512_permutexvar_ph(_mm512_set_epi16(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, - 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31), - a); -} - -// pscatter - -template <> -EIGEN_STRONG_INLINE void pscatter<half, Packet32h>(half* to, const Packet32h& from, Index stride) { - EIGEN_ALIGN64 half aux[32]; - pstore(aux, from); - - EIGEN_UNROLL_LOOP - for (int i = 0; i < 32; i++) { - to[stride * i] = aux[i]; - } -} - -// pgather - -template <> -EIGEN_STRONG_INLINE Packet32h pgather<Eigen::half, Packet32h>(const Eigen::half* from, Index stride) { - return _mm512_castsi512_ph(_mm512_set_epi16( - from[31 * stride].x, from[30 * stride].x, from[29 * stride].x, from[28 * stride].x, from[27 * stride].x, - from[26 * stride].x, from[25 * stride].x, from[24 * stride].x, from[23 * stride].x, from[22 * stride].x, - from[21 * stride].x, from[20 * stride].x, from[19 * stride].x, from[18 * stride].x, from[17 * stride].x, - from[16 * stride].x, from[15 * stride].x, from[14 * stride].x, from[13 * stride].x, from[12 * stride].x, - from[11 * stride].x, from[10 * stride].x, from[9 * stride].x, from[8 * stride].x, from[7 * stride].x, - from[6 * stride].x, from[5 * stride].x, from[4 * stride].x, from[3 * stride].x, from[2 * stride].x, - from[1 * stride].x, from[0 * stride].x)); -} - -template <> -EIGEN_STRONG_INLINE Packet16h pcos<Packet16h>(const Packet16h&); -template <> -EIGEN_STRONG_INLINE Packet16h psin<Packet16h>(const Packet16h&); -template <> -EIGEN_STRONG_INLINE Packet16h plog<Packet16h>(const Packet16h&); -template <> -EIGEN_STRONG_INLINE Packet16h plog2<Packet16h>(const Packet16h&); -template <> -EIGEN_STRONG_INLINE Packet16h plog1p<Packet16h>(const Packet16h&); -template <> -EIGEN_STRONG_INLINE Packet16h pexp<Packet16h>(const Packet16h&); -template <> -EIGEN_STRONG_INLINE Packet16h pexpm1<Packet16h>(const Packet16h&); -template <> -EIGEN_STRONG_INLINE Packet16h ptanh<Packet16h>(const Packet16h&); -template <> -EIGEN_STRONG_INLINE Packet16h pfrexp<Packet16h>(const Packet16h&, Packet16h&); -template <> -EIGEN_STRONG_INLINE Packet16h pldexp<Packet16h>(const Packet16h&, const Packet16h&); - -EIGEN_STRONG_INLINE Packet32h combine2Packet16h(const Packet16h& a, const Packet16h& b) { - __m512d result = _mm512_undefined_pd(); - result = _mm512_insertf64x4(result, _mm256_castsi256_pd(a), 0); - result = _mm512_insertf64x4(result, _mm256_castsi256_pd(b), 1); - return _mm512_castpd_ph(result); -} - -EIGEN_STRONG_INLINE void extract2Packet16h(const Packet32h& x, Packet16h& a, Packet16h& b) { - a = _mm256_castpd_si256(_mm512_extractf64x4_pd(_mm512_castph_pd(x), 0)); - b = _mm256_castpd_si256(_mm512_extractf64x4_pd(_mm512_castph_pd(x), 1)); -} - -// psin -template <> -EIGEN_STRONG_INLINE Packet32h psin<Packet32h>(const Packet32h& a) { - Packet16h low; - Packet16h high; - extract2Packet16h(a, low, high); - - Packet16h lowOut = psin(low); - Packet16h highOut = psin(high); - - return combine2Packet16h(lowOut, highOut); -} - -// pcos -template <> -EIGEN_STRONG_INLINE Packet32h pcos<Packet32h>(const Packet32h& a) { - Packet16h low; - Packet16h high; - extract2Packet16h(a, low, high); - - Packet16h lowOut = pcos(low); - Packet16h highOut = pcos(high); - - return combine2Packet16h(lowOut, highOut); -} - -// plog -template <> -EIGEN_STRONG_INLINE Packet32h plog<Packet32h>(const Packet32h& a) { - Packet16h low; - Packet16h high; - extract2Packet16h(a, low, high); - - Packet16h lowOut = plog(low); - Packet16h highOut = plog(high); - - return combine2Packet16h(lowOut, highOut); -} - -// plog2 -template <> -EIGEN_STRONG_INLINE Packet32h plog2<Packet32h>(const Packet32h& a) { - Packet16h low; - Packet16h high; - extract2Packet16h(a, low, high); - - Packet16h lowOut = plog2(low); - Packet16h highOut = plog2(high); - - return combine2Packet16h(lowOut, highOut); -} - -// plog1p -template <> -EIGEN_STRONG_INLINE Packet32h plog1p<Packet32h>(const Packet32h& a) { - Packet16h low; - Packet16h high; - extract2Packet16h(a, low, high); - - Packet16h lowOut = plog1p(low); - Packet16h highOut = plog1p(high); - - return combine2Packet16h(lowOut, highOut); -} - -// pexp -template <> -EIGEN_STRONG_INLINE Packet32h pexp<Packet32h>(const Packet32h& a) { - Packet16h low; - Packet16h high; - extract2Packet16h(a, low, high); - - Packet16h lowOut = pexp(low); - Packet16h highOut = pexp(high); - - return combine2Packet16h(lowOut, highOut); -} - -// pexpm1 -template <> -EIGEN_STRONG_INLINE Packet32h pexpm1<Packet32h>(const Packet32h& a) { - Packet16h low; - Packet16h high; - extract2Packet16h(a, low, high); - - Packet16h lowOut = pexpm1(low); - Packet16h highOut = pexpm1(high); - - return combine2Packet16h(lowOut, highOut); -} - -// ptanh -template <> -EIGEN_STRONG_INLINE Packet32h ptanh<Packet32h>(const Packet32h& a) { - Packet16h low; - Packet16h high; - extract2Packet16h(a, low, high); - - Packet16h lowOut = ptanh(low); - Packet16h highOut = ptanh(high); - - return combine2Packet16h(lowOut, highOut); -} - -// pfrexp -template <> -EIGEN_STRONG_INLINE Packet32h pfrexp<Packet32h>(const Packet32h& a, Packet32h& exponent) { - Packet16h low; - Packet16h high; - extract2Packet16h(a, low, high); - - Packet16h exp1 = _mm256_undefined_si256(); - Packet16h exp2 = _mm256_undefined_si256(); - - Packet16h lowOut = pfrexp(low, exp1); - Packet16h highOut = pfrexp(high, exp2); - - exponent = combine2Packet16h(exp1, exp2); - - return combine2Packet16h(lowOut, highOut); -} - -// pldexp -template <> -EIGEN_STRONG_INLINE Packet32h pldexp<Packet32h>(const Packet32h& a, const Packet32h& exponent) { - Packet16h low; - Packet16h high; - extract2Packet16h(a, low, high); - - Packet16h exp1; - Packet16h exp2; - extract2Packet16h(exponent, exp1, exp2); - - Packet16h lowOut = pldexp(low, exp1); - Packet16h highOut = pldexp(high, exp2); - - return combine2Packet16h(lowOut, highOut); -} - -} // end namespace internal -} // end namespace Eigen - -#endif // EIGEN_PACKET_MATH_FP16_AVX512_H +#include "../../InternalHeaderCheck.h" + +namespace Eigen { + +namespace internal { + +typedef __m512h Packet32h; +typedef eigen_packet_wrapper<__m256i, 1> Packet16h; +typedef eigen_packet_wrapper<__m128i, 2> Packet8h; + +template <> +struct is_arithmetic<Packet8h> { + enum { value = true }; +}; + +template <> +struct packet_traits<half> : default_packet_traits { + typedef Packet32h type; + typedef Packet16h half; + enum { + Vectorizable = 1, + AlignedOnScalar = 1, + size = 32, + + HasCmp = 1, + HasAdd = 1, + HasSub = 1, + HasMul = 1, + HasDiv = 1, + HasNegate = 1, + HasAbs = 1, + HasAbs2 = 0, + HasMin = 1, + HasMax = 1, + HasConj = 1, + HasSetLinear = 0, + HasLog = 1, + HasLog1p = 1, + HasExp = 1, + HasExpm1 = 1, + HasSqrt = 1, + HasRsqrt = 1, + // These ones should be implemented in future + HasBessel = 0, + HasNdtri = 0, + HasSin = EIGEN_FAST_MATH, + HasCos = EIGEN_FAST_MATH, + HasTanh = EIGEN_FAST_MATH, + HasErf = 0, // EIGEN_FAST_MATH, + HasBlend = 0, + HasRound = 1, + HasFloor = 1, + HasCeil = 1, + HasRint = 1 + }; +}; + +template <> +struct unpacket_traits<Packet32h> { + typedef Eigen::half type; + typedef Packet16h half; + enum { + size = 32, + alignment = Aligned64, + vectorizable = true, + masked_load_available = false, + masked_store_available = false + }; +}; + +template <> +struct unpacket_traits<Packet16h> { + typedef Eigen::half type; + typedef Packet8h half; + enum { + size = 16, + alignment = Aligned32, + vectorizable = true, + masked_load_available = false, + masked_store_available = false + }; +}; + +template <> +struct unpacket_traits<Packet8h> { + typedef Eigen::half type; + typedef Packet8h half; + enum { + size = 8, + alignment = Aligned16, + vectorizable = true, + masked_load_available = false, + masked_store_available = false + }; +}; + +// Memory functions + +// pset1 + +template <> +EIGEN_STRONG_INLINE Packet32h pset1<Packet32h>(const Eigen::half& from) { + return _mm512_set1_ph(static_cast<_Float16>(from)); +} + +// pset1frombits +template <> +EIGEN_STRONG_INLINE Packet32h pset1frombits<Packet32h>(unsigned short from) { + return _mm512_castsi512_ph(_mm512_set1_epi16(from)); +} + +// pfirst + +template <> +EIGEN_STRONG_INLINE Eigen::half pfirst<Packet32h>(const Packet32h& from) { +#ifdef EIGEN_VECTORIZE_AVX512DQ + return half_impl::raw_uint16_to_half( + static_cast<unsigned short>(_mm256_extract_epi16(_mm512_extracti32x8_epi32(_mm512_castph_si512(from), 0), 0))); +#else + Eigen::half dest[32]; + _mm512_storeu_ph(dest, from); + return dest[0]; +#endif +} + +// pload + +template <> +EIGEN_STRONG_INLINE Packet32h pload<Packet32h>(const Eigen::half* from) { + EIGEN_DEBUG_ALIGNED_LOAD return _mm512_load_ph(from); +} + +// ploadu + +template <> +EIGEN_STRONG_INLINE Packet32h ploadu<Packet32h>(const Eigen::half* from) { + EIGEN_DEBUG_UNALIGNED_LOAD return _mm512_loadu_ph(from); +} + +// pstore + +template <> +EIGEN_STRONG_INLINE void pstore<half>(Eigen::half* to, const Packet32h& from) { + EIGEN_DEBUG_ALIGNED_STORE _mm512_store_ph(to, from); +} + +// pstoreu + +template <> +EIGEN_STRONG_INLINE void pstoreu<half>(Eigen::half* to, const Packet32h& from) { + EIGEN_DEBUG_UNALIGNED_STORE _mm512_storeu_ph(to, from); +} + +// ploaddup +template <> +EIGEN_STRONG_INLINE Packet32h ploaddup<Packet32h>(const Eigen::half* from) { + __m512h a = _mm512_castph256_ph512(_mm256_loadu_ph(from)); + return _mm512_permutexvar_ph(_mm512_set_epi16(15, 15, 14, 14, 13, 13, 12, 12, 11, 11, 10, 10, 9, 9, 8, 8, 7, 7, 6, 6, + 5, 5, 4, 4, 3, 3, 2, 2, 1, 1, 0, 0), + a); +} + +// ploadquad +template <> +EIGEN_STRONG_INLINE Packet32h ploadquad<Packet32h>(const Eigen::half* from) { + __m512h a = _mm512_castph128_ph512(_mm_loadu_ph(from)); + return _mm512_permutexvar_ph( + _mm512_set_epi16(7, 7, 7, 7, 6, 6, 6, 6, 5, 5, 5, 5, 4, 4, 4, 4, 3, 3, 3, 3, 2, 2, 2, 2, 1, 1, 1, 1, 0, 0, 0, 0), + a); +} + +// pabs + +template <> +EIGEN_STRONG_INLINE Packet32h pabs<Packet32h>(const Packet32h& a) { + return _mm512_abs_ph(a); +} + +// psignbit + +template <> +EIGEN_STRONG_INLINE Packet32h psignbit<Packet32h>(const Packet32h& a) { + return _mm512_castsi512_ph(_mm512_srai_epi16(_mm512_castph_si512(a), 15)); +} + +// pmin + +template <> +EIGEN_STRONG_INLINE Packet32h pmin<Packet32h>(const Packet32h& a, const Packet32h& b) { + return _mm512_min_ph(a, b); +} + +// pmax + +template <> +EIGEN_STRONG_INLINE Packet32h pmax<Packet32h>(const Packet32h& a, const Packet32h& b) { + return _mm512_max_ph(a, b); +} + +// plset +template <> +EIGEN_STRONG_INLINE Packet32h plset<Packet32h>(const half& a) { + return _mm512_add_ph(_mm512_set1_ph(a), + _mm512_set_ph(31.0f, 30.0f, 29.0f, 28.0f, 27.0f, 26.0f, 25.0f, 24.0f, 23.0f, 22.0f, 21.0f, 20.0f, + 19.0f, 18.0f, 17.0f, 16.0f, 15.0f, 14.0f, 13.0f, 12.0f, 11.0f, 10.0f, 9.0f, 8.0f, + 7.0f, 6.0f, 5.0f, 4.0f, 3.0f, 2.0f, 1.0f, 0.0f)); +} + +// por + +template <> +EIGEN_STRONG_INLINE Packet32h por(const Packet32h& a, const Packet32h& b) { + return _mm512_castsi512_ph(_mm512_or_si512(_mm512_castph_si512(a), _mm512_castph_si512(b))); +} + +// pxor + +template <> +EIGEN_STRONG_INLINE Packet32h pxor(const Packet32h& a, const Packet32h& b) { + return _mm512_castsi512_ph(_mm512_xor_si512(_mm512_castph_si512(a), _mm512_castph_si512(b))); +} + +// pand + +template <> +EIGEN_STRONG_INLINE Packet32h pand(const Packet32h& a, const Packet32h& b) { + return _mm512_castsi512_ph(_mm512_and_si512(_mm512_castph_si512(a), _mm512_castph_si512(b))); +} + +// pandnot + +template <> +EIGEN_STRONG_INLINE Packet32h pandnot(const Packet32h& a, const Packet32h& b) { + return _mm512_castsi512_ph(_mm512_andnot_si512(_mm512_castph_si512(b), _mm512_castph_si512(a))); +} + +// pselect + +template <> +EIGEN_DEVICE_FUNC inline Packet32h pselect(const Packet32h& mask, const Packet32h& a, const Packet32h& b) { + __mmask32 mask32 = _mm512_cmp_epi16_mask(_mm512_castph_si512(mask), _mm512_setzero_epi32(), _MM_CMPINT_EQ); + return _mm512_mask_blend_ph(mask32, a, b); +} + +// pcmp_eq + +template <> +EIGEN_STRONG_INLINE Packet32h pcmp_eq(const Packet32h& a, const Packet32h& b) { + __mmask32 mask = _mm512_cmp_ph_mask(a, b, _CMP_EQ_OQ); + return _mm512_castsi512_ph(_mm512_mask_set1_epi16(_mm512_set1_epi32(0), mask, 0xffffu)); +} + +// pcmp_le + +template <> +EIGEN_STRONG_INLINE Packet32h pcmp_le(const Packet32h& a, const Packet32h& b) { + __mmask32 mask = _mm512_cmp_ph_mask(a, b, _CMP_LE_OQ); + return _mm512_castsi512_ph(_mm512_mask_set1_epi16(_mm512_set1_epi32(0), mask, 0xffffu)); +} + +// pcmp_lt + +template <> +EIGEN_STRONG_INLINE Packet32h pcmp_lt(const Packet32h& a, const Packet32h& b) { + __mmask32 mask = _mm512_cmp_ph_mask(a, b, _CMP_LT_OQ); + return _mm512_castsi512_ph(_mm512_mask_set1_epi16(_mm512_set1_epi32(0), mask, 0xffffu)); +} + +// pcmp_lt_or_nan + +template <> +EIGEN_STRONG_INLINE Packet32h pcmp_lt_or_nan(const Packet32h& a, const Packet32h& b) { + __mmask32 mask = _mm512_cmp_ph_mask(a, b, _CMP_NGE_UQ); + return _mm512_castsi512_ph(_mm512_mask_set1_epi16(_mm512_set1_epi16(0), mask, 0xffffu)); +} + +// padd + +template <> +EIGEN_STRONG_INLINE Packet32h padd<Packet32h>(const Packet32h& a, const Packet32h& b) { + return _mm512_add_ph(a, b); +} + +template <> +EIGEN_STRONG_INLINE Packet16h padd<Packet16h>(const Packet16h& a, const Packet16h& b) { + return _mm256_castph_si256(_mm256_add_ph(_mm256_castsi256_ph(a), _mm256_castsi256_ph(b))); +} + +template <> +EIGEN_STRONG_INLINE Packet8h padd<Packet8h>(const Packet8h& a, const Packet8h& b) { + return _mm_castph_si128(_mm_add_ph(_mm_castsi128_ph(a), _mm_castsi128_ph(b))); +} + +// psub + +template <> +EIGEN_STRONG_INLINE Packet32h psub<Packet32h>(const Packet32h& a, const Packet32h& b) { + return _mm512_sub_ph(a, b); +} + +template <> +EIGEN_STRONG_INLINE Packet16h psub<Packet16h>(const Packet16h& a, const Packet16h& b) { + return _mm256_castph_si256(_mm256_sub_ph(_mm256_castsi256_ph(a), _mm256_castsi256_ph(b))); +} + +template <> +EIGEN_STRONG_INLINE Packet8h psub<Packet8h>(const Packet8h& a, const Packet8h& b) { + return _mm_castph_si128(_mm_sub_ph(_mm_castsi128_ph(a), _mm_castsi128_ph(b))); +} + +// pmul + +template <> +EIGEN_STRONG_INLINE Packet32h pmul<Packet32h>(const Packet32h& a, const Packet32h& b) { + return _mm512_mul_ph(a, b); +} + +template <> +EIGEN_STRONG_INLINE Packet16h pmul<Packet16h>(const Packet16h& a, const Packet16h& b) { + return _mm256_castph_si256(_mm256_mul_ph(_mm256_castsi256_ph(a), _mm256_castsi256_ph(b))); +} + +template <> +EIGEN_STRONG_INLINE Packet8h pmul<Packet8h>(const Packet8h& a, const Packet8h& b) { + return _mm_castph_si128(_mm_mul_ph(_mm_castsi128_ph(a), _mm_castsi128_ph(b))); +} + +// pdiv + +template <> +EIGEN_STRONG_INLINE Packet32h pdiv<Packet32h>(const Packet32h& a, const Packet32h& b) { + return _mm512_div_ph(a, b); +} + +template <> +EIGEN_STRONG_INLINE Packet16h pdiv<Packet16h>(const Packet16h& a, const Packet16h& b) { + return _mm256_castph_si256(_mm256_div_ph(_mm256_castsi256_ph(a), _mm256_castsi256_ph(b))); +} + +template <> +EIGEN_STRONG_INLINE Packet8h pdiv<Packet8h>(const Packet8h& a, const Packet8h& b) { + return _mm_castph_si128(_mm_div_ph(_mm_castsi128_ph(a), _mm_castsi128_ph(b))); +} + +// pround + +template <> +EIGEN_STRONG_INLINE Packet32h pround<Packet32h>(const Packet32h& a) { + // Work-around for default std::round rounding mode. + + // Mask for the sign bit + const Packet32h signMask = pset1frombits<Packet32h>(static_cast<numext::uint16_t>(0x8000u)); + // The largest half-preicision float less than 0.5 + const Packet32h prev0dot5 = pset1frombits<Packet32h>(static_cast<numext::uint16_t>(0x37FFu)); + + return _mm512_roundscale_ph(padd(por(pand(a, signMask), prev0dot5), a), _MM_FROUND_TO_ZERO); +} + +// print + +template <> +EIGEN_STRONG_INLINE Packet32h print<Packet32h>(const Packet32h& a) { + return _mm512_roundscale_ph(a, _MM_FROUND_CUR_DIRECTION); +} + +// pceil + +template <> +EIGEN_STRONG_INLINE Packet32h pceil<Packet32h>(const Packet32h& a) { + return _mm512_roundscale_ph(a, _MM_FROUND_TO_POS_INF); +} + +// pfloor + +template <> +EIGEN_STRONG_INLINE Packet32h pfloor<Packet32h>(const Packet32h& a) { + return _mm512_roundscale_ph(a, _MM_FROUND_TO_NEG_INF); +} + +// predux +template <> +EIGEN_STRONG_INLINE half predux<Packet32h>(const Packet32h& a) { + return (half)_mm512_reduce_add_ph(a); +} + +template <> +EIGEN_STRONG_INLINE half predux<Packet16h>(const Packet16h& a) { + return (half)_mm256_reduce_add_ph(_mm256_castsi256_ph(a)); +} + +template <> +EIGEN_STRONG_INLINE half predux<Packet8h>(const Packet8h& a) { + return (half)_mm_reduce_add_ph(_mm_castsi128_ph(a)); +} + +// predux_half_dowto4 +template <> +EIGEN_STRONG_INLINE Packet16h predux_half_dowto4<Packet32h>(const Packet32h& a) { +#ifdef EIGEN_VECTORIZE_AVX512DQ + __m256i lowHalf = _mm256_castps_si256(_mm512_extractf32x8_ps(_mm512_castph_ps(a), 0)); + __m256i highHalf = _mm256_castps_si256(_mm512_extractf32x8_ps(_mm512_castph_ps(a), 1)); + + return Packet16h(padd<Packet16h>(lowHalf, highHalf)); +#else + Eigen::half data[32]; + _mm512_storeu_ph(data, a); + + __m256i lowHalf = _mm256_castph_si256(_mm256_loadu_ph(data)); + __m256i highHalf = _mm256_castph_si256(_mm256_loadu_ph(data + 16)); + + return Packet16h(padd<Packet16h>(lowHalf, highHalf)); +#endif +} + +// predux_max + +// predux_min + +// predux_mul + +#ifdef EIGEN_VECTORIZE_FMA + +// pmadd + +template <> +EIGEN_STRONG_INLINE Packet32h pmadd(const Packet32h& a, const Packet32h& b, const Packet32h& c) { + return _mm512_fmadd_ph(a, b, c); +} + +template <> +EIGEN_STRONG_INLINE Packet16h pmadd(const Packet16h& a, const Packet16h& b, const Packet16h& c) { + return _mm256_castph_si256(_mm256_fmadd_ph(_mm256_castsi256_ph(a), _mm256_castsi256_ph(b), _mm256_castsi256_ph(c))); +} + +template <> +EIGEN_STRONG_INLINE Packet8h pmadd(const Packet8h& a, const Packet8h& b, const Packet8h& c) { + return _mm_castph_si128(_mm_fmadd_ph(_mm_castsi128_ph(a), _mm_castsi128_ph(b), _mm_castsi128_ph(c))); +} + +// pmsub + +template <> +EIGEN_STRONG_INLINE Packet32h pmsub(const Packet32h& a, const Packet32h& b, const Packet32h& c) { + return _mm512_fmsub_ph(a, b, c); +} + +template <> +EIGEN_STRONG_INLINE Packet16h pmsub(const Packet16h& a, const Packet16h& b, const Packet16h& c) { + return _mm256_castph_si256(_mm256_fmsub_ph(_mm256_castsi256_ph(a), _mm256_castsi256_ph(b), _mm256_castsi256_ph(c))); +} + +template <> +EIGEN_STRONG_INLINE Packet8h pmsub(const Packet8h& a, const Packet8h& b, const Packet8h& c) { + return _mm_castph_si128(_mm_fmsub_ph(_mm_castsi128_ph(a), _mm_castsi128_ph(b), _mm_castsi128_ph(c))); +} + +// pnmadd + +template <> +EIGEN_STRONG_INLINE Packet32h pnmadd(const Packet32h& a, const Packet32h& b, const Packet32h& c) { + return _mm512_fnmadd_ph(a, b, c); +} + +template <> +EIGEN_STRONG_INLINE Packet16h pnmadd(const Packet16h& a, const Packet16h& b, const Packet16h& c) { + return _mm256_castph_si256(_mm256_fnmadd_ph(_mm256_castsi256_ph(a), _mm256_castsi256_ph(b), _mm256_castsi256_ph(c))); +} + +template <> +EIGEN_STRONG_INLINE Packet8h pnmadd(const Packet8h& a, const Packet8h& b, const Packet8h& c) { + return _mm_castph_si128(_mm_fnmadd_ph(_mm_castsi128_ph(a), _mm_castsi128_ph(b), _mm_castsi128_ph(c))); +} + +// pnmsub + +template <> +EIGEN_STRONG_INLINE Packet32h pnmsub(const Packet32h& a, const Packet32h& b, const Packet32h& c) { + return _mm512_fnmsub_ph(a, b, c); +} + +template <> +EIGEN_STRONG_INLINE Packet16h pnmsub(const Packet16h& a, const Packet16h& b, const Packet16h& c) { + return _mm256_castph_si256(_mm256_fnmsub_ph(_mm256_castsi256_ph(a), _mm256_castsi256_ph(b), _mm256_castsi256_ph(c))); +} + +template <> +EIGEN_STRONG_INLINE Packet8h pnmsub(const Packet8h& a, const Packet8h& b, const Packet8h& c) { + return _mm_castph_si128(_mm_fnmsub_ph(_mm_castsi128_ph(a), _mm_castsi128_ph(b), _mm_castsi128_ph(c))); +} + +#endif + +// pnegate + +template <> +EIGEN_STRONG_INLINE Packet32h pnegate<Packet32h>(const Packet32h& a) { + return _mm512_sub_ph(_mm512_set1_ph(0.0), a); +} + +// pconj + +template <> +EIGEN_STRONG_INLINE Packet32h pconj<Packet32h>(const Packet32h& a) { + return a; +} + +// psqrt + +template <> +EIGEN_STRONG_INLINE Packet32h psqrt<Packet32h>(const Packet32h& a) { + return _mm512_sqrt_ph(a); +} + +// prsqrt + +template <> +EIGEN_STRONG_INLINE Packet32h prsqrt<Packet32h>(const Packet32h& a) { + return _mm512_rsqrt_ph(a); +} + +// preciprocal + +template <> +EIGEN_STRONG_INLINE Packet32h preciprocal<Packet32h>(const Packet32h& a) { + return _mm512_rcp_ph(a); +} + +// ptranspose + +EIGEN_DEVICE_FUNC inline void ptranspose(PacketBlock<Packet32h, 32>& a) { + __m512i t[32]; + + EIGEN_UNROLL_LOOP + for (int i = 0; i < 16; i++) { + t[2 * i] = _mm512_unpacklo_epi16(_mm512_castph_si512(a.packet[2 * i]), _mm512_castph_si512(a.packet[2 * i + 1])); + t[2 * i + 1] = + _mm512_unpackhi_epi16(_mm512_castph_si512(a.packet[2 * i]), _mm512_castph_si512(a.packet[2 * i + 1])); + } + + __m512i p[32]; + + EIGEN_UNROLL_LOOP + for (int i = 0; i < 8; i++) { + p[4 * i] = _mm512_unpacklo_epi32(t[4 * i], t[4 * i + 2]); + p[4 * i + 1] = _mm512_unpackhi_epi32(t[4 * i], t[4 * i + 2]); + p[4 * i + 2] = _mm512_unpacklo_epi32(t[4 * i + 1], t[4 * i + 3]); + p[4 * i + 3] = _mm512_unpackhi_epi32(t[4 * i + 1], t[4 * i + 3]); + } + + __m512i q[32]; + + EIGEN_UNROLL_LOOP + for (int i = 0; i < 4; i++) { + q[8 * i] = _mm512_unpacklo_epi64(p[8 * i], p[8 * i + 4]); + q[8 * i + 1] = _mm512_unpackhi_epi64(p[8 * i], p[8 * i + 4]); + q[8 * i + 2] = _mm512_unpacklo_epi64(p[8 * i + 1], p[8 * i + 5]); + q[8 * i + 3] = _mm512_unpackhi_epi64(p[8 * i + 1], p[8 * i + 5]); + q[8 * i + 4] = _mm512_unpacklo_epi64(p[8 * i + 2], p[8 * i + 6]); + q[8 * i + 5] = _mm512_unpackhi_epi64(p[8 * i + 2], p[8 * i + 6]); + q[8 * i + 6] = _mm512_unpacklo_epi64(p[8 * i + 3], p[8 * i + 7]); + q[8 * i + 7] = _mm512_unpackhi_epi64(p[8 * i + 3], p[8 * i + 7]); + } + + __m512i f[32]; + +#define PACKET32H_TRANSPOSE_HELPER(X, Y) \ + do { \ + f[Y * 8] = _mm512_inserti32x4(f[Y * 8], _mm512_extracti32x4_epi32(q[X * 8], Y), X); \ + f[Y * 8 + 1] = _mm512_inserti32x4(f[Y * 8 + 1], _mm512_extracti32x4_epi32(q[X * 8 + 1], Y), X); \ + f[Y * 8 + 2] = _mm512_inserti32x4(f[Y * 8 + 2], _mm512_extracti32x4_epi32(q[X * 8 + 2], Y), X); \ + f[Y * 8 + 3] = _mm512_inserti32x4(f[Y * 8 + 3], _mm512_extracti32x4_epi32(q[X * 8 + 3], Y), X); \ + f[Y * 8 + 4] = _mm512_inserti32x4(f[Y * 8 + 4], _mm512_extracti32x4_epi32(q[X * 8 + 4], Y), X); \ + f[Y * 8 + 5] = _mm512_inserti32x4(f[Y * 8 + 5], _mm512_extracti32x4_epi32(q[X * 8 + 5], Y), X); \ + f[Y * 8 + 6] = _mm512_inserti32x4(f[Y * 8 + 6], _mm512_extracti32x4_epi32(q[X * 8 + 6], Y), X); \ + f[Y * 8 + 7] = _mm512_inserti32x4(f[Y * 8 + 7], _mm512_extracti32x4_epi32(q[X * 8 + 7], Y), X); \ + } while (false); + + PACKET32H_TRANSPOSE_HELPER(0, 0); + PACKET32H_TRANSPOSE_HELPER(1, 1); + PACKET32H_TRANSPOSE_HELPER(2, 2); + PACKET32H_TRANSPOSE_HELPER(3, 3); + + PACKET32H_TRANSPOSE_HELPER(1, 0); + PACKET32H_TRANSPOSE_HELPER(2, 0); + PACKET32H_TRANSPOSE_HELPER(3, 0); + PACKET32H_TRANSPOSE_HELPER(2, 1); + PACKET32H_TRANSPOSE_HELPER(3, 1); + PACKET32H_TRANSPOSE_HELPER(3, 2); + + PACKET32H_TRANSPOSE_HELPER(0, 1); + PACKET32H_TRANSPOSE_HELPER(0, 2); + PACKET32H_TRANSPOSE_HELPER(0, 3); + PACKET32H_TRANSPOSE_HELPER(1, 2); + PACKET32H_TRANSPOSE_HELPER(1, 3); + PACKET32H_TRANSPOSE_HELPER(2, 3); + +#undef PACKET32H_TRANSPOSE_HELPER + + EIGEN_UNROLL_LOOP + for (int i = 0; i < 32; i++) { + a.packet[i] = _mm512_castsi512_ph(f[i]); + } +} + +EIGEN_DEVICE_FUNC inline void ptranspose(PacketBlock<Packet32h, 4>& a) { + __m512i p0, p1, p2, p3, t0, t1, t2, t3, a0, a1, a2, a3; + t0 = _mm512_unpacklo_epi16(_mm512_castph_si512(a.packet[0]), _mm512_castph_si512(a.packet[1])); + t1 = _mm512_unpackhi_epi16(_mm512_castph_si512(a.packet[0]), _mm512_castph_si512(a.packet[1])); + t2 = _mm512_unpacklo_epi16(_mm512_castph_si512(a.packet[2]), _mm512_castph_si512(a.packet[3])); + t3 = _mm512_unpackhi_epi16(_mm512_castph_si512(a.packet[2]), _mm512_castph_si512(a.packet[3])); + + p0 = _mm512_unpacklo_epi32(t0, t2); + p1 = _mm512_unpackhi_epi32(t0, t2); + p2 = _mm512_unpacklo_epi32(t1, t3); + p3 = _mm512_unpackhi_epi32(t1, t3); + + a0 = p0; + a1 = p1; + a2 = p2; + a3 = p3; + + a0 = _mm512_inserti32x4(a0, _mm512_extracti32x4_epi32(p1, 0), 1); + a1 = _mm512_inserti32x4(a1, _mm512_extracti32x4_epi32(p0, 1), 0); + + a0 = _mm512_inserti32x4(a0, _mm512_extracti32x4_epi32(p2, 0), 2); + a2 = _mm512_inserti32x4(a2, _mm512_extracti32x4_epi32(p0, 2), 0); + + a0 = _mm512_inserti32x4(a0, _mm512_extracti32x4_epi32(p3, 0), 3); + a3 = _mm512_inserti32x4(a3, _mm512_extracti32x4_epi32(p0, 3), 0); + + a1 = _mm512_inserti32x4(a1, _mm512_extracti32x4_epi32(p2, 1), 2); + a2 = _mm512_inserti32x4(a2, _mm512_extracti32x4_epi32(p1, 2), 1); + + a2 = _mm512_inserti32x4(a2, _mm512_extracti32x4_epi32(p3, 2), 3); + a3 = _mm512_inserti32x4(a3, _mm512_extracti32x4_epi32(p2, 3), 2); + + a1 = _mm512_inserti32x4(a1, _mm512_extracti32x4_epi32(p3, 1), 3); + a3 = _mm512_inserti32x4(a3, _mm512_extracti32x4_epi32(p1, 3), 1); + + a.packet[0] = _mm512_castsi512_ph(a0); + a.packet[1] = _mm512_castsi512_ph(a1); + a.packet[2] = _mm512_castsi512_ph(a2); + a.packet[3] = _mm512_castsi512_ph(a3); +} + +// preverse + +template <> +EIGEN_STRONG_INLINE Packet32h preverse(const Packet32h& a) { + return _mm512_permutexvar_ph(_mm512_set_epi16(0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, + 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31), + a); +} + +// pscatter + +template <> +EIGEN_STRONG_INLINE void pscatter<half, Packet32h>(half* to, const Packet32h& from, Index stride) { + EIGEN_ALIGN64 half aux[32]; + pstore(aux, from); + + EIGEN_UNROLL_LOOP + for (int i = 0; i < 32; i++) { + to[stride * i] = aux[i]; + } +} + +// pgather + +template <> +EIGEN_STRONG_INLINE Packet32h pgather<Eigen::half, Packet32h>(const Eigen::half* from, Index stride) { + return _mm512_castsi512_ph(_mm512_set_epi16( + from[31 * stride].x, from[30 * stride].x, from[29 * stride].x, from[28 * stride].x, from[27 * stride].x, + from[26 * stride].x, from[25 * stride].x, from[24 * stride].x, from[23 * stride].x, from[22 * stride].x, + from[21 * stride].x, from[20 * stride].x, from[19 * stride].x, from[18 * stride].x, from[17 * stride].x, + from[16 * stride].x, from[15 * stride].x, from[14 * stride].x, from[13 * stride].x, from[12 * stride].x, + from[11 * stride].x, from[10 * stride].x, from[9 * stride].x, from[8 * stride].x, from[7 * stride].x, + from[6 * stride].x, from[5 * stride].x, from[4 * stride].x, from[3 * stride].x, from[2 * stride].x, + from[1 * stride].x, from[0 * stride].x)); +} + +template <> +EIGEN_STRONG_INLINE Packet16h pcos<Packet16h>(const Packet16h&); +template <> +EIGEN_STRONG_INLINE Packet16h psin<Packet16h>(const Packet16h&); +template <> +EIGEN_STRONG_INLINE Packet16h plog<Packet16h>(const Packet16h&); +template <> +EIGEN_STRONG_INLINE Packet16h plog2<Packet16h>(const Packet16h&); +template <> +EIGEN_STRONG_INLINE Packet16h plog1p<Packet16h>(const Packet16h&); +template <> +EIGEN_STRONG_INLINE Packet16h pexp<Packet16h>(const Packet16h&); +template <> +EIGEN_STRONG_INLINE Packet16h pexpm1<Packet16h>(const Packet16h&); +template <> +EIGEN_STRONG_INLINE Packet16h ptanh<Packet16h>(const Packet16h&); +template <> +EIGEN_STRONG_INLINE Packet16h pfrexp<Packet16h>(const Packet16h&, Packet16h&); +template <> +EIGEN_STRONG_INLINE Packet16h pldexp<Packet16h>(const Packet16h&, const Packet16h&); + +EIGEN_STRONG_INLINE Packet32h combine2Packet16h(const Packet16h& a, const Packet16h& b) { + __m512d result = _mm512_undefined_pd(); + result = _mm512_insertf64x4(result, _mm256_castsi256_pd(a), 0); + result = _mm512_insertf64x4(result, _mm256_castsi256_pd(b), 1); + return _mm512_castpd_ph(result); +} + +EIGEN_STRONG_INLINE void extract2Packet16h(const Packet32h& x, Packet16h& a, Packet16h& b) { + a = _mm256_castpd_si256(_mm512_extractf64x4_pd(_mm512_castph_pd(x), 0)); + b = _mm256_castpd_si256(_mm512_extractf64x4_pd(_mm512_castph_pd(x), 1)); +} + +// psin +template <> +EIGEN_STRONG_INLINE Packet32h psin<Packet32h>(const Packet32h& a) { + Packet16h low; + Packet16h high; + extract2Packet16h(a, low, high); + + Packet16h lowOut = psin(low); + Packet16h highOut = psin(high); + + return combine2Packet16h(lowOut, highOut); +} + +// pcos +template <> +EIGEN_STRONG_INLINE Packet32h pcos<Packet32h>(const Packet32h& a) { + Packet16h low; + Packet16h high; + extract2Packet16h(a, low, high); + + Packet16h lowOut = pcos(low); + Packet16h highOut = pcos(high); + + return combine2Packet16h(lowOut, highOut); +} + +// plog +template <> +EIGEN_STRONG_INLINE Packet32h plog<Packet32h>(const Packet32h& a) { + Packet16h low; + Packet16h high; + extract2Packet16h(a, low, high); + + Packet16h lowOut = plog(low); + Packet16h highOut = plog(high); + + return combine2Packet16h(lowOut, highOut); +} + +// plog2 +template <> +EIGEN_STRONG_INLINE Packet32h plog2<Packet32h>(const Packet32h& a) { + Packet16h low; + Packet16h high; + extract2Packet16h(a, low, high); + + Packet16h lowOut = plog2(low); + Packet16h highOut = plog2(high); + + return combine2Packet16h(lowOut, highOut); +} + +// plog1p +template <> +EIGEN_STRONG_INLINE Packet32h plog1p<Packet32h>(const Packet32h& a) { + Packet16h low; + Packet16h high; + extract2Packet16h(a, low, high); + + Packet16h lowOut = plog1p(low); + Packet16h highOut = plog1p(high); + + return combine2Packet16h(lowOut, highOut); +} + +// pexp +template <> +EIGEN_STRONG_INLINE Packet32h pexp<Packet32h>(const Packet32h& a) { + Packet16h low; + Packet16h high; + extract2Packet16h(a, low, high); + + Packet16h lowOut = pexp(low); + Packet16h highOut = pexp(high); + + return combine2Packet16h(lowOut, highOut); +} + +// pexpm1 +template <> +EIGEN_STRONG_INLINE Packet32h pexpm1<Packet32h>(const Packet32h& a) { + Packet16h low; + Packet16h high; + extract2Packet16h(a, low, high); + + Packet16h lowOut = pexpm1(low); + Packet16h highOut = pexpm1(high); + + return combine2Packet16h(lowOut, highOut); +} + +// ptanh +template <> +EIGEN_STRONG_INLINE Packet32h ptanh<Packet32h>(const Packet32h& a) { + Packet16h low; + Packet16h high; + extract2Packet16h(a, low, high); + + Packet16h lowOut = ptanh(low); + Packet16h highOut = ptanh(high); + + return combine2Packet16h(lowOut, highOut); +} + +// pfrexp +template <> +EIGEN_STRONG_INLINE Packet32h pfrexp<Packet32h>(const Packet32h& a, Packet32h& exponent) { + Packet16h low; + Packet16h high; + extract2Packet16h(a, low, high); + + Packet16h exp1 = _mm256_undefined_si256(); + Packet16h exp2 = _mm256_undefined_si256(); + + Packet16h lowOut = pfrexp(low, exp1); + Packet16h highOut = pfrexp(high, exp2); + + exponent = combine2Packet16h(exp1, exp2); + + return combine2Packet16h(lowOut, highOut); +} + +// pldexp +template <> +EIGEN_STRONG_INLINE Packet32h pldexp<Packet32h>(const Packet32h& a, const Packet32h& exponent) { + Packet16h low; + Packet16h high; + extract2Packet16h(a, low, high); + + Packet16h exp1; + Packet16h exp2; + extract2Packet16h(exponent, exp1, exp2); + + Packet16h lowOut = pldexp(low, exp1); + Packet16h highOut = pldexp(high, exp2); + + return combine2Packet16h(lowOut, highOut); +} + +} // end namespace internal +} // end namespace Eigen + +#endif // EIGEN_PACKET_MATH_FP16_AVX512_H
diff --git a/Eigen/src/Core/arch/AltiVec/MatrixProduct.h b/Eigen/src/Core/arch/AltiVec/MatrixProduct.h index b232a8f..e9a9307 100644 --- a/Eigen/src/Core/arch/AltiVec/MatrixProduct.h +++ b/Eigen/src/Core/arch/AltiVec/MatrixProduct.h
@@ -49,11 +49,6 @@ #include "MatrixProductMMA.h" #endif -/************************************************************************************************** - * TODO * - * - Check StorageOrder on dhs_pack (the innermost second loop seems unvectorized when it could). * - * - Check the possibility of transposing as GETREAL and GETIMAG when needed. * - **************************************************************************************************/ // IWYU pragma: private #include "../../InternalHeaderCheck.h" @@ -120,6 +115,16 @@ 20, 21, 22, 23, 28, 29, 30, 31}; +const static Packet16uc p16uc_GETREAL32b = { 0, 1, 2, 3, + 16, 17, 18, 19, + 8, 9, 10, 11, + 24, 25, 26, 27}; + +const static Packet16uc p16uc_GETIMAG32b = { 4, 5, 6, 7, + 20, 21, 22, 23, + 12, 13, 14, 15, + 28, 29, 30, 31}; + /********************************************* * Single precision real and complex packing * * *******************************************/ @@ -440,6 +445,78 @@ // General template for lhs & rhs complex packing. template<typename Scalar, typename DataMapper, typename Packet, typename PacketC, int StorageOrder, bool Conjugate, bool PanelMode, bool UseLhs> struct dhs_cpack { + template<bool transpose> + EIGEN_ALWAYS_INLINE void dhs_cblock(PacketBlock<PacketC,8>& cblock, PacketBlock<Packet,4>& block, Packet16uc permute) + { + if (transpose) { + block.packet[0] = vec_perm(cblock.packet[0].v, cblock.packet[1].v, permute); + block.packet[1] = vec_perm(cblock.packet[2].v, cblock.packet[3].v, permute); + block.packet[2] = vec_perm(cblock.packet[4].v, cblock.packet[5].v, permute); + block.packet[3] = vec_perm(cblock.packet[6].v, cblock.packet[7].v, permute); + + Packet4f t0, t1, t2, t3; +#ifdef EIGEN_VECTORIZE_VSX + t0 = reinterpret_cast<Packet>(vec_mergeh(reinterpret_cast<Packet2ul>(block.packet[0]), reinterpret_cast<Packet2ul>(block.packet[1]))); + t1 = reinterpret_cast<Packet>(vec_mergel(reinterpret_cast<Packet2ul>(block.packet[0]), reinterpret_cast<Packet2ul>(block.packet[1]))); + t2 = reinterpret_cast<Packet>(vec_mergeh(reinterpret_cast<Packet2ul>(block.packet[2]), reinterpret_cast<Packet2ul>(block.packet[3]))); + t3 = reinterpret_cast<Packet>(vec_mergel(reinterpret_cast<Packet2ul>(block.packet[2]), reinterpret_cast<Packet2ul>(block.packet[3]))); +#else + t0 = reinterpret_cast<Packet>(vec_perm(block.packet[0], block.packet[1], p16uc_TRANSPOSE64_HI)); + t1 = reinterpret_cast<Packet>(vec_perm(block.packet[0], block.packet[1], p16uc_TRANSPOSE64_LO)); + t2 = reinterpret_cast<Packet>(vec_perm(block.packet[2], block.packet[3], p16uc_TRANSPOSE64_HI)); + t3 = reinterpret_cast<Packet>(vec_perm(block.packet[2], block.packet[3], p16uc_TRANSPOSE64_LO)); +#endif + + block.packet[0] = t0; + block.packet[1] = t1; + block.packet[2] = t2; + block.packet[3] = t3; + } else { + block.packet[0] = vec_perm(cblock.packet[0].v, cblock.packet[4].v, permute); + block.packet[1] = vec_perm(cblock.packet[1].v, cblock.packet[5].v, permute); + block.packet[2] = vec_perm(cblock.packet[2].v, cblock.packet[6].v, permute); + block.packet[3] = vec_perm(cblock.packet[3].v, cblock.packet[7].v, permute); + } + } + + EIGEN_ALWAYS_INLINE void dhs_ccopy(Scalar* blockAt, const DataMapper& lhs2, Index& i, Index& rir, Index& rii, Index depth, const Index vectorSize) + { + PacketBlock<Packet,4> blockr, blocki; + PacketBlock<PacketC,8> cblock; + + for(; i + vectorSize <= depth; i+=vectorSize) + { + if (UseLhs) { + bload<DataMapper, PacketC, 2, StorageOrder, true, 4>(cblock, lhs2, 0, i); + } else { + bload<DataMapper, PacketC, 2, StorageOrder, true, 4>(cblock, lhs2, i, 0); + } + + if(((StorageOrder == RowMajor) && UseLhs) || (((StorageOrder == ColMajor) && !UseLhs))) + { + dhs_cblock<true>(cblock, blockr, p16uc_GETREAL32b); + dhs_cblock<true>(cblock, blocki, p16uc_GETIMAG32b); + } else { + dhs_cblock<false>(cblock, blockr, p16uc_GETREAL32); + dhs_cblock<false>(cblock, blocki, p16uc_GETIMAG32); + } + + if(Conjugate) + { + blocki.packet[0] = -blocki.packet[0]; + blocki.packet[1] = -blocki.packet[1]; + blocki.packet[2] = -blocki.packet[2]; + blocki.packet[3] = -blocki.packet[3]; + } + + storeBlock<Scalar, Packet, 4>(blockAt + rir, blockr); + storeBlock<Scalar, Packet, 4>(blockAt + rii, blocki); + + rir += 4*vectorSize; + rii += 4*vectorSize; + } + } + EIGEN_STRONG_INLINE void operator()(std::complex<Scalar>* blockA, const DataMapper& lhs, Index depth, Index rows, Index stride, Index offset) { const Index vectorSize = quad_traits<Scalar>::vectorsize; @@ -455,47 +532,8 @@ rii = rir + vectorDelta; - for(; i + vectorSize <= depth; i+=vectorSize) - { - PacketBlock<Packet,4> blockr, blocki; - PacketBlock<PacketC,8> cblock; + dhs_ccopy(blockAt, lhs2, i, rir, rii, depth, vectorSize); - if (UseLhs) { - bload<DataMapper, PacketC, 2, StorageOrder, true, 4>(cblock, lhs2, 0, i); - } else { - bload<DataMapper, PacketC, 2, StorageOrder, true, 4>(cblock, lhs2, i, 0); - } - - blockr.packet[0] = vec_perm(cblock.packet[0].v, cblock.packet[4].v, p16uc_GETREAL32); - blockr.packet[1] = vec_perm(cblock.packet[1].v, cblock.packet[5].v, p16uc_GETREAL32); - blockr.packet[2] = vec_perm(cblock.packet[2].v, cblock.packet[6].v, p16uc_GETREAL32); - blockr.packet[3] = vec_perm(cblock.packet[3].v, cblock.packet[7].v, p16uc_GETREAL32); - - blocki.packet[0] = vec_perm(cblock.packet[0].v, cblock.packet[4].v, p16uc_GETIMAG32); - blocki.packet[1] = vec_perm(cblock.packet[1].v, cblock.packet[5].v, p16uc_GETIMAG32); - blocki.packet[2] = vec_perm(cblock.packet[2].v, cblock.packet[6].v, p16uc_GETIMAG32); - blocki.packet[3] = vec_perm(cblock.packet[3].v, cblock.packet[7].v, p16uc_GETIMAG32); - - if(Conjugate) - { - blocki.packet[0] = -blocki.packet[0]; - blocki.packet[1] = -blocki.packet[1]; - blocki.packet[2] = -blocki.packet[2]; - blocki.packet[3] = -blocki.packet[3]; - } - - if(((StorageOrder == RowMajor) && UseLhs) || (((StorageOrder == ColMajor) && !UseLhs))) - { - ptranspose(blockr); - ptranspose(blocki); - } - - storeBlock<Scalar, Packet, 4>(blockAt + rir, blockr); - storeBlock<Scalar, Packet, 4>(blockAt + rii, blocki); - - rir += 4*vectorSize; - rii += 4*vectorSize; - } for(; i < depth; i++) { PacketBlock<Packet,1> blockr, blocki; @@ -592,6 +630,36 @@ // General template for lhs & rhs packing. template<typename Scalar, typename DataMapper, typename Packet, int StorageOrder, bool PanelMode, bool UseLhs> struct dhs_pack{ + template<Index n> + EIGEN_ALWAYS_INLINE void dhs_copy(Scalar* blockA, const DataMapper& lhs2, Index& i, Index& ri, Index depth, const Index vectorSize) + { + PacketBlock<Packet,4> block[n]; + + for(; i + n*vectorSize <= depth; i+=n*vectorSize) + { + for (Index k = 0; k < n; k++) { + if (UseLhs) { + bload<DataMapper, Packet, 4, StorageOrder, false, 4>(block[k], lhs2, 0, i + k*vectorSize); + } else { + bload<DataMapper, Packet, 4, StorageOrder, false, 4>(block[k], lhs2, i + k*vectorSize, 0); + } + } + + if(((StorageOrder == RowMajor) && UseLhs) || ((StorageOrder == ColMajor) && !UseLhs)) + { + for (Index k = 0; k < n; k++) { + ptranspose(block[k]); + } + } + + for (Index k = 0; k < n; k++) { + storeBlock<Scalar, Packet, 4>(blockA + ri + k*4*vectorSize, block[k]); + } + + ri += n*4*vectorSize; + } + } + EIGEN_STRONG_INLINE void operator()(Scalar* blockA, const DataMapper& lhs, Index depth, Index rows, Index stride, Index offset) { const Index vectorSize = quad_traits<Scalar>::vectorsize; @@ -604,24 +672,10 @@ if(PanelMode) ri += vectorSize*offset; - for(; i + vectorSize <= depth; i+=vectorSize) - { - PacketBlock<Packet,4> block; + dhs_copy<4>(blockA, lhs2, i, ri, depth, vectorSize); + dhs_copy<2>(blockA, lhs2, i, ri, depth, vectorSize); + dhs_copy<1>(blockA, lhs2, i, ri, depth, vectorSize); - if (UseLhs) { - bload<DataMapper, Packet, 4, StorageOrder, false, 4>(block, lhs2, 0, i); - } else { - bload<DataMapper, Packet, 4, StorageOrder, false, 4>(block, lhs2, i, 0); - } - if(((StorageOrder == RowMajor) && UseLhs) || ((StorageOrder == ColMajor) && !UseLhs)) - { - ptranspose(block); - } - - storeBlock<Scalar, Packet, 4>(blockA + ri, block); - - ri += 4*vectorSize; - } for(; i < depth; i++) { if(((StorageOrder == RowMajor) && UseLhs) || ((StorageOrder == ColMajor) && !UseLhs)) @@ -691,6 +745,39 @@ template<typename DataMapper, int StorageOrder, bool PanelMode> struct dhs_pack<double, DataMapper, Packet2d, StorageOrder, PanelMode, true> { + template<Index n> + EIGEN_ALWAYS_INLINE void dhs_copy(double* blockA, const DataMapper& lhs2, Index& i, Index& ri, Index depth, const Index vectorSize) + { + PacketBlock<Packet2d,2> block[n]; + + for(; i + n*vectorSize <= depth; i+=n*vectorSize) + { + for (Index k = 0; k < n; k++) { + if(StorageOrder == RowMajor) + { + block[k].packet[0] = lhs2.template loadPacket<Packet2d>(0, i + k*vectorSize); + block[k].packet[1] = lhs2.template loadPacket<Packet2d>(1, i + k*vectorSize); + } else { + block[k].packet[0] = lhs2.template loadPacket<Packet2d>(0, i + k*vectorSize + 0); + block[k].packet[1] = lhs2.template loadPacket<Packet2d>(0, i + k*vectorSize + 1); + } + } + + if(StorageOrder == RowMajor) + { + for (Index k = 0; k < n; k++) { + ptranspose(block[k]); + } + } + + for (Index k = 0; k < n; k++) { + storeBlock<double, Packet2d, 2>(blockA + ri + k*2*vectorSize, block[k]); + } + + ri += n*2*vectorSize; + } + } + EIGEN_STRONG_INLINE void operator()(double* blockA, const DataMapper& lhs, Index depth, Index rows, Index stride, Index offset) { const Index vectorSize = quad_traits<double>::vectorsize; @@ -703,24 +790,10 @@ if(PanelMode) ri += vectorSize*offset; - for(; i + vectorSize <= depth; i+=vectorSize) - { - PacketBlock<Packet2d,2> block; - if(StorageOrder == RowMajor) - { - block.packet[0] = lhs2.template loadPacket<Packet2d>(0, i); - block.packet[1] = lhs2.template loadPacket<Packet2d>(1, i); + dhs_copy<4>(blockA, lhs2, i, ri, depth, vectorSize); + dhs_copy<2>(blockA, lhs2, i, ri, depth, vectorSize); + dhs_copy<1>(blockA, lhs2, i, ri, depth, vectorSize); - ptranspose(block); - } else { - block.packet[0] = lhs2.template loadPacket<Packet2d>(0, i + 0); - block.packet[1] = lhs2.template loadPacket<Packet2d>(0, i + 1); - } - - storeBlock<double, Packet2d, 2>(blockA + ri, block); - - ri += 2*vectorSize; - } for(; i < depth; i++) { if(StorageOrder == RowMajor) @@ -759,6 +832,53 @@ template<typename DataMapper, int StorageOrder, bool PanelMode> struct dhs_pack<double, DataMapper, Packet2d, StorageOrder, PanelMode, false> { + template<Index n> + EIGEN_ALWAYS_INLINE void dhs_copy(double* blockB, const DataMapper& rhs2, Index& i, Index& ri, Index depth, const Index vectorSize) + { + PacketBlock<Packet2d,2> block1[n], block2[n]; + PacketBlock<Packet2d,4> block3[n]; + + for(; i + n*vectorSize <= depth; i+=n*vectorSize) + { + for (Index k = 0; k < n; k++) { + if(StorageOrder == ColMajor) + { + block1[k].packet[0] = rhs2.template loadPacket<Packet2d>(i + k*vectorSize, 0); + block1[k].packet[1] = rhs2.template loadPacket<Packet2d>(i + k*vectorSize, 1); + block2[k].packet[0] = rhs2.template loadPacket<Packet2d>(i + k*vectorSize, 2); + block2[k].packet[1] = rhs2.template loadPacket<Packet2d>(i + k*vectorSize, 3); + } else { + block3[k].packet[0] = rhs2.template loadPacket<Packet2d>(i + k*vectorSize + 0, 0); //[a1 a2] + block3[k].packet[1] = rhs2.template loadPacket<Packet2d>(i + k*vectorSize + 0, 2); //[a3 a4] + block3[k].packet[2] = rhs2.template loadPacket<Packet2d>(i + k*vectorSize + 1, 0); //[b1 b2] + block3[k].packet[3] = rhs2.template loadPacket<Packet2d>(i + k*vectorSize + 1, 2); //[b3 b4] + } + } + + if(StorageOrder == ColMajor) + { + for (Index k = 0; k < n; k++) { + ptranspose(block1[k]); + ptranspose(block2[k]); + } + } + + for (Index k = 0; k < n; k++) { + if(StorageOrder == ColMajor) + { + pstore<double>(blockB + ri + k*4*vectorSize , block1[k].packet[0]); + pstore<double>(blockB + ri + k*4*vectorSize + 2, block2[k].packet[0]); + pstore<double>(blockB + ri + k*4*vectorSize + 4, block1[k].packet[1]); + pstore<double>(blockB + ri + k*4*vectorSize + 6, block2[k].packet[1]); + } else { + storeBlock<double, Packet2d, 4>(blockB + ri + k*4*vectorSize, block3[k]); + } + } + + ri += n*4*vectorSize; + } + } + EIGEN_STRONG_INLINE void operator()(double* blockB, const DataMapper& rhs, Index depth, Index cols, Index stride, Index offset) { const Index vectorSize = quad_traits<double>::vectorsize; @@ -771,35 +891,10 @@ if(PanelMode) ri += offset*(2*vectorSize); - for(; i + vectorSize <= depth; i+=vectorSize) - { - PacketBlock<Packet2d,4> block; - if(StorageOrder == ColMajor) - { - PacketBlock<Packet2d,2> block1, block2; - block1.packet[0] = rhs2.template loadPacket<Packet2d>(i, 0); - block1.packet[1] = rhs2.template loadPacket<Packet2d>(i, 1); - block2.packet[0] = rhs2.template loadPacket<Packet2d>(i, 2); - block2.packet[1] = rhs2.template loadPacket<Packet2d>(i, 3); + dhs_copy<4>(blockB, rhs2, i, ri, depth, vectorSize); + dhs_copy<2>(blockB, rhs2, i, ri, depth, vectorSize); + dhs_copy<1>(blockB, rhs2, i, ri, depth, vectorSize); - ptranspose(block1); - ptranspose(block2); - - pstore<double>(blockB + ri , block1.packet[0]); - pstore<double>(blockB + ri + 2, block2.packet[0]); - pstore<double>(blockB + ri + 4, block1.packet[1]); - pstore<double>(blockB + ri + 6, block2.packet[1]); - } else { - block.packet[0] = rhs2.template loadPacket<Packet2d>(i + 0, 0); //[a1 a2] - block.packet[1] = rhs2.template loadPacket<Packet2d>(i + 0, 2); //[a3 a4] - block.packet[2] = rhs2.template loadPacket<Packet2d>(i + 1, 0); //[b1 b2] - block.packet[3] = rhs2.template loadPacket<Packet2d>(i + 1, 2); //[b3 b4] - - storeBlock<double, Packet2d, 4>(blockB + ri, block); - } - - ri += 4*vectorSize; - } for(; i < depth; i++) { if(StorageOrder == ColMajor) @@ -1296,6 +1391,54 @@ template<typename DataMapper, typename Packet, typename PacketC, int StorageOrder, bool Conjugate, bool PanelMode> struct dhs_cpack<double, DataMapper, Packet, PacketC, StorageOrder, Conjugate, PanelMode, true> { + EIGEN_ALWAYS_INLINE void dhs_ccopy(double* blockAt, const DataMapper& lhs2, Index& i, Index& rir, Index& rii, Index depth, const Index vectorSize) + { + PacketBlock<Packet,2> blockr, blocki; + PacketBlock<PacketC,4> cblock; + + for(; i + vectorSize <= depth; i+=vectorSize) + { + if(StorageOrder == ColMajor) + { + cblock.packet[0] = lhs2.template loadPacket<PacketC>(0, i + 0); //[a1 a1i] + cblock.packet[1] = lhs2.template loadPacket<PacketC>(0, i + 1); //[b1 b1i] + + cblock.packet[2] = lhs2.template loadPacket<PacketC>(1, i + 0); //[a2 a2i] + cblock.packet[3] = lhs2.template loadPacket<PacketC>(1, i + 1); //[b2 b2i] + + blockr.packet[0] = vec_mergeh(cblock.packet[0].v, cblock.packet[2].v); //[a1 a2] + blockr.packet[1] = vec_mergeh(cblock.packet[1].v, cblock.packet[3].v); //[b1 b2] + + blocki.packet[0] = vec_mergel(cblock.packet[0].v, cblock.packet[2].v); + blocki.packet[1] = vec_mergel(cblock.packet[1].v, cblock.packet[3].v); + } else { + cblock.packet[0] = lhs2.template loadPacket<PacketC>(0, i); //[a1 a1i] + cblock.packet[1] = lhs2.template loadPacket<PacketC>(1, i); //[a2 a2i] + + cblock.packet[2] = lhs2.template loadPacket<PacketC>(0, i + 1); //[b1 b1i] + cblock.packet[3] = lhs2.template loadPacket<PacketC>(1, i + 1); //[b2 b2i + + blockr.packet[0] = vec_mergeh(cblock.packet[0].v, cblock.packet[1].v); //[a1 a2] + blockr.packet[1] = vec_mergeh(cblock.packet[2].v, cblock.packet[3].v); //[b1 b2] + + blocki.packet[0] = vec_mergel(cblock.packet[0].v, cblock.packet[1].v); + blocki.packet[1] = vec_mergel(cblock.packet[2].v, cblock.packet[3].v); + } + + if(Conjugate) + { + blocki.packet[0] = -blocki.packet[0]; + blocki.packet[1] = -blocki.packet[1]; + } + + storeBlock<double, Packet, 2>(blockAt + rir, blockr); + storeBlock<double, Packet, 2>(blockAt + rii, blocki); + + rir += 2*vectorSize; + rii += 2*vectorSize; + } + } + EIGEN_STRONG_INLINE void operator()(std::complex<double>* blockA, const DataMapper& lhs, Index depth, Index rows, Index stride, Index offset) { const Index vectorSize = quad_traits<double>::vectorsize; @@ -1311,50 +1454,8 @@ rii = rir + vectorDelta; - for(; i + vectorSize <= depth; i+=vectorSize) - { - PacketBlock<Packet,2> blockr, blocki; - PacketBlock<PacketC,4> cblock; + dhs_ccopy(blockAt, lhs2, i, rir, rii, depth, vectorSize); - if(StorageOrder == ColMajor) - { - cblock.packet[0] = lhs2.template loadPacket<PacketC>(0, i + 0); //[a1 a1i] - cblock.packet[1] = lhs2.template loadPacket<PacketC>(0, i + 1); //[b1 b1i] - - cblock.packet[2] = lhs2.template loadPacket<PacketC>(1, i + 0); //[a2 a2i] - cblock.packet[3] = lhs2.template loadPacket<PacketC>(1, i + 1); //[b2 b2i] - - blockr.packet[0] = vec_mergeh(cblock.packet[0].v, cblock.packet[2].v); //[a1 a2] - blockr.packet[1] = vec_mergeh(cblock.packet[1].v, cblock.packet[3].v); //[b1 b2] - - blocki.packet[0] = vec_mergel(cblock.packet[0].v, cblock.packet[2].v); - blocki.packet[1] = vec_mergel(cblock.packet[1].v, cblock.packet[3].v); - } else { - cblock.packet[0] = lhs2.template loadPacket<PacketC>(0, i); //[a1 a1i] - cblock.packet[1] = lhs2.template loadPacket<PacketC>(1, i); //[a2 a2i] - - cblock.packet[2] = lhs2.template loadPacket<PacketC>(0, i + 1); //[b1 b1i] - cblock.packet[3] = lhs2.template loadPacket<PacketC>(1, i + 1); //[b2 b2i - - blockr.packet[0] = vec_mergeh(cblock.packet[0].v, cblock.packet[1].v); //[a1 a2] - blockr.packet[1] = vec_mergeh(cblock.packet[2].v, cblock.packet[3].v); //[b1 b2] - - blocki.packet[0] = vec_mergel(cblock.packet[0].v, cblock.packet[1].v); - blocki.packet[1] = vec_mergel(cblock.packet[2].v, cblock.packet[3].v); - } - - if(Conjugate) - { - blocki.packet[0] = -blocki.packet[0]; - blocki.packet[1] = -blocki.packet[1]; - } - - storeBlock<double, Packet, 2>(blockAt + rir, blockr); - storeBlock<double, Packet, 2>(blockAt + rii, blocki); - - rir += 2*vectorSize; - rii += 2*vectorSize; - } for(; i < depth; i++) { PacketBlock<Packet,1> blockr, blocki; @@ -1410,6 +1511,35 @@ template<typename DataMapper, typename Packet, typename PacketC, int StorageOrder, bool Conjugate, bool PanelMode> struct dhs_cpack<double, DataMapper, Packet, PacketC, StorageOrder, Conjugate, PanelMode, false> { + EIGEN_ALWAYS_INLINE void dhs_ccopy(double* blockBt, const DataMapper& rhs2, Index& i, Index& rir, Index& rii, Index depth, const Index vectorSize) + { + for(; i < depth; i++) + { + PacketBlock<PacketC,4> cblock; + PacketBlock<Packet,2> blockr, blocki; + + bload<DataMapper, PacketC, 2, ColMajor, false, 4>(cblock, rhs2, i, 0); + + blockr.packet[0] = vec_mergeh(cblock.packet[0].v, cblock.packet[1].v); + blockr.packet[1] = vec_mergeh(cblock.packet[2].v, cblock.packet[3].v); + + blocki.packet[0] = vec_mergel(cblock.packet[0].v, cblock.packet[1].v); + blocki.packet[1] = vec_mergel(cblock.packet[2].v, cblock.packet[3].v); + + if(Conjugate) + { + blocki.packet[0] = -blocki.packet[0]; + blocki.packet[1] = -blocki.packet[1]; + } + + storeBlock<double, Packet, 2>(blockBt + rir, blockr); + storeBlock<double, Packet, 2>(blockBt + rii, blocki); + + rir += 2*vectorSize; + rii += 2*vectorSize; + } + } + EIGEN_STRONG_INLINE void operator()(std::complex<double>* blockB, const DataMapper& rhs, Index depth, Index cols, Index stride, Index offset) { const Index vectorSize = quad_traits<double>::vectorsize; @@ -1425,31 +1555,7 @@ rii = rir + vectorDelta; - for(; i < depth; i++) - { - PacketBlock<PacketC,4> cblock; - PacketBlock<Packet,2> blockr, blocki; - - bload<DataMapper, PacketC, 2, ColMajor, false, 4>(cblock, rhs2, i, 0); - - blockr.packet[0] = vec_mergeh(cblock.packet[0].v, cblock.packet[1].v); - blockr.packet[1] = vec_mergeh(cblock.packet[2].v, cblock.packet[3].v); - - blocki.packet[0] = vec_mergel(cblock.packet[0].v, cblock.packet[1].v); - blocki.packet[1] = vec_mergel(cblock.packet[2].v, cblock.packet[3].v); - - if(Conjugate) - { - blocki.packet[0] = -blocki.packet[0]; - blocki.packet[1] = -blocki.packet[1]; - } - - storeBlock<double, Packet, 2>(blockBt + rir, blockr); - storeBlock<double, Packet, 2>(blockBt + rii, blocki); - - rir += 2*vectorSize; - rii += 2*vectorSize; - } + dhs_ccopy(blockBt, rhs2, i, rir, rii, depth, vectorSize); rir += ((PanelMode) ? (2*vectorSize*(2*stride - depth)) : vectorDelta); }
diff --git a/Eigen/src/Core/arch/AltiVec/MatrixProductMMA.h b/Eigen/src/Core/arch/AltiVec/MatrixProductMMA.h index e83cba5..72e8c31 100644 --- a/Eigen/src/Core/arch/AltiVec/MatrixProductMMA.h +++ b/Eigen/src/Core/arch/AltiVec/MatrixProductMMA.h
@@ -42,19 +42,13 @@ __builtin_mma_xxsetaccz(acc); } -#ifdef USE_PARTIAL_PACKETS template<typename DataMapper, typename Packet, bool full> EIGEN_ALWAYS_INLINE void storeAccumulator(Index i, const DataMapper& data, const Packet& alpha, const Index elements, __vector_quad* acc) -#else -template<typename DataMapper, typename Packet, const Index accCols, const Index accCols2> -EIGEN_ALWAYS_INLINE void storeAccumulator(Index i, const DataMapper& data, const Packet& alpha, const Packet& pMask, __vector_quad* acc) -#endif { PacketBlock<Packet, 4> result; __builtin_mma_disassemble_acc(&result.packet, acc); PacketBlock<Packet, 4> tRes; -#ifdef USE_PARTIAL_PACKETS if (full) { EIGEN_UNUSED_VARIABLE(elements); bload<DataMapper, Packet, 0, ColMajor, false, 4>(tRes, data, i, 0); @@ -65,11 +59,6 @@ bscale<Packet, 4>(tRes, result, alpha); bstore_partial<DataMapper, Packet, 4>(tRes, data, i, elements); } -#else - bload<DataMapper, Packet, 0, ColMajor, false, 4>(tRes, data, i, 0); - bscale<Packet, 4, (accCols != accCols2)>(tRes, result, alpha, pMask); - bstore<DataMapper, Packet, 4>(tRes, data, i); -#endif } template<typename DataMapper, typename Packet, typename Packetc, const Index accCols, const Index accCols2> @@ -166,78 +155,118 @@ ploadRhsMMA(lhs, lhsV); } -#if (EIGEN_COMP_LLVM || (__GNUC__ >= 11)) +#define GEMM_MULTIPLE_COLS + +// Disable in GCC until unnecessary register moves are fixed +//#if (EIGEN_COMP_LLVM || (__GNUC__ >= 11)) +#if EIGEN_COMP_LLVM #define VECTOR_PAIR_LOADS_LHS #endif // PEEL_MMA loop factor. +#ifdef GEMM_MULTIPLE_COLS +#define PEEL_MMA 8 +#else +// Register spillage with GCC12+ +#if EIGEN_COMP_LLVM || (__GNUC__ < 12) || defined(VECTOR_PAIR_LOADS_LHS) #define PEEL_MMA 7 +#else +#define PEEL_MMA 6 +#endif +#endif #define MICRO_MMA_UNROLL(func) \ func(0) func(1) func(2) func(3) func(4) func(5) func(6) func(7) #define MICRO_MMA_WORK(func, type, peel) \ - func(0,type,peel) func(1,type,peel) func(2,type,peel) func(3,type,peel) \ - func(4,type,peel) func(5,type,peel) func(6,type,peel) func(7,type,peel) + if (accItr == 1) { \ + func(0,type,peel,0,0) func(1,type,peel,1,0) func(2,type,peel,2,0) func(3,type,peel,3,0) \ + func(4,type,peel,4,0) func(5,type,peel,5,0) func(6,type,peel,6,0) func(7,type,peel,7,0) \ + } else if (accItr == 2) { \ + func(0,type,peel,0,0) func(1,type,peel,0,1) func(2,type,peel,1,0) func(3,type,peel,1,1) \ + func(4,type,peel,2,0) func(5,type,peel,2,1) func(6,type,peel,3,0) func(7,type,peel,3,1) \ + } else { \ + func(0,type,peel,0,0) func(1,type,peel,0,1) func(2,type,peel,0,2) func(3,type,peel,0,3) \ + func(4,type,peel,1,0) func(5,type,peel,1,1) func(6,type,peel,1,2) func(7,type,peel,1,3) \ + } -#define MICRO_MMA_WORK_ONE(iter, type, peel) \ - if (unroll_factor > iter) { \ - pgerMMA<Packet, type, false>(&accZero##iter, rhsV[peel], lhsV##iter); \ +#define MICRO_MMA_WORK_ONE(iter, type, peel, left, right) \ + if (unroll_factor > left) { \ + pgerMMA<Packet, type, false>(&accZero##iter, rhsV##right[peel], lhsV##left); \ } #ifdef VECTOR_PAIR_LOADS_LHS -#define MICRO_MMA_WORK_TWO(iter, type, peel) \ - if (unroll_factor > iter) { \ - pgerMMA<Packet, type, false>(&accZero##iter, rhsV[peel], lhsV2##iter.packet[peel & 1]); \ +#define MICRO_MMA_WORK_TWO(iter, type, peel, left, right) \ + if (unroll_factor > left) { \ + pgerMMA<Packet, type, false>(&accZero##iter, rhsV##right[peel], lhsV2##left.packet[peel & 1]); \ } -#define MICRO_MMA_LOAD1_TWO(lhs_ptr, iter) \ - if (unroll_factor > iter) { \ - if (MICRO_NORMAL(iter)) { \ - ploadLhsMMA(reinterpret_cast<const double*>(lhs_ptr##iter), plhsV##iter); \ - __builtin_vsx_disassemble_pair(reinterpret_cast<void*>(&lhsV2##iter.packet), &plhsV##iter); \ - lhs_ptr##iter += accCols*2; \ +#define MICRO_MMA_LOAD1_TWO(lhs_ptr, left) \ + if (unroll_factor > left) { \ + if (MICRO_NORMAL(left)) { \ + ploadLhsMMA(reinterpret_cast<const double*>(lhs_ptr##left), plhsV##left); \ + __builtin_vsx_disassemble_pair(reinterpret_cast<void*>(&lhsV2##left.packet), &plhsV##left); \ + lhs_ptr##left += accCols*2; \ } else { \ - lhsV2##iter.packet[0] = ploadLhs<Packet>(lhs_ptr##iter); \ - lhsV2##iter.packet[1] = ploadLhs<Packet>(lhs_ptr##iter + accCols2); \ - lhs_ptr##iter += accCols2*2; \ - EIGEN_UNUSED_VARIABLE(plhsV##iter) \ + lhsV2##left.packet[0] = ploadLhs<Packet>(lhs_ptr##left); \ + lhsV2##left.packet[1] = ploadLhs<Packet>(lhs_ptr##left + accCols2); \ + lhs_ptr##left += accCols2*2; \ + EIGEN_UNUSED_VARIABLE(plhsV##left); \ } \ } else { \ - EIGEN_UNUSED_VARIABLE(lhsV2##iter); \ - EIGEN_UNUSED_VARIABLE(plhsV##iter) \ + EIGEN_UNUSED_VARIABLE(lhsV2##left); \ + EIGEN_UNUSED_VARIABLE(plhsV##left); \ } -#define MICRO_MMA_LOAD_TWO(iter) MICRO_MMA_LOAD1_TWO(lhs_ptr, iter) +#define MICRO_MMA_LOAD_TWO(left) MICRO_MMA_LOAD1_TWO(lhs_ptr, left) #endif +#define MICRO_MMA_UNROLL_ITER(func, val) \ + func(val,0) \ + if (accItr > 1) { \ + func(val,1) \ + if (accItr > 2) { \ + func(val,2) \ + func(val,3) \ + } \ + } + +#define MICRO_MMA_LOAD_ONE_RHS1(peel, right) \ + ploadRhsMMA(rhs_ptr##right + (accRows * peel), rhsV##right[peel]); + +#define MICRO_MMA_LOAD_ONE_RHS(peel) \ + MICRO_MMA_UNROLL_ITER(MICRO_MMA_LOAD_ONE_RHS1, peel) + #define MICRO_MMA_TYPE_PEEL(funcw, funcl, type, peel) \ if (PEEL_MMA > peel) { \ Packet lhsV0, lhsV1, lhsV2, lhsV3, lhsV4, lhsV5, lhsV6, lhsV7; \ - ploadRhsMMA(rhs_ptr + (accRows * peel), rhsV[peel]); \ + MICRO_MMA_LOAD_ONE_RHS(peel) \ MICRO_MMA_UNROLL(funcl) \ MICRO_MMA_WORK(funcw, type, peel) \ } #ifndef VECTOR_PAIR_LOADS_LHS #define MICRO_MMA_UNROLL_TYPE_PEEL(funcw, funcl, type) \ - type rhsV[8]; \ + type rhsV0[8], rhsV1[(accItr > 1) ? 8 : 1], rhsV2[(accItr > 2) ? 8 : 1], rhsV3[(accItr > 2) ? 8 : 1]; \ MICRO_MMA_TYPE_PEEL(funcw,funcl,type,0) MICRO_MMA_TYPE_PEEL(funcw,funcl,type,1) \ MICRO_MMA_TYPE_PEEL(funcw,funcl,type,2) MICRO_MMA_TYPE_PEEL(funcw,funcl,type,3) \ MICRO_MMA_TYPE_PEEL(funcw,funcl,type,4) MICRO_MMA_TYPE_PEEL(funcw,funcl,type,5) \ MICRO_MMA_TYPE_PEEL(funcw,funcl,type,6) MICRO_MMA_TYPE_PEEL(funcw,funcl,type,7) #else +#define MICRO_MMA_LOAD_TWO_RHS(peel1, right) \ + ploadRhsMMA(reinterpret_cast<const double*>(rhs_ptr##right + (accRows * peel1)), prhsV##peel1); \ + __builtin_vsx_disassemble_pair(reinterpret_cast<void*>(&rhsV##right[peel1]), &prhsV##peel1); + #define MICRO_MMA_TYPE_PEEL2(funcw1, funcl1, funcw2, funcl2, type, peel1, peel2) \ if (PEEL_MMA > peel2) { \ PacketBlock<Packet,2> lhsV20, lhsV21, lhsV22, lhsV23, lhsV24, lhsV25, lhsV26, lhsV27; \ __vector_pair plhsV0, plhsV1, plhsV2, plhsV3, plhsV4, plhsV5, plhsV6, plhsV7; \ if (sizeof(type) == 16) { \ - ploadRhsMMA(reinterpret_cast<const double*>(rhs_ptr + (accRows * peel1)), prhsV##peel1); \ - __builtin_vsx_disassemble_pair(reinterpret_cast<void*>(&rhsV[peel1]), &prhsV##peel1); \ + MICRO_MMA_UNROLL_ITER(MICRO_MMA_LOAD_TWO_RHS, peel1) \ } else { \ EIGEN_UNUSED_VARIABLE(prhsV##peel1); \ - ploadRhsMMA(rhs_ptr + (accRows * peel1), rhsV[peel1]); \ - ploadRhsMMA(rhs_ptr + (accRows * peel2), rhsV[peel2]); \ + MICRO_MMA_LOAD_ONE_RHS(peel1) \ + MICRO_MMA_LOAD_ONE_RHS(peel2) \ } \ MICRO_MMA_UNROLL(funcl2) \ MICRO_MMA_WORK(funcw2, type, peel1) \ @@ -248,7 +277,7 @@ } #define MICRO_MMA_UNROLL_TYPE_PEEL2(funcw1, funcl1, funcw2, funcl2, type) \ - type rhsV[8]; \ + type rhsV0[8], rhsV1[(accItr > 1) ? 8 : 1], rhsV2[(accItr > 2) ? 8 : 1], rhsV3[(accItr > 2) ? 8 : 1]; \ __vector_pair prhsV0, prhsV2, prhsV4, prhsV6; \ MICRO_MMA_TYPE_PEEL2(funcw1,funcl1,funcw2,funcl2,type,0,1) \ MICRO_MMA_TYPE_PEEL2(funcw1,funcl1,funcw2,funcl2,type,2,3) \ @@ -257,19 +286,25 @@ #endif #define MICRO_MMA_UNROLL_TYPE_ONE(funcw, funcl, type) \ - type rhsV[1]; \ + type rhsV0[1], rhsV1[1], rhsV2[1], rhsV3[1]; \ MICRO_MMA_TYPE_PEEL(funcw,funcl,type,0) +#define MICRO_MMA_UPDATE_RHS1(size, right) \ + rhs_ptr##right += (accRows * size); + +#define MICRO_MMA_UPDATE_RHS(size) \ + MICRO_MMA_UNROLL_ITER(MICRO_MMA_UPDATE_RHS1, size) + #define MICRO_MMA_UNROLL_TYPE(MICRO_MMA_TYPE, size) \ MICRO_MMA_TYPE(MICRO_MMA_WORK_ONE, MICRO_LOAD_ONE, RhsPacket) \ - rhs_ptr += (accRows * size); + MICRO_MMA_UPDATE_RHS(size) #ifndef VECTOR_PAIR_LOADS_LHS #define MICRO_MMA_ONE_PEEL MICRO_MMA_UNROLL_TYPE(MICRO_MMA_UNROLL_TYPE_PEEL, PEEL_MMA) #else #define MICRO_MMA_UNROLL_TYPE2(MICRO_MMA_TYPE, size) \ MICRO_MMA_TYPE(MICRO_MMA_WORK_ONE, MICRO_LOAD_ONE, MICRO_MMA_WORK_TWO, MICRO_MMA_LOAD_TWO, RhsPacket) \ - rhs_ptr += (accRows * size); + MICRO_MMA_UPDATE_RHS(size) #define MICRO_MMA_ONE_PEEL MICRO_MMA_UNROLL_TYPE2(MICRO_MMA_UNROLL_TYPE_PEEL2, PEEL_MMA) #endif @@ -277,7 +312,7 @@ #define MICRO_MMA_ONE MICRO_MMA_UNROLL_TYPE(MICRO_MMA_UNROLL_TYPE_ONE, 1) #define MICRO_MMA_DST_PTR_ONE(iter) \ - if (unroll_factor > iter) { \ + if (unroll_factor * accItr > iter) { \ bsetzeroMMA(&accZero##iter); \ } else { \ EIGEN_UNUSED_VARIABLE(accZero##iter); \ @@ -289,45 +324,69 @@ #define MICRO_MMA_PREFETCH MICRO_MMA_UNROLL(MICRO_PREFETCH_ONE) -#ifdef USE_PARTIAL_PACKETS -#define MICRO_MMA_STORE_ONE(iter) \ - if (unroll_factor > iter) { \ - storeAccumulator<DataMapper, Packet, MICRO_NORMAL_PARTIAL(iter)>(row + iter*accCols, res, pAlpha, accCols2, &accZero##iter); \ +#define MICRO_MMA_STORE_ONE(iter, left, right) \ + if (unroll_factor > left) { \ + storeAccumulator<DataMapper, Packet, MICRO_NORMAL_PARTIAL(left)>(row + left*accCols, res##right, pAlpha, accCols2, &accZero##iter); \ } -#else -#define MICRO_MMA_STORE_ONE(iter) \ - if (unroll_factor > iter) { \ - storeAccumulator<DataMapper, Packet, accCols, (unroll_factor != (iter + 1)) ? accCols : accCols2>(row + iter*accCols, res, pAlpha, pMask, &accZero##iter); \ + +#define MICRO_MMA_ITER_UNROLL(func) \ + if (accItr == 1) { \ + func(0,0,0) func(1,1,0) func(2,2,0) func(3,3,0) \ + func(4,4,0) func(5,5,0) func(6,6,0) func(7,7,0) \ + } else if (accItr == 2) { \ + func(0,0,0) func(1,0,1) func(2,1,0) func(3,1,1) \ + func(4,2,0) func(5,2,1) func(6,3,0) func(7,3,1) \ + } else { \ + func(0,0,0) func(1,0,1) func(2,0,2) func(3,0,3) \ + func(4,1,0) func(5,1,1) func(6,1,2) func(7,1,3) \ } -#endif -#define MICRO_MMA_STORE MICRO_MMA_UNROLL(MICRO_MMA_STORE_ONE) +#define MICRO_MMA_STORE MICRO_MMA_ITER_UNROLL(MICRO_MMA_STORE_ONE) -#ifdef USE_PARTIAL_PACKETS -template<int unroll_factor, typename Scalar, typename Packet, typename RhsPacket, typename DataMapper, const Index accRows, const Index accCols, bool full> -#else -template<int unroll_factor, typename Scalar, typename Packet, typename RhsPacket, typename DataMapper, const Index accRows, const Index accCols, const Index accCols2> -#endif +#define MICRO_MMA_EXTRA_ROWS(right) \ + gemm_extra_row<Scalar, Packet, DataMapper, accRows, accCols>(res3##right, blockA, rhs_base + right*accRows*strideB, depth, strideA, offsetA, strideB, row, rows, remaining_rows, pAlpha, pMask); + +#define MICRO_MMA_EXTRA_ROWS1(val, right) \ + MICRO_MMA_EXTRA_ROWS(right); + +template<int unroll_factor, typename Scalar, typename Packet, typename RhsPacket, typename DataMapper, const Index accRows, const Index accCols, bool full, const Index accItr> EIGEN_ALWAYS_INLINE void gemm_unrolled_MMA_iteration( - const DataMapper& res, + const DataMapper& res0, + const DataMapper& res1, + const DataMapper& res2, + const DataMapper& res3, const Scalar* lhs_base, const Scalar* rhs_base, Index depth, Index strideA, + Index strideB, Index offsetA, Index& row, const Packet& pAlpha, -#ifdef USE_PARTIAL_PACKETS Index accCols2 -#else - const Packet& pMask -#endif ) { - const Scalar* rhs_ptr = rhs_base; + const Scalar* rhs_ptr0 = rhs_base, * rhs_ptr1 = NULL, * rhs_ptr2 = NULL, * rhs_ptr3 = NULL; const Scalar* lhs_ptr0 = NULL, * lhs_ptr1 = NULL, * lhs_ptr2 = NULL, * lhs_ptr3 = NULL, * lhs_ptr4 = NULL, * lhs_ptr5 = NULL, * lhs_ptr6 = NULL, * lhs_ptr7 = NULL; __vector_quad accZero0, accZero1, accZero2, accZero3, accZero4, accZero5, accZero6, accZero7; + if (accItr > 1) { + rhs_ptr1 = rhs_base + (accRows * strideB); + } else { + EIGEN_UNUSED_VARIABLE(strideB); + EIGEN_UNUSED_VARIABLE(rhs_ptr1); + EIGEN_UNUSED_VARIABLE(res1); + } + if (accItr > 2) { + rhs_ptr2 = rhs_base + (2 * accRows * strideB); + rhs_ptr3 = rhs_base + (3 * accRows * strideB); + } else { + EIGEN_UNUSED_VARIABLE(rhs_ptr2); + EIGEN_UNUSED_VARIABLE(rhs_ptr3); + EIGEN_UNUSED_VARIABLE(res2); + EIGEN_UNUSED_VARIABLE(res3); + } + MICRO_MMA_SRC_PTR MICRO_MMA_DST_PTR @@ -347,17 +406,16 @@ MICRO_UPDATE } -#ifdef USE_PARTIAL_PACKETS #define MICRO_MMA_UNROLL_ITER2(N, M) \ - gemm_unrolled_MMA_iteration<N + (M ? 1 : 0), Scalar, Packet, RhsPacket, DataMapper, accRows, accCols, !M>(res3, lhs_base, rhs_base, depth, strideA, offsetA, row, pAlpha, M ? remaining_rows : accCols); \ + gemm_unrolled_MMA_iteration<N + (M ? 1 : 0), Scalar, Packet, RhsPacket, DataMapper, accRows, accCols, !M, accItr>(res30, res31, res32, res33, lhs_base, rhs_base, depth, strideA, strideB, offsetA, row, pAlpha, M ? remaining_rows : accCols); \ if (M) return; -#else -#define MICRO_MMA_UNROLL_ITER2(N, M) \ - gemm_unrolled_MMA_iteration<N + (M ? 1 : 0), Scalar, Packet, RhsPacket, DataMapper, accRows, accCols, M ? M : accCols>(res3, lhs_base, rhs_base, depth, strideA, offsetA, row, pAlpha, pMask); \ - if (M) return; -#endif -template<typename Scalar, typename Packet, typename RhsPacket, typename DataMapper, const Index accRows, const Index accCols> +#define MICRO_MMA_ROWS(n) \ + while(row + n*accCols <= rows) { \ + MICRO_MMA_UNROLL_ITER2(n, 0); \ + } + +template<typename Scalar, typename Packet, typename RhsPacket, typename DataMapper, const Index accRows, const Index accCols, const Index accItr> EIGEN_ALWAYS_INLINE void gemmMMA_cols( const DataMapper& res, const Scalar* blockA, @@ -373,45 +431,71 @@ const Packet& pAlpha, const Packet& pMask) { - const DataMapper res3 = res.getSubMapper(0, col); + const DataMapper res30 = res.getSubMapper(0, col); + const DataMapper res31 = (accItr > 1) ? res30.getSubMapper(0, accRows*1) : res30; + const DataMapper res32 = (accItr > 2) ? res30.getSubMapper(0, accRows*2) : res30; + const DataMapper res33 = (accItr > 2) ? res30.getSubMapper(0, accRows*3) : res30; const Scalar* rhs_base = blockB + col*strideB + accRows*offsetB; const Scalar* lhs_base = blockA + accCols*offsetA; Index row = 0; #define MAX_MMA_UNROLL 7 - while(row + MAX_MMA_UNROLL*accCols <= rows) { - MICRO_MMA_UNROLL_ITER2(MAX_MMA_UNROLL, 0); + +#if MAX_MMA_UNROLL < 2 + if (1) { +#elif MAX_MMA_UNROLL < 4 + if (accItr <= 2) { +#else + if (accItr == 1) { +#endif + MICRO_MMA_ROWS(MAX_MMA_UNROLL); + } else if (accItr == 2) { + MICRO_MMA_ROWS(4); + } else { + MICRO_MMA_ROWS(2); } switch( (rows-row)/accCols ) { #if MAX_MMA_UNROLL > 7 case 7: - MICRO_UNROLL_ITER(MICRO_MMA_UNROLL_ITER2, 7) + if (accItr == 1) { + MICRO_UNROLL_ITER(MICRO_MMA_UNROLL_ITER2, 7) + } break; #endif #if MAX_MMA_UNROLL > 6 case 6: - MICRO_UNROLL_ITER(MICRO_MMA_UNROLL_ITER2, 6) + if (accItr == 1) { + MICRO_UNROLL_ITER(MICRO_MMA_UNROLL_ITER2, 6) + } break; #endif #if MAX_MMA_UNROLL > 5 case 5: - MICRO_UNROLL_ITER(MICRO_MMA_UNROLL_ITER2, 5) + if (accItr == 1) { + MICRO_UNROLL_ITER(MICRO_MMA_UNROLL_ITER2, 5) + } break; #endif #if MAX_MMA_UNROLL > 4 case 4: - MICRO_UNROLL_ITER(MICRO_MMA_UNROLL_ITER2, 4) + if (accItr == 1) { + MICRO_UNROLL_ITER(MICRO_MMA_UNROLL_ITER2, 4) + } break; #endif #if MAX_MMA_UNROLL > 3 case 3: - MICRO_UNROLL_ITER(MICRO_MMA_UNROLL_ITER2, 3) + if (accItr <= 2) { + MICRO_UNROLL_ITER(MICRO_MMA_UNROLL_ITER2, 3) + } break; #endif #if MAX_MMA_UNROLL > 2 case 2: - MICRO_UNROLL_ITER(MICRO_MMA_UNROLL_ITER2, 2) + if (accItr <= 2) { + MICRO_UNROLL_ITER(MICRO_MMA_UNROLL_ITER2, 2) + } break; #endif #if MAX_MMA_UNROLL > 1 @@ -426,10 +510,16 @@ if(remaining_rows > 0) { - gemm_extra_row<Scalar, Packet, DataMapper, accRows, accCols>(res3, blockA, rhs_base, depth, strideA, offsetA, strideB, row, rows, remaining_rows, pAlpha, pMask); + MICRO_MMA_UNROLL_ITER(MICRO_MMA_EXTRA_ROWS1, 0) } } +#define MICRO_MMA_COLS(n) \ + for(; col + n*accRows <= cols; col += n*accRows) \ + { \ + gemmMMA_cols<Scalar, Packet, RhsPacket2, DataMapper, accRows, accCols, n>(res, blockA, blockB, depth, strideA, offsetA, strideB, offsetB, col, rows, remaining_rows, pAlpha, pMask); \ + } + template<typename Scalar, typename Packet, typename RhsPacket, typename DataMapper, const Index accRows, const Index accCols> void gemmMMA(const DataMapper& res, const Scalar* blockA, const Scalar* blockB, Index rows, Index depth, Index cols, Scalar alpha, Index strideA, Index strideB, Index offsetA, Index offsetB) { @@ -444,10 +534,11 @@ typedef typename std::conditional_t<(sizeof(Scalar) == sizeof(float)), RhsPacket, __vector_pair> RhsPacket2; Index col = 0; - for(; col + accRows <= cols; col += accRows) - { - gemmMMA_cols<Scalar, Packet, RhsPacket2, DataMapper, accRows, accCols>(res, blockA, blockB, depth, strideA, offsetA, strideB, offsetB, col, rows, remaining_rows, pAlpha, pMask); - } +#ifdef GEMM_MULTIPLE_COLS + MICRO_MMA_COLS(4); + MICRO_MMA_COLS(2); +#endif + MICRO_MMA_COLS(1); if (col != cols) { @@ -459,62 +550,88 @@ #define advanceCols ((RhsIsReal) ? 1 : 2) // PEEL_COMPLEX_MMA loop factor. +#ifdef GEMM_MULTIPLE_COLS +#define PEEL_COMPLEX_MMA 4 +#else #define PEEL_COMPLEX_MMA 3 +#endif #define MICRO_COMPLEX_MMA_UNROLL(func) \ func(0) func(1) func(2) func(3) #define MICRO_COMPLEX_MMA_WORK(func, type, peel) \ - func(0,type,peel) func(1,type,peel) func(2,type,peel) func(3,type,peel) + if (accItr == 1) { \ + func(0,type,peel,0,0) func(1,type,peel,1,0) func(2,type,peel,2,0) func(3,type,peel,3,0) \ + } else if (accItr == 2) { \ + func(0,type,peel,0,0) func(1,type,peel,0,1) func(2,type,peel,1,0) func(3,type,peel,1,1) \ + } else { \ + func(0,type,peel,0,0) func(1,type,peel,0,1) func(2,type,peel,0,2) func(3,type,peel,0,3) \ + } -#define MICRO_COMPLEX_MMA_WORK_ONE(iter, type, peel) \ - if (unroll_factor > iter) { \ - pgercMMA<Packet, type, ConjugateLhs, ConjugateRhs, LhsIsReal, RhsIsReal>(&accReal##iter, &accImag##iter, lhsV##iter, lhsVi##iter, rhsV[peel], rhsVi[peel]); \ +#define MICRO_COMPLEX_MMA_WORK_ONE(iter, type, peel, left, right) \ + if (unroll_factor > left) { \ + pgercMMA<Packet, type, ConjugateLhs, ConjugateRhs, LhsIsReal, RhsIsReal>(&accReal##iter, &accImag##iter, lhsV##left, lhsVi##left, rhsV##right[peel], rhsVi##right[peel]); \ } #ifdef VECTOR_PAIR_LOADS_LHS -#define MICRO_COMPLEX_MMA_WORK_TWO(iter, type, peel) \ - if (unroll_factor > iter) { \ - pgercMMA<Packet, type, ConjugateLhs, ConjugateRhs, LhsIsReal, RhsIsReal>(&accReal##iter, &accImag##iter, lhsV2##iter.packet[peel & 1], lhsVi2##iter.packet[peel & 1], rhsV[peel], rhsVi[peel]); \ +#define MICRO_COMPLEX_MMA_WORK_TWO(iter, type, peel, left, right) \ + if (unroll_factor > left) { \ + pgercMMA<Packet, type, ConjugateLhs, ConjugateRhs, LhsIsReal, RhsIsReal>(&accReal##iter, &accImag##iter, lhsV2##left.packet[peel & 1], lhsVi2##left.packet[peel & 1], rhsV##right[peel], rhsVi##right[peel]); \ } -#define MICRO_COMPLEX_MMA_LOAD1_TWO(lhs_ptr, iter) \ - if (!LhsIsReal && (unroll_factor > iter)) { \ - if (MICRO_NORMAL(iter)) { \ - ploadLhsMMA(reinterpret_cast<const double*>(lhs_ptr_real##iter + imag_delta), plhsVi##iter); \ - __builtin_vsx_disassemble_pair(reinterpret_cast<void*>(&lhsVi2##iter.packet), &plhsVi##iter); \ +#define MICRO_COMPLEX_MMA_LOAD1_TWO(lhs_ptr, left) \ + if (!LhsIsReal && (unroll_factor > left)) { \ + if (MICRO_NORMAL(left)) { \ + ploadLhsMMA(reinterpret_cast<const double*>(lhs_ptr_real##left + imag_delta), plhsVi##left); \ + __builtin_vsx_disassemble_pair(reinterpret_cast<void*>(&lhsVi2##left.packet), &plhsVi##left); \ } else { \ - lhsVi2##iter.packet[0] = ploadLhs<Packet>(lhs_ptr_real##iter + imag_delta2); \ - lhsVi2##iter.packet[1] = ploadLhs<Packet>(lhs_ptr_real##iter + imag_delta2 + accCols2); \ - EIGEN_UNUSED_VARIABLE(plhsVi##iter) \ + lhsVi2##left.packet[0] = ploadLhs<Packet>(lhs_ptr_real##left + imag_delta2); \ + lhsVi2##left.packet[1] = ploadLhs<Packet>(lhs_ptr_real##left + imag_delta2 + accCols2); \ + EIGEN_UNUSED_VARIABLE(plhsVi##left); \ } \ } else { \ - EIGEN_UNUSED_VARIABLE(lhsVi2##iter); \ - EIGEN_UNUSED_VARIABLE(plhsVi##iter) \ + EIGEN_UNUSED_VARIABLE(lhsVi2##left); \ + EIGEN_UNUSED_VARIABLE(plhsVi##left); \ } \ - MICRO_MMA_LOAD1_TWO(lhs_ptr_real, iter) + MICRO_MMA_LOAD1_TWO(lhs_ptr_real, left) -#define MICRO_COMPLEX_MMA_LOAD_TWO(iter) MICRO_COMPLEX_MMA_LOAD1_TWO(lhs_ptr, iter) +#define MICRO_COMPLEX_MMA_LOAD_TWO(left) MICRO_COMPLEX_MMA_LOAD1_TWO(lhs_ptr, left) #endif +#define MICRO_COMPLEX_MMA_LOAD_RHS1(peel, right) \ + ploadRhsMMA(rhs_ptr_real##right + (accRows * peel), rhsV##right[peel]); \ + if (!RhsIsReal) { \ + ploadRhsMMA(rhs_ptr_imag##right + (accRows * peel), rhsVi##right[peel]); \ + } + +#define MICRO_COMPLEX_MMA_LOAD_ONE_RHS(peel) \ + MICRO_MMA_UNROLL_ITER(MICRO_COMPLEX_MMA_LOAD_RHS1, peel) + #define MICRO_COMPLEX_MMA_TYPE_PEEL(funcw, funcl, type, peel) \ if (PEEL_COMPLEX_MMA > peel) { \ Packet lhsV0, lhsV1, lhsV2, lhsV3; \ Packet lhsVi0, lhsVi1, lhsVi2, lhsVi3; \ - ploadRhsMMA(rhs_ptr_real + (accRows * peel), rhsV[peel]); \ - if(!RhsIsReal) { \ - ploadRhsMMA(rhs_ptr_imag + (accRows * peel), rhsVi[peel]); \ - } \ + MICRO_COMPLEX_MMA_LOAD_ONE_RHS(peel) \ MICRO_COMPLEX_MMA_UNROLL(funcl) \ MICRO_COMPLEX_MMA_WORK(funcw, type, peel) \ } #ifndef VECTOR_PAIR_LOADS_LHS #define MICRO_COMPLEX_MMA_UNROLL_TYPE_PEEL(funcw, funcl, type) \ - type rhsV[4], rhsVi[4]; \ + type rhsV0[4], rhsVi0[4], rhsV1[(accItr > 1) ? 4 : 1], rhsVi1[(accItr > 1) ? 4 : 1], rhsV2[(accItr > 2) ? 4 : 1], rhsVi2[(accItr > 2) ? 4 : 1], rhsV3[(accItr > 2) ? 4 : 1], rhsVi3[(accItr > 2) ? 4 : 1]; \ MICRO_COMPLEX_MMA_TYPE_PEEL(funcw,funcl,type,0) MICRO_COMPLEX_MMA_TYPE_PEEL(funcw,funcl,type,1) \ MICRO_COMPLEX_MMA_TYPE_PEEL(funcw,funcl,type,2) MICRO_COMPLEX_MMA_TYPE_PEEL(funcw,funcl,type,3) #else +#define MICRO_COMPLEX_MMA_LOAD_TWO_RHS(peel1, right) \ + ploadRhsMMA(reinterpret_cast<const double*>(rhs_ptr_real##right + (accRows * peel1)), prhsV##peel1); \ + __builtin_vsx_disassemble_pair(reinterpret_cast<void*>(&rhsV##right[peel1]), &prhsV##peel1); \ + if(!RhsIsReal) { \ + ploadRhsMMA(reinterpret_cast<const double*>(rhs_ptr_imag##right + (accRows * peel1)), prhsVi##peel1); \ + __builtin_vsx_disassemble_pair(reinterpret_cast<void*>(&rhsVi##right[peel1]), &prhsVi##peel1); \ + } else { \ + EIGEN_UNUSED_VARIABLE(prhsVi##peel1); \ + } + #define MICRO_COMPLEX_MMA_TYPE_PEEL2(funcw1, funcl1, funcw2, funcl2, type, peel1, peel2) \ if (PEEL_COMPLEX_MMA > peel2) { \ PacketBlock<Packet,2> lhsV20, lhsV21, lhsV22, lhsV23; \ @@ -522,23 +639,12 @@ __vector_pair plhsV0, plhsV1, plhsV2, plhsV3; \ __vector_pair plhsVi0, plhsVi1, plhsVi2, plhsVi3; \ if (sizeof(type) == 16) { \ - ploadRhsMMA(reinterpret_cast<const double*>(rhs_ptr_real + (accRows * peel1)), prhsV##peel1); \ - __builtin_vsx_disassemble_pair(reinterpret_cast<void*>(&rhsV[peel1]), &prhsV##peel1); \ - if(!RhsIsReal) { \ - ploadRhsMMA(reinterpret_cast<const double*>(rhs_ptr_imag + (accRows * peel1)), prhsVi##peel1); \ - __builtin_vsx_disassemble_pair(reinterpret_cast<void*>(&rhsVi[peel1]), &prhsVi##peel1); \ - } else { \ - EIGEN_UNUSED_VARIABLE(prhsVi##peel1); \ - } \ + MICRO_MMA_UNROLL_ITER(MICRO_COMPLEX_MMA_LOAD_TWO_RHS, peel1) \ } else { \ EIGEN_UNUSED_VARIABLE(prhsV##peel1); \ EIGEN_UNUSED_VARIABLE(prhsVi##peel1); \ - ploadRhsMMA(rhs_ptr_real + (accRows * peel1), rhsV[peel1]); \ - ploadRhsMMA(rhs_ptr_real + (accRows * peel2), rhsV[peel2]); \ - if(!RhsIsReal) { \ - ploadRhsMMA(rhs_ptr_imag + (accRows * peel1), rhsVi[peel1]); \ - ploadRhsMMA(rhs_ptr_imag + (accRows * peel2), rhsVi[peel2]); \ - } \ + MICRO_COMPLEX_MMA_LOAD_ONE_RHS(peel1); \ + MICRO_COMPLEX_MMA_LOAD_ONE_RHS(peel2); \ } \ MICRO_COMPLEX_MMA_UNROLL(funcl2) \ MICRO_COMPLEX_MMA_WORK(funcw2, type, peel1) \ @@ -550,7 +656,7 @@ } #define MICRO_COMPLEX_MMA_UNROLL_TYPE_PEEL2(funcw1, funcl1, funcw2, funcl2, type) \ - type rhsV[4], rhsVi[4]; \ + type rhsV0[4], rhsVi0[4], rhsV1[(accItr > 1) ? 4 : 1], rhsVi1[(accItr > 1) ? 4 : 1], rhsV2[(accItr > 2) ? 4 : 1], rhsVi2[(accItr > 2) ? 4 : 1], rhsV3[(accItr > 2) ? 4 : 1], rhsVi3[(accItr > 2) ? 4 : 1]; \ __vector_pair prhsV0, prhsV2; \ __vector_pair prhsVi0, prhsVi2; \ MICRO_COMPLEX_MMA_TYPE_PEEL2(funcw1,funcl1,funcw2,funcl2,type,0,1) \ @@ -558,21 +664,26 @@ #endif #define MICRO_COMPLEX_MMA_UNROLL_TYPE_ONE(funcw, funcl, type) \ - type rhsV[1], rhsVi[1]; \ + type rhsV0[1], rhsVi0[1], rhsV1[1], rhsVi1[1], rhsV2[1], rhsVi2[1], rhsV3[1], rhsVi3[1]; \ MICRO_COMPLEX_MMA_TYPE_PEEL(funcw,funcl,type,0) +#define MICRO_COMPLEX_MMA_UPDATE_RHS1(size, right) \ + rhs_ptr_real##right += (accRows * size); \ + if(!RhsIsReal) rhs_ptr_imag##right += (accRows * size); + +#define MICRO_COMPLEX_MMA_UPDATE_RHS(size) \ + MICRO_MMA_UNROLL_ITER(MICRO_COMPLEX_MMA_UPDATE_RHS1, size) + #define MICRO_COMPLEX_MMA_UNROLL_TYPE(MICRO_COMPLEX_MMA_TYPE, size) \ MICRO_COMPLEX_MMA_TYPE(MICRO_COMPLEX_MMA_WORK_ONE, MICRO_COMPLEX_LOAD_ONE, RhsPacket) \ - rhs_ptr_real += (accRows * size); \ - if(!RhsIsReal) rhs_ptr_imag += (accRows * size); + MICRO_COMPLEX_MMA_UPDATE_RHS(size); #ifndef VECTOR_PAIR_LOADS_LHS #define MICRO_COMPLEX_MMA_ONE_PEEL MICRO_COMPLEX_MMA_UNROLL_TYPE(MICRO_COMPLEX_MMA_UNROLL_TYPE_PEEL, PEEL_COMPLEX_MMA) #else #define MICRO_COMPLEX_MMA_UNROLL_TYPE2(MICRO_COMPLEX_MMA_TYPE, size) \ MICRO_COMPLEX_MMA_TYPE(MICRO_COMPLEX_MMA_WORK_ONE, MICRO_COMPLEX_LOAD_ONE, MICRO_COMPLEX_MMA_WORK_TWO, MICRO_COMPLEX_MMA_LOAD_TWO, RhsPacket) \ - rhs_ptr_real += (accRows * size); \ - if(!RhsIsReal) rhs_ptr_imag += (accRows * size); + MICRO_COMPLEX_MMA_UPDATE_RHS(size); #define MICRO_COMPLEX_MMA_ONE_PEEL MICRO_COMPLEX_MMA_UNROLL_TYPE2(MICRO_COMPLEX_MMA_UNROLL_TYPE_PEEL2, PEEL_COMPLEX_MMA) #endif @@ -580,7 +691,7 @@ #define MICRO_COMPLEX_MMA_ONE MICRO_COMPLEX_MMA_UNROLL_TYPE(MICRO_COMPLEX_MMA_UNROLL_TYPE_ONE, 1) #define MICRO_COMPLEX_MMA_DST_PTR_ONE(iter) \ - if (unroll_factor > iter) { \ + if (unroll_factor * accItr > iter) { \ bsetzeroMMA(&accReal##iter); \ bsetzeroMMA(&accImag##iter); \ } else { \ @@ -594,16 +705,34 @@ #define MICRO_COMPLEX_MMA_PREFETCH MICRO_COMPLEX_MMA_UNROLL(MICRO_COMPLEX_PREFETCH_ONE) -#define MICRO_COMPLEX_MMA_STORE_ONE(iter) \ - if (unroll_factor > iter) { \ - storeComplexAccumulator<DataMapper, Packet, Packetc, accCols, (unroll_factor != (iter + 1)) ? accCols : accCols2>(row + iter*accCols, res, pAlphaReal, pAlphaImag, pMask, &accReal##iter, &accImag##iter); \ +#define MICRO_COMPLEX_MMA_STORE_ONE(iter, left, right) \ + if (unroll_factor > left) { \ + storeComplexAccumulator<DataMapper, Packet, Packetc, accCols, (unroll_factor != (left + 1)) ? accCols : accCols2>(row + left*accCols, res##right, pAlphaReal, pAlphaImag, pMask, &accReal##iter, &accImag##iter); \ } -#define MICRO_COMPLEX_MMA_STORE MICRO_COMPLEX_MMA_UNROLL(MICRO_COMPLEX_MMA_STORE_ONE) +#define MICRO_COMPLEX_MMA_ITER_UNROLL(func) \ + if (accItr == 1) { \ + func(0,0,0) func(1,1,0) func(2,2,0) func(3,3,0) \ + } else if (accItr == 2) { \ + func(0,0,0) func(1,0,1) func(2,1,0) func(3,1,1) \ + } else { \ + func(0,0,0) func(1,0,1) func(2,0,2) func(3,0,3) \ + } -template<int unroll_factor, typename Scalar, typename Packet, typename Packetc, typename RhsPacket, typename DataMapper, const Index accRows, const Index accCols, const Index accCols2, bool ConjugateLhs, bool ConjugateRhs, bool LhsIsReal, bool RhsIsReal> +#define MICRO_COMPLEX_MMA_STORE MICRO_COMPLEX_MMA_ITER_UNROLL(MICRO_COMPLEX_MMA_STORE_ONE) + +#define MICRO_COMPLEX_MMA_EXTRA_ROWS(right) \ + gemm_complex_extra_row<Scalar, Packet, Packetc, DataMapper, accRows, accCols, ConjugateLhs, ConjugateRhs, LhsIsReal, RhsIsReal>(res3##right, blockA, rhs_base + right*accRows*(RhsIsReal ? 1 : 2)*strideB, depth, strideA, offsetA, strideB, row, rows, remaining_rows, pAlphaReal, pAlphaImag, pMask); + +#define MICRO_COMPLEX_MMA_EXTRA_ROWS1(val, right) \ + MICRO_COMPLEX_MMA_EXTRA_ROWS(right); + +template<int unroll_factor, typename Scalar, typename Packet, typename Packetc, typename RhsPacket, typename DataMapper, const Index accRows, const Index accCols, const Index accCols2, bool ConjugateLhs, bool ConjugateRhs, bool LhsIsReal, bool RhsIsReal, const Index accItr> EIGEN_ALWAYS_INLINE void gemm_complex_unrolled_MMA_iteration( - const DataMapper& res, + const DataMapper& res0, + const DataMapper& res1, + const DataMapper& res2, + const DataMapper& res3, const Scalar* lhs_base, const Scalar* rhs_base, Index depth, @@ -615,14 +744,48 @@ const Packet& pAlphaImag, const Packet& pMask) { - const Scalar* rhs_ptr_real = rhs_base; - const Scalar* rhs_ptr_imag = NULL; + const Scalar* rhs_ptr_real0 = rhs_base, * rhs_ptr_real1 = NULL, * rhs_ptr_real2 = NULL, * rhs_ptr_real3 = NULL; + const Scalar* rhs_ptr_imag0 = NULL, * rhs_ptr_imag1 = NULL, * rhs_ptr_imag2 = NULL, * rhs_ptr_imag3 = NULL; const Index imag_delta = accCols*strideA; const Index imag_delta2 = accCols2*strideA; + if(!RhsIsReal) { - rhs_ptr_imag = rhs_base + accRows*strideB; + rhs_ptr_imag0 = rhs_base + accRows*strideB; } else { - EIGEN_UNUSED_VARIABLE(rhs_ptr_imag); + EIGEN_UNUSED_VARIABLE(rhs_ptr_imag0); + } + if (accItr > 1) { + if(!RhsIsReal) { + rhs_ptr_real1 = rhs_base + (2*accRows*strideB); + rhs_ptr_imag1 = rhs_base + (3*accRows*strideB); + } else { + rhs_ptr_real1 = rhs_base + accRows*strideB; + EIGEN_UNUSED_VARIABLE(rhs_ptr_imag1); + } + } else { + EIGEN_UNUSED_VARIABLE(rhs_ptr_real1); + EIGEN_UNUSED_VARIABLE(rhs_ptr_imag1); + EIGEN_UNUSED_VARIABLE(res1); + } + if (accItr > 2) { + if(!RhsIsReal) { + rhs_ptr_real2 = rhs_base + (4*accRows*strideB); + rhs_ptr_imag2 = rhs_base + (5*accRows*strideB); + rhs_ptr_real3 = rhs_base + (6*accRows*strideB); + rhs_ptr_imag3 = rhs_base + (7*accRows*strideB); + } else { + rhs_ptr_real2 = rhs_base + (2*accRows*strideB); + rhs_ptr_real3 = rhs_base + (3*accRows*strideB); + EIGEN_UNUSED_VARIABLE(rhs_ptr_imag2); + EIGEN_UNUSED_VARIABLE(rhs_ptr_imag3); + } + } else { + EIGEN_UNUSED_VARIABLE(rhs_ptr_real2); + EIGEN_UNUSED_VARIABLE(rhs_ptr_real3); + EIGEN_UNUSED_VARIABLE(rhs_ptr_imag2); + EIGEN_UNUSED_VARIABLE(rhs_ptr_imag3); + EIGEN_UNUSED_VARIABLE(res2); + EIGEN_UNUSED_VARIABLE(res3); } const Scalar* lhs_ptr_real0 = NULL, * lhs_ptr_real1 = NULL; const Scalar* lhs_ptr_real2 = NULL, * lhs_ptr_real3 = NULL; @@ -651,10 +814,15 @@ } #define MICRO_COMPLEX_MMA_UNROLL_ITER2(N, M) \ - gemm_complex_unrolled_MMA_iteration<N + (M ? 1 : 0), Scalar, Packet, Packetc, RhsPacket, DataMapper, accRows, accCols, M ? M : accCols, ConjugateLhs, ConjugateRhs, LhsIsReal, RhsIsReal>(res3, lhs_base, rhs_base, depth, strideA, offsetA, strideB, row, pAlphaReal, pAlphaImag, pMask); \ + gemm_complex_unrolled_MMA_iteration<N + (M ? 1 : 0), Scalar, Packet, Packetc, RhsPacket, DataMapper, accRows, accCols, M ? M : accCols, ConjugateLhs, ConjugateRhs, LhsIsReal, RhsIsReal, accItr>(res30, res31, res32, res33, lhs_base, rhs_base, depth, strideA, offsetA, strideB, row, pAlphaReal, pAlphaImag, pMask); \ if (M) return; -template<typename Scalar, typename Packet, typename Packetc, typename RhsPacket, typename DataMapper, const Index accRows, const Index accCols, bool ConjugateLhs, bool ConjugateRhs, bool LhsIsReal, bool RhsIsReal> +#define MICRO_COMPLEX_MMA_ROWS(n) \ + while(row + n*accCols <= rows) { \ + MICRO_COMPLEX_MMA_UNROLL_ITER2(n, 0); \ + } + +template<typename Scalar, typename Packet, typename Packetc, typename RhsPacket, typename DataMapper, const Index accRows, const Index accCols, bool ConjugateLhs, bool ConjugateRhs, bool LhsIsReal, bool RhsIsReal, const Index accItr> EIGEN_ALWAYS_INLINE void gemmMMA_complex_cols( const DataMapper& res, const Scalar* blockA, @@ -671,35 +839,50 @@ const Packet& pAlphaImag, const Packet& pMask) { - const DataMapper res3 = res.getSubMapper(0, col); + const DataMapper res30 = res.getSubMapper(0, col); + const DataMapper res31 = (accItr > 1) ? res30.getSubMapper(0, accRows*1) : res30; + const DataMapper res32 = (accItr > 2) ? res30.getSubMapper(0, accRows*2) : res30; + const DataMapper res33 = (accItr > 2) ? res30.getSubMapper(0, accRows*3) : res30; const Scalar* rhs_base = blockB + advanceCols*col*strideB + accRows*offsetB; const Scalar* lhs_base = blockA + accCols*offsetA; Index row = 0; #define MAX_COMPLEX_MMA_UNROLL 4 - while(row + MAX_COMPLEX_MMA_UNROLL*accCols <= rows) { - MICRO_COMPLEX_MMA_UNROLL_ITER2(MAX_COMPLEX_MMA_UNROLL, 0); + +#if MAX_COMPLEX_MMA_UNROLL < 2 + if (1) { +#elif MAX_COMPLEX_MMA_UNROLL < 4 + if (accItr <= 2) { +#else + if (accItr == 1) { +#endif + MICRO_COMPLEX_MMA_ROWS(MAX_COMPLEX_MMA_UNROLL); + } else if (accItr == 2) { + MICRO_COMPLEX_MMA_ROWS(2); + } else { + MICRO_COMPLEX_MMA_ROWS(1); } switch( (rows-row)/accCols ) { -#if MAX_COMPLEX_MMA_UNROLL > 4 - case 4: - MICRO_COMPLEX_UNROLL_ITER(MICRO_COMPLEX_MMA_UNROLL_ITER2, 4) - break; -#endif #if MAX_COMPLEX_MMA_UNROLL > 3 case 3: - MICRO_COMPLEX_UNROLL_ITER(MICRO_COMPLEX_MMA_UNROLL_ITER2, 3) + if (accItr == 1) { + MICRO_COMPLEX_UNROLL_ITER(MICRO_COMPLEX_MMA_UNROLL_ITER2, 3) + } break; #endif #if MAX_COMPLEX_MMA_UNROLL > 2 case 2: - MICRO_COMPLEX_UNROLL_ITER(MICRO_COMPLEX_MMA_UNROLL_ITER2, 2) + if (accItr == 1) { + MICRO_COMPLEX_UNROLL_ITER(MICRO_COMPLEX_MMA_UNROLL_ITER2, 2) + } break; #endif #if MAX_COMPLEX_MMA_UNROLL > 1 case 1: - MICRO_COMPLEX_UNROLL_ITER(MICRO_COMPLEX_MMA_UNROLL_ITER2, 1) + if (accItr <= 2) { + MICRO_COMPLEX_UNROLL_ITER(MICRO_COMPLEX_MMA_UNROLL_ITER2, 1) + } break; #endif default: @@ -709,10 +892,16 @@ if(remaining_rows > 0) { - gemm_complex_extra_row<Scalar, Packet, Packetc, DataMapper, accRows, accCols, ConjugateLhs, ConjugateRhs, LhsIsReal, RhsIsReal>(res3, blockA, rhs_base, depth, strideA, offsetA, strideB, row, rows, remaining_rows, pAlphaReal, pAlphaImag, pMask); + MICRO_MMA_UNROLL_ITER(MICRO_COMPLEX_MMA_EXTRA_ROWS1, 0) } } +#define MICRO_COMPLEX_MMA_COLS(n) \ + for(; col + n*accRows <= cols; col += n*accRows) \ + { \ + gemmMMA_complex_cols<Scalar, Packet, Packetc, RhsPacket2, DataMapper, accRows, accCols, ConjugateLhs, ConjugateRhs, LhsIsReal, RhsIsReal, n>(res, blockA, blockB, depth, strideA, offsetA, strideB, offsetB, col, rows, remaining_rows, pAlphaReal, pAlphaImag, pMask); \ + } + template<typename LhsScalar, typename RhsScalar, typename Scalarc, typename Scalar, typename Packet, typename Packetc, typename RhsPacket, typename DataMapper, const Index accRows, const Index accCols, bool ConjugateLhs, bool ConjugateRhs, bool LhsIsReal, bool RhsIsReal> void gemm_complexMMA(const DataMapper& res, const LhsScalar* blockAc, const RhsScalar* blockBc, Index rows, Index depth, Index cols, Scalarc alpha, Index strideA, Index strideB, Index offsetA, Index offsetB) { @@ -731,10 +920,11 @@ typedef typename std::conditional_t<(sizeof(Scalar) == sizeof(float)), RhsPacket, __vector_pair> RhsPacket2; Index col = 0; - for(; col + accRows <= cols; col += accRows) - { - gemmMMA_complex_cols<Scalar, Packet, Packetc, RhsPacket2, DataMapper, accRows, accCols, ConjugateLhs, ConjugateRhs, LhsIsReal, RhsIsReal>(res, blockA, blockB, depth, strideA, offsetA, strideB, offsetB, col, rows, remaining_rows, pAlphaReal, pAlphaImag, pMask); - } +#ifdef GEMM_MULTIPLE_COLS + MICRO_COMPLEX_MMA_COLS(4); + MICRO_COMPLEX_MMA_COLS(2); +#endif + MICRO_COMPLEX_MMA_COLS(1); if (col != cols) {
diff --git a/Eigen/src/Core/arch/SYCL/InteropHeaders.h b/Eigen/src/Core/arch/SYCL/InteropHeaders.h index ed4edc1..27d9a82 100644 --- a/Eigen/src/Core/arch/SYCL/InteropHeaders.h +++ b/Eigen/src/Core/arch/SYCL/InteropHeaders.h
@@ -86,6 +86,8 @@ typedef packet_type half; \ }; +SYCL_PACKET_TRAITS(cl::sycl::cl_half8, 1, Eigen::half, 8) +SYCL_PACKET_TRAITS(cl::sycl::cl_half8, 1, const Eigen::half, 8) SYCL_PACKET_TRAITS(cl::sycl::cl_float4, 1, float, 4) SYCL_PACKET_TRAITS(cl::sycl::cl_float4, 1, const float, 4) SYCL_PACKET_TRAITS(cl::sycl::cl_double2, 0, double, 2) @@ -100,6 +102,7 @@ struct is_arithmetic<packet_type> { \ enum { value = true }; \ }; +SYCL_ARITHMETIC(cl::sycl::cl_half8) SYCL_ARITHMETIC(cl::sycl::cl_float4) SYCL_ARITHMETIC(cl::sycl::cl_double2) #undef SYCL_ARITHMETIC @@ -111,6 +114,7 @@ enum { size = lengths, vectorizable = true, alignment = Aligned16 }; \ typedef packet_type half; \ }; +SYCL_UNPACKET_TRAITS(cl::sycl::cl_half8, Eigen::half, 8) SYCL_UNPACKET_TRAITS(cl::sycl::cl_float4, float, 4) SYCL_UNPACKET_TRAITS(cl::sycl::cl_double2, double, 2)
diff --git a/Eigen/src/Core/arch/SYCL/MathFunctions.h b/Eigen/src/Core/arch/SYCL/MathFunctions.h index bdbb21f..a8adc46 100644 --- a/Eigen/src/Core/arch/SYCL/MathFunctions.h +++ b/Eigen/src/Core/arch/SYCL/MathFunctions.h
@@ -38,6 +38,7 @@ return cl::sycl::log(a); \ } +SYCL_PLOG(cl::sycl::cl_half8) SYCL_PLOG(cl::sycl::cl_float4) SYCL_PLOG(cl::sycl::cl_double2) #undef SYCL_PLOG @@ -49,6 +50,7 @@ return cl::sycl::log1p(a); \ } +SYCL_PLOG1P(cl::sycl::cl_half8) SYCL_PLOG1P(cl::sycl::cl_float4) SYCL_PLOG1P(cl::sycl::cl_double2) #undef SYCL_PLOG1P @@ -60,6 +62,7 @@ return cl::sycl::log10(a); \ } +SYCL_PLOG10(cl::sycl::cl_half8) SYCL_PLOG10(cl::sycl::cl_float4) SYCL_PLOG10(cl::sycl::cl_double2) #undef SYCL_PLOG10 @@ -71,6 +74,8 @@ return cl::sycl::exp(a); \ } +SYCL_PEXP(cl::sycl::cl_half8) +SYCL_PEXP(cl::sycl::cl_half) SYCL_PEXP(cl::sycl::cl_float4) SYCL_PEXP(cl::sycl::cl_float) SYCL_PEXP(cl::sycl::cl_double2) @@ -83,6 +88,7 @@ return cl::sycl::expm1(a); \ } +SYCL_PEXPM1(cl::sycl::cl_half8) SYCL_PEXPM1(cl::sycl::cl_float4) SYCL_PEXPM1(cl::sycl::cl_double2) #undef SYCL_PEXPM1 @@ -94,6 +100,7 @@ return cl::sycl::sqrt(a); \ } +SYCL_PSQRT(cl::sycl::cl_half8) SYCL_PSQRT(cl::sycl::cl_float4) SYCL_PSQRT(cl::sycl::cl_double2) #undef SYCL_PSQRT @@ -105,6 +112,7 @@ return cl::sycl::rsqrt(a); \ } +SYCL_PRSQRT(cl::sycl::cl_half8) SYCL_PRSQRT(cl::sycl::cl_float4) SYCL_PRSQRT(cl::sycl::cl_double2) #undef SYCL_PRSQRT @@ -117,6 +125,7 @@ return cl::sycl::sin(a); \ } +SYCL_PSIN(cl::sycl::cl_half8) SYCL_PSIN(cl::sycl::cl_float4) SYCL_PSIN(cl::sycl::cl_double2) #undef SYCL_PSIN @@ -129,6 +138,7 @@ return cl::sycl::cos(a); \ } +SYCL_PCOS(cl::sycl::cl_half8) SYCL_PCOS(cl::sycl::cl_float4) SYCL_PCOS(cl::sycl::cl_double2) #undef SYCL_PCOS @@ -141,6 +151,7 @@ return cl::sycl::tan(a); \ } +SYCL_PTAN(cl::sycl::cl_half8) SYCL_PTAN(cl::sycl::cl_float4) SYCL_PTAN(cl::sycl::cl_double2) #undef SYCL_PTAN @@ -153,6 +164,7 @@ return cl::sycl::asin(a); \ } +SYCL_PASIN(cl::sycl::cl_half8) SYCL_PASIN(cl::sycl::cl_float4) SYCL_PASIN(cl::sycl::cl_double2) #undef SYCL_PASIN @@ -165,6 +177,7 @@ return cl::sycl::acos(a); \ } +SYCL_PACOS(cl::sycl::cl_half8) SYCL_PACOS(cl::sycl::cl_float4) SYCL_PACOS(cl::sycl::cl_double2) #undef SYCL_PACOS @@ -177,6 +190,7 @@ return cl::sycl::atan(a); \ } +SYCL_PATAN(cl::sycl::cl_half8) SYCL_PATAN(cl::sycl::cl_float4) SYCL_PATAN(cl::sycl::cl_double2) #undef SYCL_PATAN @@ -189,6 +203,7 @@ return cl::sycl::sinh(a); \ } +SYCL_PSINH(cl::sycl::cl_half8) SYCL_PSINH(cl::sycl::cl_float4) SYCL_PSINH(cl::sycl::cl_double2) #undef SYCL_PSINH @@ -201,6 +216,7 @@ return cl::sycl::cosh(a); \ } +SYCL_PCOSH(cl::sycl::cl_half8) SYCL_PCOSH(cl::sycl::cl_float4) SYCL_PCOSH(cl::sycl::cl_double2) #undef SYCL_PCOSH @@ -213,6 +229,7 @@ return cl::sycl::tanh(a); \ } +SYCL_PTANH(cl::sycl::cl_half8) SYCL_PTANH(cl::sycl::cl_float4) SYCL_PTANH(cl::sycl::cl_double2) #undef SYCL_PTANH @@ -224,6 +241,7 @@ return cl::sycl::ceil(a); \ } +SYCL_PCEIL(cl::sycl::cl_half) SYCL_PCEIL(cl::sycl::cl_float4) SYCL_PCEIL(cl::sycl::cl_double2) #undef SYCL_PCEIL @@ -235,6 +253,7 @@ return cl::sycl::round(a); \ } +SYCL_PROUND(cl::sycl::cl_half8) SYCL_PROUND(cl::sycl::cl_float4) SYCL_PROUND(cl::sycl::cl_double2) #undef SYCL_PROUND @@ -246,6 +265,7 @@ return cl::sycl::rint(a); \ } +SYCL_PRINT(cl::sycl::cl_half8) SYCL_PRINT(cl::sycl::cl_float4) SYCL_PRINT(cl::sycl::cl_double2) #undef SYCL_PRINT @@ -257,6 +277,7 @@ return cl::sycl::floor(a); \ } +SYCL_FLOOR(cl::sycl::cl_half8) SYCL_FLOOR(cl::sycl::cl_float4) SYCL_FLOOR(cl::sycl::cl_double2) #undef SYCL_FLOOR @@ -268,6 +289,7 @@ return expr; \ } +SYCL_PMIN(cl::sycl::cl_half8, cl::sycl::fmin(a, b)) SYCL_PMIN(cl::sycl::cl_float4, cl::sycl::fmin(a, b)) SYCL_PMIN(cl::sycl::cl_double2, cl::sycl::fmin(a, b)) #undef SYCL_PMIN @@ -279,6 +301,7 @@ return expr; \ } +SYCL_PMAX(cl::sycl::cl_half8, cl::sycl::fmax(a, b)) SYCL_PMAX(cl::sycl::cl_float4, cl::sycl::fmax(a, b)) SYCL_PMAX(cl::sycl::cl_double2, cl::sycl::fmax(a, b)) #undef SYCL_PMAX @@ -292,6 +315,7 @@ cl::sycl::rounding_mode::automatic>()); \ } +SYCL_PLDEXP(cl::sycl::cl_half8) SYCL_PLDEXP(cl::sycl::cl_float4) SYCL_PLDEXP(cl::sycl::cl_double2) #undef SYCL_PLDEXP
diff --git a/Eigen/src/Core/arch/SYCL/PacketMath.h b/Eigen/src/Core/arch/SYCL/PacketMath.h index fbdcdc0..4b0b1c6 100644 --- a/Eigen/src/Core/arch/SYCL/PacketMath.h +++ b/Eigen/src/Core/arch/SYCL/PacketMath.h
@@ -44,9 +44,34 @@ SYCL_PLOAD(cl::sycl::cl_float4, ) SYCL_PLOAD(cl::sycl::cl_double2, u) SYCL_PLOAD(cl::sycl::cl_double2, ) - #undef SYCL_PLOAD +template <> +EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_half8 + pload<cl::sycl::cl_half8>( + const typename unpacket_traits<cl::sycl::cl_half8>::type* from) { + auto ptr = cl::sycl::address_space_cast< + cl::sycl::access::address_space::generic_space, + cl::sycl::access::decorated::no>( + reinterpret_cast<const cl::sycl::cl_half*>(from)); + cl::sycl::cl_half8 res{}; + res.load(0, ptr); + return res; +} + +template <> +EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_half8 +ploadu<cl::sycl::cl_half8>( + const typename unpacket_traits<cl::sycl::cl_half8>::type* from) { + auto ptr = cl::sycl::address_space_cast< + cl::sycl::access::address_space::generic_space, + cl::sycl::access::decorated::no>( + reinterpret_cast<const cl::sycl::cl_half*>(from)); + cl::sycl::cl_half8 res{}; + res.load(0, ptr); + return res; +} + #define SYCL_PSTORE(scalar, packet_type, alignment) \ template <> \ EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void pstore##alignment( \ @@ -59,9 +84,28 @@ SYCL_PSTORE(float, cl::sycl::cl_float4, u) SYCL_PSTORE(double, cl::sycl::cl_double2, ) SYCL_PSTORE(double, cl::sycl::cl_double2, u) - #undef SYCL_PSTORE +template <> +EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void pstoreu( + Eigen::half* to, const cl::sycl::cl_half8& from) { + auto ptr = cl::sycl::address_space_cast< + cl::sycl::access::address_space::generic_space, + cl::sycl::access::decorated::no>( + reinterpret_cast<cl::sycl::cl_half*>(to)); + from.store(0, ptr); +} + +template <> +EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void pstore( + Eigen::half* to, const cl::sycl::cl_half8& from) { + auto ptr = cl::sycl::address_space_cast< + cl::sycl::access::address_space::generic_space, + cl::sycl::access::decorated::no>( + reinterpret_cast<cl::sycl::cl_half*>(to)); + from.store(0, ptr); +} + #define SYCL_PSET1(packet_type) \ template <> \ EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE packet_type pset1<packet_type>( \ @@ -70,6 +114,7 @@ } // global space +SYCL_PSET1(cl::sycl::cl_half8) SYCL_PSET1(cl::sycl::cl_float4) SYCL_PSET1(cl::sycl::cl_double2) @@ -87,6 +132,58 @@ }; template <> +struct get_base_packet<cl::sycl::cl_half8> { + template <typename sycl_multi_pointer> + static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_half8 get_ploaddup( + sycl_multi_pointer from) { + return cl::sycl::cl_half8(static_cast<cl::sycl::half>(from[0]), + static_cast<cl::sycl::half>(from[0]), + static_cast<cl::sycl::half>(from[1]), + static_cast<cl::sycl::half>(from[1]), + static_cast<cl::sycl::half>(from[2]), + static_cast<cl::sycl::half>(from[2]), + static_cast<cl::sycl::half>(from[3]), + static_cast<cl::sycl::half>(from[3])); + } + template <typename sycl_multi_pointer> + static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_half8 get_pgather( + sycl_multi_pointer from, Index stride) { + return cl::sycl::cl_half8(static_cast<cl::sycl::half>(from[0 * stride]), + static_cast<cl::sycl::half>(from[1 * stride]), + static_cast<cl::sycl::half>(from[2 * stride]), + static_cast<cl::sycl::half>(from[3 * stride]), + static_cast<cl::sycl::half>(from[4 * stride]), + static_cast<cl::sycl::half>(from[5 * stride]), + static_cast<cl::sycl::half>(from[6 * stride]), + static_cast<cl::sycl::half>(from[7 * stride])); + } + + template <typename sycl_multi_pointer> + static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void set_pscatter( + sycl_multi_pointer to, const cl::sycl::cl_half8& from, Index stride) { + auto tmp = stride; + to[0] = Eigen::half(from.s0()); + to[tmp] = Eigen::half(from.s1()); + to[tmp += stride] = Eigen::half(from.s2()); + to[tmp += stride] = Eigen::half(from.s3()); + to[tmp += stride] = Eigen::half(from.s4()); + to[tmp += stride] = Eigen::half(from.s5()); + to[tmp += stride] = Eigen::half(from.s6()); + to[tmp += stride] = Eigen::half(from.s7()); + } + static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_half8 set_plset( + const cl::sycl::half& a) { + return cl::sycl::cl_half8(static_cast<cl::sycl::half>(a), static_cast<cl::sycl::half>(a + 1), + static_cast<cl::sycl::half>(a + 2), + static_cast<cl::sycl::half>(a + 3), + static_cast<cl::sycl::half>(a + 4), + static_cast<cl::sycl::half>(a + 5), + static_cast<cl::sycl::half>(a + 6), + static_cast<cl::sycl::half>(a + 7)); + } +}; + +template <> struct get_base_packet<cl::sycl::cl_float4> { template <typename sycl_multi_pointer> static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_float4 get_ploaddup( @@ -152,6 +249,7 @@ return get_base_packet<packet_type>::get_ploaddup(from); \ } +SYCL_PLOAD_DUP_SPECILIZE(cl::sycl::cl_half8) SYCL_PLOAD_DUP_SPECILIZE(cl::sycl::cl_float4) SYCL_PLOAD_DUP_SPECILIZE(cl::sycl::cl_double2) @@ -165,9 +263,14 @@ } SYCL_PLSET(cl::sycl::cl_float4) SYCL_PLSET(cl::sycl::cl_double2) - #undef SYCL_PLSET +template <> +EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_half8 plset<cl::sycl::cl_half8>( + const typename unpacket_traits<cl::sycl::cl_half8>::type& a) { + return get_base_packet<cl::sycl::cl_half8>::set_plset((const cl::sycl::half &) a); +} + #define SYCL_PGATHER_SPECILIZE(scalar, packet_type) \ template <> \ EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type \ @@ -176,9 +279,9 @@ return get_base_packet<packet_type>::get_pgather(from, stride); \ } +SYCL_PGATHER_SPECILIZE(Eigen::half, cl::sycl::cl_half8) SYCL_PGATHER_SPECILIZE(float, cl::sycl::cl_float4) SYCL_PGATHER_SPECILIZE(double, cl::sycl::cl_double2) - #undef SYCL_PGATHER_SPECILIZE #define SYCL_PSCATTER_SPECILIZE(scalar, packet_type) \ @@ -189,6 +292,7 @@ get_base_packet<packet_type>::set_pscatter(to, from, stride); \ } +SYCL_PSCATTER_SPECILIZE(Eigen::half, cl::sycl::cl_half8) SYCL_PSCATTER_SPECILIZE(float, cl::sycl::cl_float4) SYCL_PSCATTER_SPECILIZE(double, cl::sycl::cl_double2) @@ -201,11 +305,17 @@ return cl::sycl::mad(a, b, c); \ } +SYCL_PMAD(cl::sycl::cl_half8) SYCL_PMAD(cl::sycl::cl_float4) SYCL_PMAD(cl::sycl::cl_double2) #undef SYCL_PMAD template <> +EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half pfirst<cl::sycl::cl_half8>( + const cl::sycl::cl_half8& a) { + return Eigen::half(a.s0()); +} +template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float pfirst<cl::sycl::cl_float4>( const cl::sycl::cl_float4& a) { return a.x(); @@ -217,6 +327,13 @@ } template <> +EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half predux<cl::sycl::cl_half8>( + const cl::sycl::cl_half8& a) { + return Eigen::half(a.s0() + a.s1() + a.s2() + a.s3() + a.s4() + a.s5() + + a.s6() + a.s7()); +} + +template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float predux<cl::sycl::cl_float4>( const cl::sycl::cl_float4& a) { return a.x() + a.y() + a.z() + a.w(); @@ -229,6 +346,17 @@ } template <> +EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half predux_max<cl::sycl::cl_half8>( + const cl::sycl::cl_half8& a) { + return Eigen::half(cl::sycl::fmax( + cl::sycl::fmax( + cl::sycl::fmax(a.s0(), a.s1()), + cl::sycl::fmax(a.s2(), a.s3())), + cl::sycl::fmax( + cl::sycl::fmax(a.s4(), a.s5()), + cl::sycl::fmax(a.s6(), a.s7())))); +} +template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float predux_max<cl::sycl::cl_float4>( const cl::sycl::cl_float4& a) { return cl::sycl::fmax(cl::sycl::fmax(a.x(), a.y()), @@ -241,6 +369,17 @@ } template <> +EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half predux_min<cl::sycl::cl_half8>( + const cl::sycl::cl_half8& a) { + return Eigen::half(cl::sycl::fmin( + cl::sycl::fmin( + cl::sycl::fmin(a.s0(), a.s1()), + cl::sycl::fmin(a.s2(), a.s3())), + cl::sycl::fmin( + cl::sycl::fmin(a.s4(), a.s5()), + cl::sycl::fmin(a.s6(), a.s7())))); +} +template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float predux_min<cl::sycl::cl_float4>( const cl::sycl::cl_float4& a) { return cl::sycl::fmin(cl::sycl::fmin(a.x(), a.y()), @@ -253,6 +392,12 @@ } template <> +EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half predux_mul<cl::sycl::cl_half8>( + const cl::sycl::cl_half8& a) { + return Eigen::half(a.s0() * a.s1() * a.s2() * a.s3() * a.s4() * a.s5() * + a.s6() * a.s7()); +} +template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float predux_mul<cl::sycl::cl_float4>( const cl::sycl::cl_float4& a) { return a.x() * a.y() * a.z() * a.w(); @@ -264,6 +409,14 @@ } template <> +EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_half8 +pabs<cl::sycl::cl_half8>(const cl::sycl::cl_half8& a) { + return cl::sycl::cl_half8(cl::sycl::fabs(a.s0()), cl::sycl::fabs(a.s1()), + cl::sycl::fabs(a.s2()), cl::sycl::fabs(a.s3()), + cl::sycl::fabs(a.s4()), cl::sycl::fabs(a.s5()), + cl::sycl::fabs(a.s6()), cl::sycl::fabs(a.s7())); +} +template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_float4 pabs<cl::sycl::cl_float4>(const cl::sycl::cl_float4& a) { return cl::sycl::cl_float4(cl::sycl::fabs(a.x()), cl::sycl::fabs(a.y()), @@ -300,6 +453,9 @@ return sycl_pcmp_##OP<TYPE>(a, b); \ } +SYCL_PCMP(le, cl::sycl::cl_half8) +SYCL_PCMP(lt, cl::sycl::cl_half8) +SYCL_PCMP(eq, cl::sycl::cl_half8) SYCL_PCMP(le, cl::sycl::cl_float4) SYCL_PCMP(lt, cl::sycl::cl_float4) SYCL_PCMP(eq, cl::sycl::cl_float4) @@ -309,6 +465,121 @@ #undef SYCL_PCMP EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void ptranspose( + PacketBlock<cl::sycl::cl_half8, 8>& kernel) { + cl::sycl::cl_half tmp = kernel.packet[0].s1(); + kernel.packet[0].s1() = kernel.packet[1].s0(); + kernel.packet[1].s0() = tmp; + + tmp = kernel.packet[0].s2(); + kernel.packet[0].s2() = kernel.packet[2].s0(); + kernel.packet[2].s0() = tmp; + + tmp = kernel.packet[0].s3(); + kernel.packet[0].s3() = kernel.packet[3].s0(); + kernel.packet[3].s0() = tmp; + + tmp = kernel.packet[0].s4(); + kernel.packet[0].s4() = kernel.packet[4].s0(); + kernel.packet[4].s0() = tmp; + + tmp = kernel.packet[0].s5(); + kernel.packet[0].s5() = kernel.packet[5].s0(); + kernel.packet[5].s0() = tmp; + + tmp = kernel.packet[0].s6(); + kernel.packet[0].s6() = kernel.packet[6].s0(); + kernel.packet[6].s0() = tmp; + + tmp = kernel.packet[0].s7(); + kernel.packet[0].s7() = kernel.packet[7].s0(); + kernel.packet[7].s0() = tmp; + + tmp = kernel.packet[1].s2(); + kernel.packet[1].s2() = kernel.packet[2].s1(); + kernel.packet[2].s1() = tmp; + + tmp = kernel.packet[1].s3(); + kernel.packet[1].s3() = kernel.packet[3].s1(); + kernel.packet[3].s1() = tmp; + + tmp = kernel.packet[1].s4(); + kernel.packet[1].s4() = kernel.packet[4].s1(); + kernel.packet[4].s1() = tmp; + + tmp = kernel.packet[1].s5(); + kernel.packet[1].s5() = kernel.packet[5].s1(); + kernel.packet[5].s1() = tmp; + + tmp = kernel.packet[1].s6(); + kernel.packet[1].s6() = kernel.packet[6].s1(); + kernel.packet[6].s1() = tmp; + + tmp = kernel.packet[1].s7(); + kernel.packet[1].s7() = kernel.packet[7].s1(); + kernel.packet[7].s1() = tmp; + + tmp = kernel.packet[2].s3(); + kernel.packet[2].s3() = kernel.packet[3].s2(); + kernel.packet[3].s2() = tmp; + + tmp = kernel.packet[2].s4(); + kernel.packet[2].s4() = kernel.packet[4].s2(); + kernel.packet[4].s2() = tmp; + + tmp = kernel.packet[2].s5(); + kernel.packet[2].s5() = kernel.packet[5].s2(); + kernel.packet[5].s2() = tmp; + + tmp = kernel.packet[2].s6(); + kernel.packet[2].s6() = kernel.packet[6].s2(); + kernel.packet[6].s2() = tmp; + + tmp = kernel.packet[2].s7(); + kernel.packet[2].s7() = kernel.packet[7].s2(); + kernel.packet[7].s2() = tmp; + + tmp = kernel.packet[3].s4(); + kernel.packet[3].s4() = kernel.packet[4].s3(); + kernel.packet[4].s3() = tmp; + + tmp = kernel.packet[3].s5(); + kernel.packet[3].s5() = kernel.packet[5].s3(); + kernel.packet[5].s3() = tmp; + + tmp = kernel.packet[3].s6(); + kernel.packet[3].s6() = kernel.packet[6].s3(); + kernel.packet[6].s3() = tmp; + + tmp = kernel.packet[3].s7(); + kernel.packet[3].s7() = kernel.packet[7].s3(); + kernel.packet[7].s3() = tmp; + + tmp = kernel.packet[4].s5(); + kernel.packet[4].s5() = kernel.packet[5].s4(); + kernel.packet[5].s4() = tmp; + + tmp = kernel.packet[4].s6(); + kernel.packet[4].s6() = kernel.packet[6].s4(); + kernel.packet[6].s4() = tmp; + + tmp = kernel.packet[4].s7(); + kernel.packet[4].s7() = kernel.packet[7].s4(); + kernel.packet[7].s4() = tmp; + + tmp = kernel.packet[5].s6(); + kernel.packet[5].s6() = kernel.packet[6].s5(); + kernel.packet[6].s5() = tmp; + + tmp = kernel.packet[5].s7(); + kernel.packet[5].s7() = kernel.packet[7].s5(); + kernel.packet[7].s5() = tmp; + + tmp = kernel.packet[6].s7(); + kernel.packet[6].s7() = kernel.packet[7].s6(); + kernel.packet[7].s6() = tmp; +} + +EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void ptranspose( PacketBlock<cl::sycl::cl_float4, 4>& kernel) { float tmp = kernel.packet[0].y(); kernel.packet[0].y() = kernel.packet[1].x(); @@ -343,6 +614,19 @@ } template <> +EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_half8 pblend( + const Selector<unpacket_traits<cl::sycl::cl_half8>::size>& ifPacket, + const cl::sycl::cl_half8& thenPacket, + const cl::sycl::cl_half8& elsePacket) { + cl::sycl::cl_short8 condition( + ifPacket.select[0] ? 0 : -1, ifPacket.select[1] ? 0 : -1, + ifPacket.select[2] ? 0 : -1, ifPacket.select[3] ? 0 : -1, + ifPacket.select[4] ? 0 : -1, ifPacket.select[5] ? 0 : -1, + ifPacket.select[6] ? 0 : -1, ifPacket.select[7] ? 0 : -1); + return cl::sycl::select(thenPacket, elsePacket, condition); +} + +template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_float4 pblend( const Selector<unpacket_traits<cl::sycl::cl_float4>::size>& ifPacket, const cl::sycl::cl_float4& thenPacket,
diff --git a/Eigen/src/Core/util/XprHelper.h b/Eigen/src/Core/util/XprHelper.h index bb382a7..d05e7d1 100644 --- a/Eigen/src/Core/util/XprHelper.h +++ b/Eigen/src/Core/util/XprHelper.h
@@ -825,7 +825,7 @@ // For regular block expressions, simply forward along the InnerPanel argument, // which is set when calling row/column expressions. - static constexpr bool is_inner_panel(bool inner_panel) { return inner_panel; }; + static constexpr bool is_inner_panel(bool inner_panel) { return inner_panel; } // Only enable non-const base function if XprType is not const (otherwise we get a duplicate definition). template<typename T = XprType, typename EnableIf=std::enable_if_t<!std::is_const<T>::value>>
diff --git a/Eigen/src/Geometry/Transform.h b/Eigen/src/Geometry/Transform.h index 8f64863..1cfff1b 100644 --- a/Eigen/src/Geometry/Transform.h +++ b/Eigen/src/Geometry/Transform.h
@@ -1432,7 +1432,7 @@ typedef Transform<typename Other::Scalar,Dim,Mode,Options> TransformType; typedef typename TransformType::MatrixType MatrixType; typedef Transform<typename Other::Scalar,Dim,Projective,Options> ResultType; - static ResultType run(const Other& other,const TransformType& tr) + static EIGEN_DEVICE_FUNC ResultType run(const Other& other,const TransformType& tr) { return ResultType(other * tr.matrix()); } }; @@ -1443,7 +1443,7 @@ typedef Transform<typename Other::Scalar,Dim,AffineCompact,Options> TransformType; typedef typename TransformType::MatrixType MatrixType; typedef Transform<typename Other::Scalar,Dim,Projective,Options> ResultType; - static ResultType run(const Other& other,const TransformType& tr) + static EIGEN_DEVICE_FUNC ResultType run(const Other& other,const TransformType& tr) { ResultType res; res.matrix().noalias() = other.template block<HDim,Dim>(0,0) * tr.matrix(); @@ -1459,7 +1459,7 @@ typedef Transform<typename Other::Scalar,Dim,Mode,Options> TransformType; typedef typename TransformType::MatrixType MatrixType; typedef TransformType ResultType; - static ResultType run(const Other& other,const TransformType& tr) + static EIGEN_DEVICE_FUNC ResultType run(const Other& other,const TransformType& tr) { ResultType res; res.affine().noalias() = other * tr.matrix(); @@ -1475,7 +1475,7 @@ typedef Transform<typename Other::Scalar,Dim,AffineCompact,Options> TransformType; typedef typename TransformType::MatrixType MatrixType; typedef TransformType ResultType; - static ResultType run(const Other& other,const TransformType& tr) + static EIGEN_DEVICE_FUNC ResultType run(const Other& other,const TransformType& tr) { ResultType res; res.matrix().noalias() = other.template block<Dim,Dim>(0,0) * tr.matrix(); @@ -1491,7 +1491,7 @@ typedef Transform<typename Other::Scalar,Dim,Mode,Options> TransformType; typedef typename TransformType::MatrixType MatrixType; typedef TransformType ResultType; - static ResultType run(const Other& other, const TransformType& tr) + static EIGEN_DEVICE_FUNC ResultType run(const Other& other, const TransformType& tr) { TransformType res; if(Mode!=int(AffineCompact)) @@ -1513,7 +1513,7 @@ typedef Transform<Scalar,Dim,LhsMode,LhsOptions> Lhs; typedef Transform<Scalar,Dim,RhsMode,RhsOptions> Rhs; typedef Transform<Scalar,Dim,ResultMode,LhsOptions> ResultType; - static ResultType run(const Lhs& lhs, const Rhs& rhs) + static EIGEN_DEVICE_FUNC ResultType run(const Lhs& lhs, const Rhs& rhs) { ResultType res; res.linear() = lhs.linear() * rhs.linear(); @@ -1529,7 +1529,7 @@ typedef Transform<Scalar,Dim,LhsMode,LhsOptions> Lhs; typedef Transform<Scalar,Dim,RhsMode,RhsOptions> Rhs; typedef Transform<Scalar,Dim,Projective> ResultType; - static ResultType run(const Lhs& lhs, const Rhs& rhs) + static EIGEN_DEVICE_FUNC ResultType run(const Lhs& lhs, const Rhs& rhs) { return ResultType( lhs.matrix() * rhs.matrix() ); } @@ -1541,7 +1541,7 @@ typedef Transform<Scalar,Dim,AffineCompact,LhsOptions> Lhs; typedef Transform<Scalar,Dim,Projective,RhsOptions> Rhs; typedef Transform<Scalar,Dim,Projective> ResultType; - static ResultType run(const Lhs& lhs, const Rhs& rhs) + static EIGEN_DEVICE_FUNC ResultType run(const Lhs& lhs, const Rhs& rhs) { ResultType res; res.matrix().template topRows<Dim>() = lhs.matrix() * rhs.matrix(); @@ -1556,7 +1556,7 @@ typedef Transform<Scalar,Dim,Projective,LhsOptions> Lhs; typedef Transform<Scalar,Dim,AffineCompact,RhsOptions> Rhs; typedef Transform<Scalar,Dim,Projective> ResultType; - static ResultType run(const Lhs& lhs, const Rhs& rhs) + static EIGEN_DEVICE_FUNC ResultType run(const Lhs& lhs, const Rhs& rhs) { ResultType res(lhs.matrix().template leftCols<Dim>() * rhs.matrix()); res.matrix().col(Dim) += lhs.matrix().col(Dim);
diff --git a/Eigen/src/PardisoSupport/PardisoSupport.h b/Eigen/src/PardisoSupport/PardisoSupport.h index ced7595..f9f7979 100644 --- a/Eigen/src/PardisoSupport/PardisoSupport.h +++ b/Eigen/src/PardisoSupport/PardisoSupport.h
@@ -274,8 +274,8 @@ m_matrix.valuePtr(), m_matrix.outerIndexPtr(), m_matrix.innerIndexPtr(), m_perm.data(), 0, m_iparm.data(), m_msglvl, NULL, NULL); manageErrorCode(error); - m_analysisIsOk = true; - m_factorizationIsOk = true; + m_analysisIsOk = m_info == Eigen::Success; + m_factorizationIsOk = m_info == Eigen::Success; m_isInitialized = true; return derived(); } @@ -296,7 +296,7 @@ m_perm.data(), 0, m_iparm.data(), m_msglvl, NULL, NULL); manageErrorCode(error); - m_analysisIsOk = true; + m_analysisIsOk = m_info == Eigen::Success; m_factorizationIsOk = false; m_isInitialized = true; return derived(); @@ -316,7 +316,7 @@ m_perm.data(), 0, m_iparm.data(), m_msglvl, NULL, NULL); manageErrorCode(error); - m_factorizationIsOk = true; + m_factorizationIsOk = m_info == Eigen::Success; return derived(); }
diff --git a/Eigen/src/SparseCore/SparseTriangularView.h b/Eigen/src/SparseCore/SparseTriangularView.h index 5cac825..2e53114 100644 --- a/Eigen/src/SparseCore/SparseTriangularView.h +++ b/Eigen/src/SparseCore/SparseTriangularView.h
@@ -152,8 +152,8 @@ } } -// inline Index row() const { return (ArgType::Flags&RowMajorBit ? Base::outer() : this->index()); } -// inline Index col() const { return (ArgType::Flags&RowMajorBit ? this->index() : Base::outer()); } + inline Index row() const { return (ArgType::Flags&RowMajorBit ? Base::outer() : this->index()); } + inline Index col() const { return (ArgType::Flags&RowMajorBit ? this->index() : Base::outer()); } inline StorageIndex index() const { if(HasUnitDiag && m_returnOne) return internal::convert_index<StorageIndex>(Base::outer());
diff --git a/debug/msvc/eigen.natvis b/debug/msvc/eigen.natvis index 22cf346..da89857 100644 --- a/debug/msvc/eigen.natvis +++ b/debug/msvc/eigen.natvis
@@ -1,235 +1,235 @@ -<?xml version="1.0" encoding="utf-8"?> - -<AutoVisualizer xmlns="http://schemas.microsoft.com/vstudio/debugger/natvis/2010"> - - <!-- Fixed x Fixed Matrix --> - <Type Name="Eigen::Matrix<*,*,*,*,*,*>"> - <AlternativeType Name="Eigen::Array<*,-1,-1,*,*,*>"/> - <DisplayString>[{$T2}, {$T3}] (fixed matrix)</DisplayString> - <Expand> - <ArrayItems Condition="Flags%2"> <!-- row major layout --> - <Rank>2</Rank> - <Size>$i==0 ? $T2 : $T3</Size> - <ValuePointer>m_storage.m_data.array</ValuePointer> - </ArrayItems> - <ArrayItems Condition="!(Flags%2)"> <!-- column major layout --> - <Direction>Backward</Direction> - <Rank>2</Rank> - <Size>$i==0 ? $T2 : $T3</Size> - <ValuePointer>m_storage.m_data.array</ValuePointer> - </ArrayItems> - </Expand> - </Type> - - <!-- 2 x 2 Matrix --> - <Type Name="Eigen::Matrix<*,2,2,*,*,*>"> - <AlternativeType Name="Eigen::Array<*,2,2,*,*,*>"/> - <DisplayString>[2, 2] (fixed matrix)</DisplayString> - <Expand> - <Synthetic Name="[row 0]" Condition="Flags%2"> - <DisplayString>({m_storage.m_data.array[0]}, {m_storage.m_data.array[1]})</DisplayString> - </Synthetic> - <Synthetic Name="[row 0]" Condition="!(Flags%2)"> - <DisplayString>({m_storage.m_data.array[0]}, {m_storage.m_data.array[2]})</DisplayString> - </Synthetic> - <Synthetic Name="[row 1]" Condition="Flags%2"> - <DisplayString>({m_storage.m_data.array[2]}, {m_storage.m_data.array[3]})</DisplayString> - </Synthetic> - <Synthetic Name="[row 1]" Condition="!(Flags%2)"> - <DisplayString>({m_storage.m_data.array[1]}, {m_storage.m_data.array[3]})</DisplayString> - </Synthetic> - </Expand> - </Type> - - <!-- 3 x 3 Matrix --> - <Type Name="Eigen::Matrix<*,3,3,*,*,*>"> - <AlternativeType Name="Eigen::Array<*,3,3,*,*,*>"/> - <DisplayString>[3, 3] (fixed matrix)</DisplayString> - <Expand> - <Synthetic Name="[row 0]" Condition="Flags%2"> - <DisplayString>({m_storage.m_data.array[0]}, {m_storage.m_data.array[1]}, {m_storage.m_data.array[2]})</DisplayString> - </Synthetic> - <Synthetic Name="[row 0]" Condition="!(Flags%2)"> - <DisplayString>({m_storage.m_data.array[0]}, {m_storage.m_data.array[3]}, {m_storage.m_data.array[6]})</DisplayString> - </Synthetic> - <Synthetic Name="[row 1]" Condition="Flags%2"> - <DisplayString>({m_storage.m_data.array[3]}, {m_storage.m_data.array[4]}, {m_storage.m_data.array[5]})</DisplayString> - </Synthetic> - <Synthetic Name="[row 1]" Condition="!(Flags%2)"> - <DisplayString>({m_storage.m_data.array[1]}, {m_storage.m_data.array[4]}, {m_storage.m_data.array[7]})</DisplayString> - </Synthetic> - <Synthetic Name="[row 2]" Condition="Flags%2"> - <DisplayString>({m_storage.m_data.array[6]}, {m_storage.m_data.array[7]}, {m_storage.m_data.array[8]})</DisplayString> - </Synthetic> - <Synthetic Name="[row 2]" Condition="!(Flags%2)"> - <DisplayString>({m_storage.m_data.array[2]}, {m_storage.m_data.array[5]}, {m_storage.m_data.array[8]})</DisplayString> - </Synthetic> - </Expand> - </Type> - - <!-- 4 x 4 Matrix --> - <Type Name="Eigen::Matrix<*,4,4,*,*,*>"> - <AlternativeType Name="Eigen::Array<*,4,4,*,*,*>"/> - <DisplayString>[4, 4] (fixed matrix)</DisplayString> - <Expand> - <Synthetic Name="[row 0]" Condition="Flags%2"> - <DisplayString>({m_storage.m_data.array[0]}, {m_storage.m_data.array[1]}, {m_storage.m_data.array[2]}, {m_storage.m_data.array[3]})</DisplayString> - </Synthetic> - <Synthetic Name="[row 0]" Condition="!(Flags%2)"> - <DisplayString>({m_storage.m_data.array[0]}, {m_storage.m_data.array[4]}, {m_storage.m_data.array[8]}, {m_storage.m_data.array[12]})</DisplayString> - </Synthetic> - <Synthetic Name="[row 1]" Condition="Flags%2"> - <DisplayString>({m_storage.m_data.array[4]}, {m_storage.m_data.array[5]}, {m_storage.m_data.array[6]}, {m_storage.m_data.array[7]})</DisplayString> - </Synthetic> - <Synthetic Name="[row 1]" Condition="!(Flags%2)"> - <DisplayString>({m_storage.m_data.array[1]}, {m_storage.m_data.array[5]}, {m_storage.m_data.array[9]}, {m_storage.m_data.array[13]})</DisplayString> - </Synthetic> - <Synthetic Name="[row 2]" Condition="Flags%2"> - <DisplayString>({m_storage.m_data.array[8]}, {m_storage.m_data.array[9]}, {m_storage.m_data.array[10]}, {m_storage.m_data.array[11]})</DisplayString> - </Synthetic> - <Synthetic Name="[row 2]" Condition="!(Flags%2)"> - <DisplayString>({m_storage.m_data.array[2]}, {m_storage.m_data.array[6]}, {m_storage.m_data.array[10]}, {m_storage.m_data.array[14]})</DisplayString> - </Synthetic> - <Synthetic Name="[row 3]" Condition="Flags%2"> - <DisplayString>({m_storage.m_data.array[12]}, {m_storage.m_data.array[13]}, {m_storage.m_data.array[14]}, {m_storage.m_data.array[15]})</DisplayString> - </Synthetic> - <Synthetic Name="[row 3]" Condition="!(Flags%2)"> - <DisplayString>({m_storage.m_data.array[3]}, {m_storage.m_data.array[7]}, {m_storage.m_data.array[11]}, {m_storage.m_data.array[15]})</DisplayString> - </Synthetic> - </Expand> - </Type> - - <!-- Dynamic x Dynamic Matrix --> - <Type Name="Eigen::Matrix<*,-1,-1,*,*,*>"> - <AlternativeType Name="Eigen::Array<*,-1,-1,*,*,*>"/> - <DisplayString Condition="m_storage.m_data == 0">empty</DisplayString> - <DisplayString Condition="m_storage.m_data != 0">[{m_storage.m_rows}, {m_storage.m_cols}] (dynamic matrix)</DisplayString> - <Expand> - <ArrayItems Condition="Flags%2"> <!-- row major layout --> - <Rank>2</Rank> - <Size>$i==0 ? m_storage.m_rows : m_storage.m_cols</Size> - <ValuePointer>m_storage.m_data</ValuePointer> - </ArrayItems> - <ArrayItems Condition="!(Flags%2)"> <!-- column major layout --> - <Direction>Backward</Direction> - <Rank>2</Rank> - <Size>$i==0 ? m_storage.m_rows : m_storage.m_cols</Size> - <ValuePointer>m_storage.m_data</ValuePointer> - </ArrayItems> - </Expand> - </Type> - - <!-- Fixed x Dynamic Matrix --> - <Type Name="Eigen::Matrix<*,*,-1,*,*,*>"> - <AlternativeType Name="Eigen::Array<*,*,-1,*,*,*>"/> - <DisplayString Condition="m_storage.m_data == 0">empty</DisplayString> - <DisplayString Condition="m_storage.m_data != 0">[{$T2}, {m_storage.m_cols}] (dynamic column matrix)</DisplayString> - <Expand> - <ArrayItems Condition="Flags%2"> <!-- row major layout --> - <Rank>2</Rank> - <Size>$i==0 ? $T2 : m_storage.m_cols</Size> - <ValuePointer>m_storage.m_data</ValuePointer> - </ArrayItems> - <ArrayItems Condition="!(Flags%2)"> <!-- column major layout --> - <Direction>Backward</Direction> - <Rank>2</Rank> - <Size>$i==0 ? $T2 : m_storage.m_cols</Size> - <ValuePointer>m_storage.m_data</ValuePointer> - </ArrayItems> - </Expand> - </Type> - - <!-- Dynamic x Fixed Matrix --> - <Type Name="Eigen::Matrix<*,-1,*,*,*,*>"> - <AlternativeType Name="Eigen::Array<*,-1,*,*,*,*>"/> - <DisplayString Condition="m_storage.m_data == 0">empty</DisplayString> - <DisplayString Condition="m_storage.m_data != 0">[{m_storage.m_rows}, {$T2}] (dynamic row matrix)</DisplayString> - <Expand> - <ArrayItems Condition="Flags%2"> <!-- row major layout --> - <Rank>2</Rank> - <Size>$i==0 ? m_storage.m_rows : $T2</Size> - <ValuePointer>m_storage.m_data</ValuePointer> - </ArrayItems> - <ArrayItems Condition="!(Flags%2)"> <!-- column major layout --> - <Direction>Backward</Direction> - <Rank>2</Rank> - <Size>$i==0 ? m_storage.m_rows : $T2</Size> - <ValuePointer>m_storage.m_data</ValuePointer> - </ArrayItems> - </Expand> - </Type> - - <!-- Dynamic Column Vector --> - <Type Name="Eigen::Matrix<*,1,-1,*,*,*>"> - <AlternativeType Name="Eigen::Array<*,1,-1,*,*,*>"/> - <DisplayString Condition="m_storage.m_data == 0">empty</DisplayString> - <DisplayString Condition="m_storage.m_data != 0">[{m_storage.m_cols}] (dynamic column vector)</DisplayString> - <Expand> - <Item Name="[size]">m_storage.m_cols</Item> - <ArrayItems> - <Size>m_storage.m_cols</Size> - <ValuePointer>m_storage.m_data</ValuePointer> - </ArrayItems> - </Expand> - </Type> - - <!-- Dynamic Row Vector --> - <Type Name="Eigen::Matrix<*,-1,1,*,*,*>"> - <AlternativeType Name="Eigen::Array<*,-1,1,*,*,*>"/> - <DisplayString Condition="m_storage.m_data == 0">empty</DisplayString> - <DisplayString Condition="m_storage.m_data != 0">[{m_storage.m_rows}] (dynamic row vector)</DisplayString> - <Expand> - <Item Name="[size]">m_storage.m_rows</Item> - <ArrayItems> - <Size>m_storage.m_rows</Size> - <ValuePointer>m_storage.m_data</ValuePointer> - </ArrayItems> - </Expand> - </Type> - - <!-- Fixed Vector --> - <Type Name="Eigen::Matrix<*,1,1,*,*,*>"> - <AlternativeType Name="Eigen::Array<*,1,1,*,*,*>"/> - <DisplayString>[1] ({m_storage.m_data.array[0]})</DisplayString> - <Expand> - <Item Name="[x]">m_storage.m_data.array[0]</Item> - </Expand> - </Type> - - <Type Name="Eigen::Matrix<*,2,1,*,*,*>"> - <AlternativeType Name="Eigen::Matrix<*,1,2,*,*,*>"/> - <AlternativeType Name="Eigen::Array<*,2,1,*,*,*>"/> - <AlternativeType Name="Eigen::Array<*,1,2,*,*,*>"/> - <DisplayString>[2] ({m_storage.m_data.array[0]}, {m_storage.m_data.array[1]})</DisplayString> - <Expand> - <Item Name="[x]">m_storage.m_data.array[0]</Item> - <Item Name="[y]">m_storage.m_data.array[1]</Item> - </Expand> - </Type> - - <Type Name="Eigen::Matrix<*,3,1,*,*,*>"> - <AlternativeType Name="Eigen::Matrix<*,1,3,*,*,*>"/> - <AlternativeType Name="Eigen::Array<*,3,1,*,*,*>"/> - <AlternativeType Name="Eigen::Array<*,1,3,*,*,*>"/> - <DisplayString>[3] ({m_storage.m_data.array[0]}, {m_storage.m_data.array[1]}, {m_storage.m_data.array[2]})</DisplayString> - <Expand> - <Item Name="[x]">m_storage.m_data.array[0]</Item> - <Item Name="[y]">m_storage.m_data.array[1]</Item> - <Item Name="[z]">m_storage.m_data.array[2]</Item> - </Expand> - </Type> - - <Type Name="Eigen::Matrix<*,4,1,*,*,*>"> - <AlternativeType Name="Eigen::Matrix<*,1,4,*,*,*>"/> - <AlternativeType Name="Eigen::Array<*,4,1,*,*,*>"/> - <AlternativeType Name="Eigen::Array<*,1,4,*,*,*>"/> - <DisplayString>[4] ({m_storage.m_data.array[0]}, {m_storage.m_data.array[1]}, {m_storage.m_data.array[2]}, {m_storage.m_data.array[3]})</DisplayString> - <Expand> - <Item Name="[x]">m_storage.m_data.array[0]</Item> - <Item Name="[y]">m_storage.m_data.array[1]</Item> - <Item Name="[z]">m_storage.m_data.array[2]</Item> - <Item Name="[w]">m_storage.m_data.array[3]</Item> - </Expand> - </Type> - -</AutoVisualizer> +<?xml version="1.0" encoding="utf-8"?> + +<AutoVisualizer xmlns="http://schemas.microsoft.com/vstudio/debugger/natvis/2010"> + + <!-- Fixed x Fixed Matrix --> + <Type Name="Eigen::Matrix<*,*,*,*,*,*>"> + <AlternativeType Name="Eigen::Array<*,-1,-1,*,*,*>"/> + <DisplayString>[{$T2}, {$T3}] (fixed matrix)</DisplayString> + <Expand> + <ArrayItems Condition="Flags%2"> <!-- row major layout --> + <Rank>2</Rank> + <Size>$i==0 ? $T2 : $T3</Size> + <ValuePointer>m_storage.m_data.array</ValuePointer> + </ArrayItems> + <ArrayItems Condition="!(Flags%2)"> <!-- column major layout --> + <Direction>Backward</Direction> + <Rank>2</Rank> + <Size>$i==0 ? $T2 : $T3</Size> + <ValuePointer>m_storage.m_data.array</ValuePointer> + </ArrayItems> + </Expand> + </Type> + + <!-- 2 x 2 Matrix --> + <Type Name="Eigen::Matrix<*,2,2,*,*,*>"> + <AlternativeType Name="Eigen::Array<*,2,2,*,*,*>"/> + <DisplayString>[2, 2] (fixed matrix)</DisplayString> + <Expand> + <Synthetic Name="[row 0]" Condition="Flags%2"> + <DisplayString>({m_storage.m_data.array[0]}, {m_storage.m_data.array[1]})</DisplayString> + </Synthetic> + <Synthetic Name="[row 0]" Condition="!(Flags%2)"> + <DisplayString>({m_storage.m_data.array[0]}, {m_storage.m_data.array[2]})</DisplayString> + </Synthetic> + <Synthetic Name="[row 1]" Condition="Flags%2"> + <DisplayString>({m_storage.m_data.array[2]}, {m_storage.m_data.array[3]})</DisplayString> + </Synthetic> + <Synthetic Name="[row 1]" Condition="!(Flags%2)"> + <DisplayString>({m_storage.m_data.array[1]}, {m_storage.m_data.array[3]})</DisplayString> + </Synthetic> + </Expand> + </Type> + + <!-- 3 x 3 Matrix --> + <Type Name="Eigen::Matrix<*,3,3,*,*,*>"> + <AlternativeType Name="Eigen::Array<*,3,3,*,*,*>"/> + <DisplayString>[3, 3] (fixed matrix)</DisplayString> + <Expand> + <Synthetic Name="[row 0]" Condition="Flags%2"> + <DisplayString>({m_storage.m_data.array[0]}, {m_storage.m_data.array[1]}, {m_storage.m_data.array[2]})</DisplayString> + </Synthetic> + <Synthetic Name="[row 0]" Condition="!(Flags%2)"> + <DisplayString>({m_storage.m_data.array[0]}, {m_storage.m_data.array[3]}, {m_storage.m_data.array[6]})</DisplayString> + </Synthetic> + <Synthetic Name="[row 1]" Condition="Flags%2"> + <DisplayString>({m_storage.m_data.array[3]}, {m_storage.m_data.array[4]}, {m_storage.m_data.array[5]})</DisplayString> + </Synthetic> + <Synthetic Name="[row 1]" Condition="!(Flags%2)"> + <DisplayString>({m_storage.m_data.array[1]}, {m_storage.m_data.array[4]}, {m_storage.m_data.array[7]})</DisplayString> + </Synthetic> + <Synthetic Name="[row 2]" Condition="Flags%2"> + <DisplayString>({m_storage.m_data.array[6]}, {m_storage.m_data.array[7]}, {m_storage.m_data.array[8]})</DisplayString> + </Synthetic> + <Synthetic Name="[row 2]" Condition="!(Flags%2)"> + <DisplayString>({m_storage.m_data.array[2]}, {m_storage.m_data.array[5]}, {m_storage.m_data.array[8]})</DisplayString> + </Synthetic> + </Expand> + </Type> + + <!-- 4 x 4 Matrix --> + <Type Name="Eigen::Matrix<*,4,4,*,*,*>"> + <AlternativeType Name="Eigen::Array<*,4,4,*,*,*>"/> + <DisplayString>[4, 4] (fixed matrix)</DisplayString> + <Expand> + <Synthetic Name="[row 0]" Condition="Flags%2"> + <DisplayString>({m_storage.m_data.array[0]}, {m_storage.m_data.array[1]}, {m_storage.m_data.array[2]}, {m_storage.m_data.array[3]})</DisplayString> + </Synthetic> + <Synthetic Name="[row 0]" Condition="!(Flags%2)"> + <DisplayString>({m_storage.m_data.array[0]}, {m_storage.m_data.array[4]}, {m_storage.m_data.array[8]}, {m_storage.m_data.array[12]})</DisplayString> + </Synthetic> + <Synthetic Name="[row 1]" Condition="Flags%2"> + <DisplayString>({m_storage.m_data.array[4]}, {m_storage.m_data.array[5]}, {m_storage.m_data.array[6]}, {m_storage.m_data.array[7]})</DisplayString> + </Synthetic> + <Synthetic Name="[row 1]" Condition="!(Flags%2)"> + <DisplayString>({m_storage.m_data.array[1]}, {m_storage.m_data.array[5]}, {m_storage.m_data.array[9]}, {m_storage.m_data.array[13]})</DisplayString> + </Synthetic> + <Synthetic Name="[row 2]" Condition="Flags%2"> + <DisplayString>({m_storage.m_data.array[8]}, {m_storage.m_data.array[9]}, {m_storage.m_data.array[10]}, {m_storage.m_data.array[11]})</DisplayString> + </Synthetic> + <Synthetic Name="[row 2]" Condition="!(Flags%2)"> + <DisplayString>({m_storage.m_data.array[2]}, {m_storage.m_data.array[6]}, {m_storage.m_data.array[10]}, {m_storage.m_data.array[14]})</DisplayString> + </Synthetic> + <Synthetic Name="[row 3]" Condition="Flags%2"> + <DisplayString>({m_storage.m_data.array[12]}, {m_storage.m_data.array[13]}, {m_storage.m_data.array[14]}, {m_storage.m_data.array[15]})</DisplayString> + </Synthetic> + <Synthetic Name="[row 3]" Condition="!(Flags%2)"> + <DisplayString>({m_storage.m_data.array[3]}, {m_storage.m_data.array[7]}, {m_storage.m_data.array[11]}, {m_storage.m_data.array[15]})</DisplayString> + </Synthetic> + </Expand> + </Type> + + <!-- Dynamic x Dynamic Matrix --> + <Type Name="Eigen::Matrix<*,-1,-1,*,*,*>"> + <AlternativeType Name="Eigen::Array<*,-1,-1,*,*,*>"/> + <DisplayString Condition="m_storage.m_data == 0">empty</DisplayString> + <DisplayString Condition="m_storage.m_data != 0">[{m_storage.m_rows}, {m_storage.m_cols}] (dynamic matrix)</DisplayString> + <Expand> + <ArrayItems Condition="Flags%2"> <!-- row major layout --> + <Rank>2</Rank> + <Size>$i==0 ? m_storage.m_rows : m_storage.m_cols</Size> + <ValuePointer>m_storage.m_data</ValuePointer> + </ArrayItems> + <ArrayItems Condition="!(Flags%2)"> <!-- column major layout --> + <Direction>Backward</Direction> + <Rank>2</Rank> + <Size>$i==0 ? m_storage.m_rows : m_storage.m_cols</Size> + <ValuePointer>m_storage.m_data</ValuePointer> + </ArrayItems> + </Expand> + </Type> + + <!-- Fixed x Dynamic Matrix --> + <Type Name="Eigen::Matrix<*,*,-1,*,*,*>"> + <AlternativeType Name="Eigen::Array<*,*,-1,*,*,*>"/> + <DisplayString Condition="m_storage.m_data == 0">empty</DisplayString> + <DisplayString Condition="m_storage.m_data != 0">[{$T2}, {m_storage.m_cols}] (dynamic column matrix)</DisplayString> + <Expand> + <ArrayItems Condition="Flags%2"> <!-- row major layout --> + <Rank>2</Rank> + <Size>$i==0 ? $T2 : m_storage.m_cols</Size> + <ValuePointer>m_storage.m_data</ValuePointer> + </ArrayItems> + <ArrayItems Condition="!(Flags%2)"> <!-- column major layout --> + <Direction>Backward</Direction> + <Rank>2</Rank> + <Size>$i==0 ? $T2 : m_storage.m_cols</Size> + <ValuePointer>m_storage.m_data</ValuePointer> + </ArrayItems> + </Expand> + </Type> + + <!-- Dynamic x Fixed Matrix --> + <Type Name="Eigen::Matrix<*,-1,*,*,*,*>"> + <AlternativeType Name="Eigen::Array<*,-1,*,*,*,*>"/> + <DisplayString Condition="m_storage.m_data == 0">empty</DisplayString> + <DisplayString Condition="m_storage.m_data != 0">[{m_storage.m_rows}, {$T2}] (dynamic row matrix)</DisplayString> + <Expand> + <ArrayItems Condition="Flags%2"> <!-- row major layout --> + <Rank>2</Rank> + <Size>$i==0 ? m_storage.m_rows : $T2</Size> + <ValuePointer>m_storage.m_data</ValuePointer> + </ArrayItems> + <ArrayItems Condition="!(Flags%2)"> <!-- column major layout --> + <Direction>Backward</Direction> + <Rank>2</Rank> + <Size>$i==0 ? m_storage.m_rows : $T2</Size> + <ValuePointer>m_storage.m_data</ValuePointer> + </ArrayItems> + </Expand> + </Type> + + <!-- Dynamic Column Vector --> + <Type Name="Eigen::Matrix<*,1,-1,*,*,*>"> + <AlternativeType Name="Eigen::Array<*,1,-1,*,*,*>"/> + <DisplayString Condition="m_storage.m_data == 0">empty</DisplayString> + <DisplayString Condition="m_storage.m_data != 0">[{m_storage.m_cols}] (dynamic column vector)</DisplayString> + <Expand> + <Item Name="[size]">m_storage.m_cols</Item> + <ArrayItems> + <Size>m_storage.m_cols</Size> + <ValuePointer>m_storage.m_data</ValuePointer> + </ArrayItems> + </Expand> + </Type> + + <!-- Dynamic Row Vector --> + <Type Name="Eigen::Matrix<*,-1,1,*,*,*>"> + <AlternativeType Name="Eigen::Array<*,-1,1,*,*,*>"/> + <DisplayString Condition="m_storage.m_data == 0">empty</DisplayString> + <DisplayString Condition="m_storage.m_data != 0">[{m_storage.m_rows}] (dynamic row vector)</DisplayString> + <Expand> + <Item Name="[size]">m_storage.m_rows</Item> + <ArrayItems> + <Size>m_storage.m_rows</Size> + <ValuePointer>m_storage.m_data</ValuePointer> + </ArrayItems> + </Expand> + </Type> + + <!-- Fixed Vector --> + <Type Name="Eigen::Matrix<*,1,1,*,*,*>"> + <AlternativeType Name="Eigen::Array<*,1,1,*,*,*>"/> + <DisplayString>[1] ({m_storage.m_data.array[0]})</DisplayString> + <Expand> + <Item Name="[x]">m_storage.m_data.array[0]</Item> + </Expand> + </Type> + + <Type Name="Eigen::Matrix<*,2,1,*,*,*>"> + <AlternativeType Name="Eigen::Matrix<*,1,2,*,*,*>"/> + <AlternativeType Name="Eigen::Array<*,2,1,*,*,*>"/> + <AlternativeType Name="Eigen::Array<*,1,2,*,*,*>"/> + <DisplayString>[2] ({m_storage.m_data.array[0]}, {m_storage.m_data.array[1]})</DisplayString> + <Expand> + <Item Name="[x]">m_storage.m_data.array[0]</Item> + <Item Name="[y]">m_storage.m_data.array[1]</Item> + </Expand> + </Type> + + <Type Name="Eigen::Matrix<*,3,1,*,*,*>"> + <AlternativeType Name="Eigen::Matrix<*,1,3,*,*,*>"/> + <AlternativeType Name="Eigen::Array<*,3,1,*,*,*>"/> + <AlternativeType Name="Eigen::Array<*,1,3,*,*,*>"/> + <DisplayString>[3] ({m_storage.m_data.array[0]}, {m_storage.m_data.array[1]}, {m_storage.m_data.array[2]})</DisplayString> + <Expand> + <Item Name="[x]">m_storage.m_data.array[0]</Item> + <Item Name="[y]">m_storage.m_data.array[1]</Item> + <Item Name="[z]">m_storage.m_data.array[2]</Item> + </Expand> + </Type> + + <Type Name="Eigen::Matrix<*,4,1,*,*,*>"> + <AlternativeType Name="Eigen::Matrix<*,1,4,*,*,*>"/> + <AlternativeType Name="Eigen::Array<*,4,1,*,*,*>"/> + <AlternativeType Name="Eigen::Array<*,1,4,*,*,*>"/> + <DisplayString>[4] ({m_storage.m_data.array[0]}, {m_storage.m_data.array[1]}, {m_storage.m_data.array[2]}, {m_storage.m_data.array[3]})</DisplayString> + <Expand> + <Item Name="[x]">m_storage.m_data.array[0]</Item> + <Item Name="[y]">m_storage.m_data.array[1]</Item> + <Item Name="[z]">m_storage.m_data.array[2]</Item> + <Item Name="[w]">m_storage.m_data.array[3]</Item> + </Expand> + </Type> + +</AutoVisualizer>
diff --git a/debug/msvc/eigen_autoexp_part.dat b/debug/msvc/eigen_autoexp_part.dat index 273c10d..35ef580 100644 --- a/debug/msvc/eigen_autoexp_part.dat +++ b/debug/msvc/eigen_autoexp_part.dat
@@ -1,295 +1,295 @@ -; *************************************************************** -; * Eigen Visualizer -; * -; * Author: Hauke Heibel <hauke.heibel@gmail.com> -; * -; * Support the enhanced debugging of the following Eigen -; * types (*: any, +:fixed dimension) : -; * -; * - Eigen::Matrix<*,4,1,*,*,*> and Eigen::Matrix<*,1,4,*,*,*> -; * - Eigen::Matrix<*,3,1,*,*,*> and Eigen::Matrix<*,1,3,*,*,*> -; * - Eigen::Matrix<*,2,1,*,*,*> and Eigen::Matrix<*,1,2,*,*,*> -; * - Eigen::Matrix<*,-1,-1,*,*,*> -; * - Eigen::Matrix<*,+,-1,*,*,*> -; * - Eigen::Matrix<*,-1,+,*,*,*> -; * - Eigen::Matrix<*,+,+,*,*,*> -; * -; * Matrices are displayed properly independently of the memory -; * alignment (RowMajor vs. ColMajor). -; * -; * This file is distributed WITHOUT ANY WARRANTY. Please ensure -; * that your original autoexp.dat file is copied to a safe -; * place before proceeding with its modification. -; *************************************************************** - -[Visualizer] - -; Fixed size 4-vectors -Eigen::Matrix<*,4,1,*,*,*>|Eigen::Matrix<*,1,4,*,*,*>{ - children - ( - #( - [internals]: [$c,!], - x : ($c.m_storage.m_data.array)[0], - y : ($c.m_storage.m_data.array)[1], - z : ($c.m_storage.m_data.array)[2], - w : ($c.m_storage.m_data.array)[3] - ) - ) - - preview - ( - #( - "[", - 4, - "](", - #array(expr: $e.m_storage.m_data.array[$i], size: 4), - ")" - ) - ) -} - -; Fixed size 3-vectors -Eigen::Matrix<*,3,1,*,*,*>|Eigen::Matrix<*,1,3,*,*,*>{ - children - ( - #( - [internals]: [$c,!], - x : ($c.m_storage.m_data.array)[0], - y : ($c.m_storage.m_data.array)[1], - z : ($c.m_storage.m_data.array)[2] - ) - ) - - preview - ( - #( - "[", - 3, - "](", - #array(expr: $e.m_storage.m_data.array[$i], size: 3), - ")" - ) - ) -} - -; Fixed size 2-vectors -Eigen::Matrix<*,2,1,*,*,*>|Eigen::Matrix<*,1,2,*,*,*>{ - children - ( - #( - [internals]: [$c,!], - x : ($c.m_storage.m_data.array)[0], - y : ($c.m_storage.m_data.array)[1] - ) - ) - - preview - ( - #( - "[", - 2, - "](", - #array(expr: $e.m_storage.m_data.array[$i], size: 2), - ")" - ) - ) -} - -; Fixed size 1-vectors -Eigen::Matrix<*,1,1,*,*,*>|Eigen::Matrix<*,1,1,*,*,*>{ - children - ( - #( - [internals]: [$c,!], - x : ($c.m_storage.m_data.array)[0] - ) - ) - - preview - ( - #( - "[", - 1, - "](", - #array(expr: $e.m_storage.m_data.array[$i], size: 1), - ")" - ) - ) -} - -; Dynamic matrices (ColMajor and RowMajor support) -Eigen::Matrix<*,-1,-1,*,*,*>{ - children - ( - #( - [internals]: [$c,!], - rows: $c.m_storage.m_rows, - cols: $c.m_storage.m_cols, - ; Check for RowMajorBit - #if ($c.Flags & 0x1) ( - #array( - rank: 2, - base: 0, - expr: ($c.m_storage.m_data)[($i % $c.m_storage.m_rows)*$c.m_storage.m_cols + (($i- $i % $c.m_storage.m_rows)/$c.m_storage.m_rows)], - size: ($r==1)*$c.m_storage.m_rows+($r==0)*$c.m_storage.m_cols - ) - ) #else ( - #array( - rank: 2, - base: 0, - expr: ($c.m_storage.m_data)[$i], - size: ($r==1)*$c.m_storage.m_rows+($r==0)*$c.m_storage.m_cols - ) - ) - ) - ) - - preview - ( - #( - "[", - $c.m_storage.m_rows, - ",", - $c.m_storage.m_cols, - "](", - #array( - expr : [($c.m_storage.m_data)[$i],g], - size : $c.m_storage.m_rows*$c.m_storage.m_cols - ), - ")" - ) - ) -} - -; Fixed rows, dynamic columns matrix (ColMajor and RowMajor support) -Eigen::Matrix<*,*,-1,*,*,*>{ - children - ( - #( - [internals]: [$c,!], - rows: $c.RowsAtCompileTime, - cols: $c.m_storage.m_cols, - ; Check for RowMajorBit - #if ($c.Flags & 0x1) ( - #array( - rank: 2, - base: 0, - expr: ($c.m_storage.m_data)[($i % $c.RowsAtCompileTime)*$c.m_storage.m_cols + (($i- $i % $c.RowsAtCompileTime)/$c.RowsAtCompileTime)], - size: ($r==1)*$c.RowsAtCompileTime+($r==0)*$c.m_storage.m_cols - ) - ) #else ( - #array( - rank: 2, - base: 0, - expr: ($c.m_storage.m_data)[$i], - size: ($r==1)*$c.RowsAtCompileTime+($r==0)*$c.m_storage.m_cols - ) - ) - ) - ) - - preview - ( - #( - "[", - $c.RowsAtCompileTime, - ",", - $c.m_storage.m_cols, - "](", - #array( - expr : [($c.m_storage.m_data)[$i],g], - size : $c.RowsAtCompileTime*$c.m_storage.m_cols - ), - ")" - ) - ) -} - -; Dynamic rows, fixed columns matrix (ColMajor and RowMajor support) -Eigen::Matrix<*,-1,*,*,*,*>{ - children - ( - #( - [internals]: [$c,!], - rows: $c.m_storage.m_rows, - cols: $c.ColsAtCompileTime, - ; Check for RowMajorBit - #if ($c.Flags & 0x1) ( - #array( - rank: 2, - base: 0, - expr: ($c.m_storage.m_data)[($i % $c.m_storage.m_rows)*$c.ColsAtCompileTime + (($i- $i % $c.m_storage.m_rows)/$c.m_storage.m_rows)], - size: ($r==1)*$c.m_storage.m_rows+($r==0)*$c.ColsAtCompileTime - ) - ) #else ( - #array( - rank: 2, - base: 0, - expr: ($c.m_storage.m_data)[$i], - size: ($r==1)*$c.m_storage.m_rows+($r==0)*$c.ColsAtCompileTime - ) - ) - ) - ) - - preview - ( - #( - "[", - $c.m_storage.m_rows, - ",", - $c.ColsAtCompileTime, - "](", - #array( - expr : [($c.m_storage.m_data)[$i],g], - size : $c.m_storage.m_rows*$c.ColsAtCompileTime - ), - ")" - ) - ) -} - -; Fixed size matrix (ColMajor and RowMajor support) -Eigen::Matrix<*,*,*,*,*,*>{ - children - ( - #( - [internals]: [$c,!], - rows: $c.RowsAtCompileTime, - cols: $c.ColsAtCompileTime, - ; Check for RowMajorBit - #if ($c.Flags & 0x1) ( - #array( - rank: 2, - base: 0, - expr: ($c.m_storage.m_data.array)[($i % $c.RowsAtCompileTime)*$c.ColsAtCompileTime + (($i- $i % $c.RowsAtCompileTime)/$c.RowsAtCompileTime)], - size: ($r==1)*$c.RowsAtCompileTime+($r==0)*$c.ColsAtCompileTime - ) - ) #else ( - #array( - rank: 2, - base: 0, - expr: ($c.m_storage.m_data.array)[$i], - size: ($r==1)*$c.RowsAtCompileTime+($r==0)*$c.ColsAtCompileTime - ) - ) - ) - ) - - preview - ( - #( - "[", - $c.RowsAtCompileTime, - ",", - $c.ColsAtCompileTime, - "](", - #array( - expr : [($c.m_storage.m_data.array)[$i],g], - size : $c.RowsAtCompileTime*$c.ColsAtCompileTime - ), - ")" - ) - ) -} +; *************************************************************** +; * Eigen Visualizer +; * +; * Author: Hauke Heibel <hauke.heibel@gmail.com> +; * +; * Support the enhanced debugging of the following Eigen +; * types (*: any, +:fixed dimension) : +; * +; * - Eigen::Matrix<*,4,1,*,*,*> and Eigen::Matrix<*,1,4,*,*,*> +; * - Eigen::Matrix<*,3,1,*,*,*> and Eigen::Matrix<*,1,3,*,*,*> +; * - Eigen::Matrix<*,2,1,*,*,*> and Eigen::Matrix<*,1,2,*,*,*> +; * - Eigen::Matrix<*,-1,-1,*,*,*> +; * - Eigen::Matrix<*,+,-1,*,*,*> +; * - Eigen::Matrix<*,-1,+,*,*,*> +; * - Eigen::Matrix<*,+,+,*,*,*> +; * +; * Matrices are displayed properly independently of the memory +; * alignment (RowMajor vs. ColMajor). +; * +; * This file is distributed WITHOUT ANY WARRANTY. Please ensure +; * that your original autoexp.dat file is copied to a safe +; * place before proceeding with its modification. +; *************************************************************** + +[Visualizer] + +; Fixed size 4-vectors +Eigen::Matrix<*,4,1,*,*,*>|Eigen::Matrix<*,1,4,*,*,*>{ + children + ( + #( + [internals]: [$c,!], + x : ($c.m_storage.m_data.array)[0], + y : ($c.m_storage.m_data.array)[1], + z : ($c.m_storage.m_data.array)[2], + w : ($c.m_storage.m_data.array)[3] + ) + ) + + preview + ( + #( + "[", + 4, + "](", + #array(expr: $e.m_storage.m_data.array[$i], size: 4), + ")" + ) + ) +} + +; Fixed size 3-vectors +Eigen::Matrix<*,3,1,*,*,*>|Eigen::Matrix<*,1,3,*,*,*>{ + children + ( + #( + [internals]: [$c,!], + x : ($c.m_storage.m_data.array)[0], + y : ($c.m_storage.m_data.array)[1], + z : ($c.m_storage.m_data.array)[2] + ) + ) + + preview + ( + #( + "[", + 3, + "](", + #array(expr: $e.m_storage.m_data.array[$i], size: 3), + ")" + ) + ) +} + +; Fixed size 2-vectors +Eigen::Matrix<*,2,1,*,*,*>|Eigen::Matrix<*,1,2,*,*,*>{ + children + ( + #( + [internals]: [$c,!], + x : ($c.m_storage.m_data.array)[0], + y : ($c.m_storage.m_data.array)[1] + ) + ) + + preview + ( + #( + "[", + 2, + "](", + #array(expr: $e.m_storage.m_data.array[$i], size: 2), + ")" + ) + ) +} + +; Fixed size 1-vectors +Eigen::Matrix<*,1,1,*,*,*>|Eigen::Matrix<*,1,1,*,*,*>{ + children + ( + #( + [internals]: [$c,!], + x : ($c.m_storage.m_data.array)[0] + ) + ) + + preview + ( + #( + "[", + 1, + "](", + #array(expr: $e.m_storage.m_data.array[$i], size: 1), + ")" + ) + ) +} + +; Dynamic matrices (ColMajor and RowMajor support) +Eigen::Matrix<*,-1,-1,*,*,*>{ + children + ( + #( + [internals]: [$c,!], + rows: $c.m_storage.m_rows, + cols: $c.m_storage.m_cols, + ; Check for RowMajorBit + #if ($c.Flags & 0x1) ( + #array( + rank: 2, + base: 0, + expr: ($c.m_storage.m_data)[($i % $c.m_storage.m_rows)*$c.m_storage.m_cols + (($i- $i % $c.m_storage.m_rows)/$c.m_storage.m_rows)], + size: ($r==1)*$c.m_storage.m_rows+($r==0)*$c.m_storage.m_cols + ) + ) #else ( + #array( + rank: 2, + base: 0, + expr: ($c.m_storage.m_data)[$i], + size: ($r==1)*$c.m_storage.m_rows+($r==0)*$c.m_storage.m_cols + ) + ) + ) + ) + + preview + ( + #( + "[", + $c.m_storage.m_rows, + ",", + $c.m_storage.m_cols, + "](", + #array( + expr : [($c.m_storage.m_data)[$i],g], + size : $c.m_storage.m_rows*$c.m_storage.m_cols + ), + ")" + ) + ) +} + +; Fixed rows, dynamic columns matrix (ColMajor and RowMajor support) +Eigen::Matrix<*,*,-1,*,*,*>{ + children + ( + #( + [internals]: [$c,!], + rows: $c.RowsAtCompileTime, + cols: $c.m_storage.m_cols, + ; Check for RowMajorBit + #if ($c.Flags & 0x1) ( + #array( + rank: 2, + base: 0, + expr: ($c.m_storage.m_data)[($i % $c.RowsAtCompileTime)*$c.m_storage.m_cols + (($i- $i % $c.RowsAtCompileTime)/$c.RowsAtCompileTime)], + size: ($r==1)*$c.RowsAtCompileTime+($r==0)*$c.m_storage.m_cols + ) + ) #else ( + #array( + rank: 2, + base: 0, + expr: ($c.m_storage.m_data)[$i], + size: ($r==1)*$c.RowsAtCompileTime+($r==0)*$c.m_storage.m_cols + ) + ) + ) + ) + + preview + ( + #( + "[", + $c.RowsAtCompileTime, + ",", + $c.m_storage.m_cols, + "](", + #array( + expr : [($c.m_storage.m_data)[$i],g], + size : $c.RowsAtCompileTime*$c.m_storage.m_cols + ), + ")" + ) + ) +} + +; Dynamic rows, fixed columns matrix (ColMajor and RowMajor support) +Eigen::Matrix<*,-1,*,*,*,*>{ + children + ( + #( + [internals]: [$c,!], + rows: $c.m_storage.m_rows, + cols: $c.ColsAtCompileTime, + ; Check for RowMajorBit + #if ($c.Flags & 0x1) ( + #array( + rank: 2, + base: 0, + expr: ($c.m_storage.m_data)[($i % $c.m_storage.m_rows)*$c.ColsAtCompileTime + (($i- $i % $c.m_storage.m_rows)/$c.m_storage.m_rows)], + size: ($r==1)*$c.m_storage.m_rows+($r==0)*$c.ColsAtCompileTime + ) + ) #else ( + #array( + rank: 2, + base: 0, + expr: ($c.m_storage.m_data)[$i], + size: ($r==1)*$c.m_storage.m_rows+($r==0)*$c.ColsAtCompileTime + ) + ) + ) + ) + + preview + ( + #( + "[", + $c.m_storage.m_rows, + ",", + $c.ColsAtCompileTime, + "](", + #array( + expr : [($c.m_storage.m_data)[$i],g], + size : $c.m_storage.m_rows*$c.ColsAtCompileTime + ), + ")" + ) + ) +} + +; Fixed size matrix (ColMajor and RowMajor support) +Eigen::Matrix<*,*,*,*,*,*>{ + children + ( + #( + [internals]: [$c,!], + rows: $c.RowsAtCompileTime, + cols: $c.ColsAtCompileTime, + ; Check for RowMajorBit + #if ($c.Flags & 0x1) ( + #array( + rank: 2, + base: 0, + expr: ($c.m_storage.m_data.array)[($i % $c.RowsAtCompileTime)*$c.ColsAtCompileTime + (($i- $i % $c.RowsAtCompileTime)/$c.RowsAtCompileTime)], + size: ($r==1)*$c.RowsAtCompileTime+($r==0)*$c.ColsAtCompileTime + ) + ) #else ( + #array( + rank: 2, + base: 0, + expr: ($c.m_storage.m_data.array)[$i], + size: ($r==1)*$c.RowsAtCompileTime+($r==0)*$c.ColsAtCompileTime + ) + ) + ) + ) + + preview + ( + #( + "[", + $c.RowsAtCompileTime, + ",", + $c.ColsAtCompileTime, + "](", + #array( + expr : [($c.m_storage.m_data.array)[$i],g], + size : $c.RowsAtCompileTime*$c.ColsAtCompileTime + ), + ")" + ) + ) +}
diff --git a/test/sparse_basic.cpp b/test/sparse_basic.cpp index 9e102b7..298619e 100644 --- a/test/sparse_basic.cpp +++ b/test/sparse_basic.cpp
@@ -746,6 +746,11 @@ // check sparse-triangular to dense refMat3 = m2.template triangularView<StrictlyUpper>(); VERIFY_IS_APPROX(refMat3, DenseMatrix(refMat2.template triangularView<StrictlyUpper>())); + + // check sparse triangular view iteration-based evaluation + m2.setZero(); + VERIFY_IS_CWISE_EQUAL(m2.template triangularView<UnitLower>().toDense(), DenseMatrix::Identity(rows, cols)); + VERIFY_IS_CWISE_EQUAL(m2.template triangularView<UnitUpper>().toDense(), DenseMatrix::Identity(rows, cols)); } // test selfadjointView
diff --git a/unsupported/Eigen/CXX11/ThreadPool b/unsupported/Eigen/CXX11/ThreadPool index 9cdf2c0..d487333 100644 --- a/unsupported/Eigen/CXX11/ThreadPool +++ b/unsupported/Eigen/CXX11/ThreadPool
@@ -1 +1 @@ -#include "../../../Eigen/ThreadPool" \ No newline at end of file +#include "../../../Eigen/ThreadPool" // IWYU pragma: export
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorContractionSycl.h b/unsupported/Eigen/CXX11/src/Tensor/TensorContractionSycl.h index cb74031..480e3ac 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorContractionSycl.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorContractionSycl.h
@@ -1118,7 +1118,7 @@ : globalContractDimOffset + privateOffsetC) : OutScalar(0); - outScalar[j] = cl::sycl::mad(matScalar, vecScalar, outScalar[j]); + outScalar[j] = ::Eigen::internal::pmadd(matScalar, vecScalar, outScalar[j]); privateOffsetNC += Properties::LocalThreadSizeNC; } privateOffsetC += Properties::LocalThreadSizeC; @@ -1263,7 +1263,7 @@ StorageIndex localid = itemID.get_local_id(0); OutScalar accumulator = OutScalar(0); for (StorageIndex i = globalid; i < rng; i += itemID.get_global_range(0)) { - accumulator = cl::sycl::mad(lhs(0, i), rhs(i, 0), accumulator); + accumulator = Eigen::internal::pmadd(lhs(0, i), rhs(i, 0), accumulator); } auto out_scratch_ptr = scratch_ptr + localid; *out_scratch_ptr = accumulator;
diff --git a/unsupported/test/CMakeLists.txt b/unsupported/test/CMakeLists.txt index 1d40ae5..d41baf2 100644 --- a/unsupported/test/CMakeLists.txt +++ b/unsupported/test/CMakeLists.txt
@@ -151,6 +151,7 @@ ei_add_test(cxx11_tensor_argmax_sycl) ei_add_test(cxx11_tensor_custom_op_sycl) ei_add_test(cxx11_tensor_scan_sycl) + ei_add_test(cxx11_tensor_of_float16_sycl) set(EIGEN_SYCL OFF) endif()
diff --git a/unsupported/test/cxx11_tensor_argmax_sycl.cpp b/unsupported/test/cxx11_tensor_argmax_sycl.cpp index 41ea3cf..8f4e095 100644 --- a/unsupported/test/cxx11_tensor_argmax_sycl.cpp +++ b/unsupported/test/cxx11_tensor_argmax_sycl.cpp
@@ -32,9 +32,9 @@ Tensor<DenseIndex, 0, Layout, DenseIndex> out_max; Tensor<DenseIndex, 0, Layout, DenseIndex> out_min; in.setRandom(); - in *= in.constant(100.0); - in(0, 0, 0) = -1000.0; - in(1, 1, 1) = 1000.0; + in *= in.constant(static_cast<DataType>(100.0)); + in(0, 0, 0) = static_cast<DataType>(-1000.0); + in(1, 1, 1) = static_cast<DataType>(1000.0); std::size_t in_bytes = in.size() * sizeof(DataType); std::size_t out_bytes = out_max.size() * sizeof(DenseIndex); @@ -93,7 +93,7 @@ ix[3] = l; // suppose dim == 1, then for all i, k, l, set tensor(i, 0, k, l) // = 10.0 - tensor(ix) = (ix[dim] != 0) ? -1.0 : 10.0; + tensor(ix) = static_cast<DataType>((ix[dim] != 0) ? -1.0 : 10.0); } } } @@ -132,7 +132,7 @@ ix[2] = k; ix[3] = l; // suppose dim == 1, then for all i, k, l, set tensor(i, 2, k, l) = 20.0 - tensor(ix) = (ix[dim] != tensor.dimension(dim) - 1) ? -1.0 : 20.0; + tensor(ix) = static_cast<DataType>((ix[dim] != tensor.dimension(dim) - 1) ? -1.0 : 20.0); } } } @@ -180,7 +180,7 @@ ix[2] = k; ix[3] = l; // suppose dim == 1, then for all i, k, l, set tensor(i, 0, k, l) = -10.0 - tensor(ix) = (ix[dim] != 0) ? 1.0 : -10.0; + tensor(ix) = static_cast<DataType>((ix[dim] != 0) ? 1.0 : -10.0); } } } @@ -219,7 +219,7 @@ ix[2] = k; ix[3] = l; // suppose dim == 1, then for all i, k, l, set tensor(i, 2, k, l) = -20.0 - tensor(ix) = (ix[dim] != tensor.dimension(dim) - 1) ? 1.0 : -20.0; + tensor(ix) = static_cast<DataType>((ix[dim] != tensor.dimension(dim) - 1) ? 1.0 : -20.0); } } } @@ -252,6 +252,7 @@ EIGEN_DECLARE_TEST(cxx11_tensor_argmax_sycl) { for (const auto& device : Eigen::get_sycl_supported_devices()) { + CALL_SUBTEST(sycl_argmax_test_per_device<half>(device)); CALL_SUBTEST(sycl_argmax_test_per_device<float>(device)); } }
diff --git a/unsupported/test/cxx11_tensor_broadcast_sycl.cpp b/unsupported/test/cxx11_tensor_broadcast_sycl.cpp index 20f84b8..6ca5ac7 100644 --- a/unsupported/test/cxx11_tensor_broadcast_sycl.cpp +++ b/unsupported/test/cxx11_tensor_broadcast_sycl.cpp
@@ -139,6 +139,7 @@ EIGEN_DECLARE_TEST(cxx11_tensor_broadcast_sycl) { for (const auto& device :Eigen::get_sycl_supported_devices()) { + CALL_SUBTEST(sycl_broadcast_test_per_device<half>(device)); CALL_SUBTEST(sycl_broadcast_test_per_device<float>(device)); } }
diff --git a/unsupported/test/cxx11_tensor_builtins_sycl.cpp b/unsupported/test/cxx11_tensor_builtins_sycl.cpp index be7afd2..c8668d8 100644 --- a/unsupported/test/cxx11_tensor_builtins_sycl.cpp +++ b/unsupported/test/cxx11_tensor_builtins_sycl.cpp
@@ -28,26 +28,135 @@ // Functions used to compare the TensorMap implementation on the device with // the equivalent on the host namespace SYCL { -template <typename T> T abs(T x) { + +template <typename T> T abs(T x) { return cl::sycl::abs(x); } +template <> Eigen::half abs(Eigen::half x) { + return Eigen::half(cl::sycl::fabs(static_cast<cl::sycl::half>(x))); +} -template <> float abs(float x) { +template <> float abs(float x) { return cl::sycl::fabs(x); } -template <> double abs(double x) { +template <> double abs(double x) { return cl::sycl::fabs(x); } template <typename T> T square(T x) { return x * x; } template <typename T> T cube(T x) { return x * x * x; } template <typename T> T inverse(T x) { return T(1) / x; } -template <typename T> T cwiseMax(T x, T y) { -return cl::sycl::max(x, y); +template <typename T> T cwiseMax(T x, T y) { + return cl::sycl::max(x, y); } -template <typename T> T cwiseMin(T x, T y) { - return cl::sycl::min(x, y); +template <> Eigen::half cwiseMax(Eigen::half x, Eigen::half y) { +return Eigen::half(cl::sycl::max(static_cast<cl::sycl::half>(x), static_cast<cl::sycl::half>(y))); +} + +template <typename T> T cwiseMin(T x, T y) { + return cl::sycl::min(x, y); +} +template <> Eigen::half cwiseMin(Eigen::half x, Eigen::half y) { + return Eigen::half(cl::sycl::min(static_cast<cl::sycl::half>(x), static_cast<cl::sycl::half>(y))); +} + +template <typename T> T sqrt(T x) { + return cl::sycl::sqrt(x); +} +template <> Eigen::half sqrt(Eigen::half x) { + return Eigen::half(cl::sycl::sqrt(static_cast<cl::sycl::half>(x))); +} + +template <typename T> T rsqrt(T x) { + return cl::sycl::rsqrt(x); +} +template <> Eigen::half rsqrt(Eigen::half x) { + return Eigen::half(cl::sycl::rsqrt(static_cast<cl::sycl::half>(x))); +} + +template <typename T> T tanh(T x) { + return cl::sycl::tanh(x); +} +template <> Eigen::half tanh(Eigen::half x) { + return Eigen::half(cl::sycl::tanh(static_cast<cl::sycl::half>(x))); +} + +template <typename T> T exp(T x) { + return cl::sycl::exp(x); +} +template <> Eigen::half exp(Eigen::half x) { + return Eigen::half(cl::sycl::exp(static_cast<cl::sycl::half>(x))); +} + +template <typename T> T expm1(T x) { + return cl::sycl::expm1(x); +} +template <> Eigen::half expm1(Eigen::half x) { + return Eigen::half(cl::sycl::expm1(static_cast<cl::sycl::half>(x))); +} + +template <typename T> T log(T x) { + return cl::sycl::log(x); +} +template <> Eigen::half log(Eigen::half x) { + return Eigen::half(cl::sycl::log(static_cast<cl::sycl::half>(x))); +} + +template <typename T> T ceil(T x) { + return cl::sycl::ceil(x); +} +template <> Eigen::half ceil(Eigen::half x) { + return Eigen::half(cl::sycl::ceil(static_cast<cl::sycl::half>(x))); +} + +template <typename T> T floor(T x) { + return cl::sycl::floor(x); +} +template <> Eigen::half floor(Eigen::half x) { + return Eigen::half(cl::sycl::floor(static_cast<cl::sycl::half>(x))); +} + +template <typename T> T round(T x) { + return cl::sycl::round(x); +} +template <> Eigen::half round(Eigen::half x) { + return Eigen::half(cl::sycl::round(static_cast<cl::sycl::half>(x))); +} + +template <typename T> T log1p(T x) { + return cl::sycl::log1p(x); +} +template <> Eigen::half log1p(Eigen::half x) { + return Eigen::half(cl::sycl::log1p(static_cast<cl::sycl::half>(x))); +} + +template <typename T> T sign(T x) { + return cl::sycl::sign(x); +} +template <> Eigen::half sign(Eigen::half x) { + return Eigen::half(cl::sycl::sign(static_cast<cl::sycl::half>(x))); +} + +template <typename T> T isnan(T x) { + return cl::sycl::isnan(x); +} +template <> Eigen::half isnan(Eigen::half x) { + return Eigen::half(cl::sycl::isnan(static_cast<cl::sycl::half>(x))); +} + +template <typename T> T isfinite(T x) { + return cl::sycl::isfinite(x); +} +template <> Eigen::half isfinite(Eigen::half x) { + return Eigen::half(cl::sycl::isfinite(static_cast<cl::sycl::half>(x))); +} + +template <typename T> T isinf(T x) { + return cl::sycl::isinf(x); +} +template <> Eigen::half isinf(Eigen::half x) { + return Eigen::half(cl::sycl::isinf(static_cast<cl::sycl::half>(x))); } } @@ -157,15 +266,14 @@ #define DECLARE_UNARY_STRUCT(FUNC) \ struct op_##FUNC { \ template <typename T> \ - auto operator()(const T& x) -> decltype(cl::sycl::FUNC(x)) { \ - return cl::sycl::FUNC(x); \ + auto operator()(const T& x) -> decltype(SYCL::FUNC(x)) { \ + return SYCL::FUNC(x); \ } \ template <typename T> \ auto operator()(const TensorMap<T>& x) -> decltype(x.FUNC()) { \ return x.FUNC(); \ } \ - }; - +}; DECLARE_UNARY_STRUCT(sqrt) DECLARE_UNARY_STRUCT(rsqrt) @@ -390,8 +498,10 @@ for (const auto& device :Eigen::get_sycl_supported_devices()) { QueueInterface queueInterface(device); Eigen::SyclDevice sycl_device(&queueInterface); - CALL_SUBTEST_1(test_builtin_unary_sycl<float>(sycl_device)); - CALL_SUBTEST_2(test_floating_builtin_binary_sycl<float>(sycl_device)); - CALL_SUBTEST_3(test_integer_builtin_binary_sycl<int>(sycl_device)); + CALL_SUBTEST_1(test_builtin_unary_sycl<half>(sycl_device)); + CALL_SUBTEST_2(test_floating_builtin_binary_sycl<half>(sycl_device)); + CALL_SUBTEST_3(test_builtin_unary_sycl<float>(sycl_device)); + CALL_SUBTEST_4(test_floating_builtin_binary_sycl<float>(sycl_device)); + CALL_SUBTEST_5(test_integer_builtin_binary_sycl<int>(sycl_device)); } }
diff --git a/unsupported/test/cxx11_tensor_chipping_sycl.cpp b/unsupported/test/cxx11_tensor_chipping_sycl.cpp index 1e70931..b018da8 100644 --- a/unsupported/test/cxx11_tensor_chipping_sycl.cpp +++ b/unsupported/test/cxx11_tensor_chipping_sycl.cpp
@@ -619,5 +619,6 @@ { for (const auto& device :Eigen::get_sycl_supported_devices()) { CALL_SUBTEST(sycl_chipping_test_per_device<float>(device)); + CALL_SUBTEST(sycl_chipping_test_per_device<half>(device)); } }
diff --git a/unsupported/test/cxx11_tensor_concatenation_sycl.cpp b/unsupported/test/cxx11_tensor_concatenation_sycl.cpp index 765991b..18dc95b 100644 --- a/unsupported/test/cxx11_tensor_concatenation_sycl.cpp +++ b/unsupported/test/cxx11_tensor_concatenation_sycl.cpp
@@ -175,6 +175,7 @@ } EIGEN_DECLARE_TEST(cxx11_tensor_concatenation_sycl) { for (const auto& device :Eigen::get_sycl_supported_devices()) { + CALL_SUBTEST(tensorConcat_perDevice<half>(device)); CALL_SUBTEST(tensorConcat_perDevice<float>(device)); } }
diff --git a/unsupported/test/cxx11_tensor_device_sycl.cpp b/unsupported/test/cxx11_tensor_device_sycl.cpp index 2f75ef8..66a13e8 100644 --- a/unsupported/test/cxx11_tensor_device_sycl.cpp +++ b/unsupported/test/cxx11_tensor_device_sycl.cpp
@@ -82,6 +82,7 @@ EIGEN_DECLARE_TEST(cxx11_tensor_device_sycl) { for (const auto& device :Eigen::get_sycl_supported_devices()) { + CALL_SUBTEST(sycl_device_test_per_device<half>(device)); CALL_SUBTEST(sycl_device_test_per_device<float>(device)); CALL_SUBTEST(sycl_device_test_per_device<OffByOneScalar<int>>(device)); }
diff --git a/unsupported/test/cxx11_tensor_image_op_sycl.cpp b/unsupported/test/cxx11_tensor_image_op_sycl.cpp index db1c020..083c004 100644 --- a/unsupported/test/cxx11_tensor_image_op_sycl.cpp +++ b/unsupported/test/cxx11_tensor_image_op_sycl.cpp
@@ -95,6 +95,7 @@ EIGEN_DECLARE_TEST(cxx11_tensor_image_op_sycl) { for (const auto& device :Eigen::get_sycl_supported_devices()) { + CALL_SUBTEST(sycl_computing_test_per_device<half>(device)); CALL_SUBTEST(sycl_computing_test_per_device<float>(device)); #ifdef EIGEN_SYCL_DOUBLE_SUPPORT CALL_SUBTEST(sycl_computing_test_per_device<double>(device));
diff --git a/unsupported/test/cxx11_tensor_inflation_sycl.cpp b/unsupported/test/cxx11_tensor_inflation_sycl.cpp index 521ae0c..75c2c0e 100644 --- a/unsupported/test/cxx11_tensor_inflation_sycl.cpp +++ b/unsupported/test/cxx11_tensor_inflation_sycl.cpp
@@ -131,6 +131,7 @@ EIGEN_DECLARE_TEST(cxx11_tensor_inflation_sycl) { for (const auto& device :Eigen::get_sycl_supported_devices()) { + CALL_SUBTEST(sycl_inflation_test_per_device<half>(device)); CALL_SUBTEST(sycl_inflation_test_per_device<float>(device)); } }
diff --git a/unsupported/test/cxx11_tensor_layout_swap_sycl.cpp b/unsupported/test/cxx11_tensor_layout_swap_sycl.cpp index 9546b91..f556d84 100644 --- a/unsupported/test/cxx11_tensor_layout_swap_sycl.cpp +++ b/unsupported/test/cxx11_tensor_layout_swap_sycl.cpp
@@ -121,6 +121,7 @@ EIGEN_DECLARE_TEST(cxx11_tensor_layout_swap_sycl) { for (const auto& device :Eigen::get_sycl_supported_devices()) { + CALL_SUBTEST(sycl_tensor_layout_swap_test_per_device<half>(device)); CALL_SUBTEST(sycl_tensor_layout_swap_test_per_device<float>(device)); } }
diff --git a/unsupported/test/cxx11_tensor_math_sycl.cpp b/unsupported/test/cxx11_tensor_math_sycl.cpp index 029653e..84d638e 100644 --- a/unsupported/test/cxx11_tensor_math_sycl.cpp +++ b/unsupported/test/cxx11_tensor_math_sycl.cpp
@@ -100,6 +100,7 @@ EIGEN_DECLARE_TEST(cxx11_tensor_math_sycl) { for (const auto& device :Eigen::get_sycl_supported_devices()) { + CALL_SUBTEST(sycl_computing_test_per_device<half>(device)); CALL_SUBTEST(sycl_computing_test_per_device<float>(device)); } }
diff --git a/unsupported/test/cxx11_tensor_morphing_sycl.cpp b/unsupported/test/cxx11_tensor_morphing_sycl.cpp index bf001b4..a1545d4 100644 --- a/unsupported/test/cxx11_tensor_morphing_sycl.cpp +++ b/unsupported/test/cxx11_tensor_morphing_sycl.cpp
@@ -381,6 +381,7 @@ EIGEN_DECLARE_TEST(cxx11_tensor_morphing_sycl) { for (const auto& device :Eigen::get_sycl_supported_devices()) { + CALL_SUBTEST(sycl_morphing_test_per_device<half>(device)); CALL_SUBTEST(sycl_morphing_test_per_device<float>(device)); } }
diff --git a/unsupported/test/cxx11_tensor_of_float16_sycl.cpp b/unsupported/test/cxx11_tensor_of_float16_sycl.cpp new file mode 100644 index 0000000..1624acc --- /dev/null +++ b/unsupported/test/cxx11_tensor_of_float16_sycl.cpp
@@ -0,0 +1,406 @@ +// This file is part of Eigen, a lightweight C++ template library +// for linear algebra. +// +// Copyright (C) 2023 +// Alejandro Acosta Codeplay Software Ltd. +// Contact: <eigen@codeplay.com> +// Benoit Steiner <benoit.steiner.goog@gmail.com> +// +// This Source Code Form is subject to the terms of the Mozilla +// Public License v. 2.0. If a copy of the MPL was not distributed +// with this file, You can obtain one at http://mozilla.org/MPL/2.0/. + +#define EIGEN_TEST_NO_LONGDOUBLE +#define EIGEN_TEST_NO_COMPLEX + +#define EIGEN_DEFAULT_DENSE_INDEX_TYPE int +#define EIGEN_USE_SYCL +#define EIGEN_SYCL_HALF_SUPPORT + +#include "main.h" +#include <unsupported/Eigen/CXX11/Tensor> + +using Eigen::Tensor; +using Eigen::SyclDevice; + +void test_gpu_numext(const Eigen::SyclDevice &sycl_device) { + int num_elem = 101; + + float* d_float = static_cast<float*>(sycl_device.allocate(num_elem * sizeof(float))); + bool* d_res_half = static_cast<bool*>(sycl_device.allocate(num_elem * sizeof(bool))); + bool* d_res_float = static_cast<bool*>(sycl_device.allocate(num_elem * sizeof(bool))); + + Eigen::TensorMap<Tensor<float, 1>, Eigen::Aligned> gpu_float(d_float, num_elem); + Eigen::TensorMap<Tensor<bool, 1>, Eigen::Aligned> gpu_res_half(d_res_half, num_elem); + Eigen::TensorMap<Tensor<bool, 1>, Eigen::Aligned> gpu_res_float(d_res_float, num_elem); + + gpu_float.device(sycl_device) = gpu_float.random() - gpu_float.constant(0.5f); + gpu_res_float.device(sycl_device) = gpu_float.unaryExpr(Eigen::internal::scalar_isnan_op<float>()); + gpu_res_half.device(sycl_device) = gpu_float.cast<Eigen::half>().unaryExpr(Eigen::internal::scalar_isnan_op<Eigen::half>()); + + Tensor<bool, 1> half_prec(num_elem); + Tensor<bool, 1> full_prec(num_elem); + + sycl_device.memcpyDeviceToHost(half_prec.data(), d_res_half,num_elem * sizeof(bool)); + sycl_device.memcpyDeviceToHost(full_prec.data(), d_res_float,num_elem * sizeof(bool)); + + for (int i = 0; i < num_elem; ++i) { + std::cout << "Checking numext " << i << std::endl; + VERIFY_IS_EQUAL(full_prec(i), half_prec(i)); + } +} + +void test_gpu_conversion(const Eigen::SyclDevice &sycl_device) { + int num_elem = 101; + + float* d_float = static_cast<float*>(sycl_device.allocate(num_elem * sizeof(float))); + Eigen::half* d_half = static_cast<Eigen::half*>(sycl_device.allocate(num_elem * sizeof(Eigen::half))); + float* d_conv = static_cast<float*>(sycl_device.allocate(num_elem * sizeof(float))); + + Eigen::TensorMap<Eigen::Tensor<float, 1>, Eigen::Aligned> gpu_float( + d_float, num_elem); + Eigen::TensorMap<Eigen::Tensor<Eigen::half, 1>, Eigen::Aligned> gpu_half( + d_half, num_elem); + Eigen::TensorMap<Eigen::Tensor<float, 1>, Eigen::Aligned> gpu_conv( + d_conv, num_elem); + + gpu_float.device(sycl_device) = gpu_float.random(); + gpu_half.device(sycl_device) = gpu_float.cast<Eigen::half>(); + gpu_conv.device(sycl_device) = gpu_half.cast<float>(); + + Tensor<float, 1> initial(num_elem); + Tensor<float, 1> final(num_elem); + sycl_device.memcpyDeviceToHost(initial.data(), d_float, num_elem*sizeof(float)); + sycl_device.memcpyDeviceToHost(final.data(), d_conv, num_elem*sizeof(float)); + + for (int i = 0; i < num_elem; ++i) { + VERIFY_IS_APPROX(initial(i), final(i)); + } +} + +void test_gpu_unary(const Eigen::SyclDevice &sycl_device) { + int num_elem = 101; + + float* d_float = (float*)sycl_device.allocate(num_elem * sizeof(float)); + float* d_res_half = (float*)sycl_device.allocate(num_elem * sizeof(float)); + float* d_res_float = (float*)sycl_device.allocate(num_elem * sizeof(float)); + + Eigen::TensorMap<Eigen::Tensor<float, 1>, Eigen::Aligned> gpu_float( + d_float, num_elem); + Eigen::TensorMap<Eigen::Tensor<float, 1>, Eigen::Aligned> gpu_res_half( + d_res_half, num_elem); + Eigen::TensorMap<Eigen::Tensor<float, 1>, Eigen::Aligned> gpu_res_float( + d_res_float, num_elem); + + gpu_float.device(sycl_device) = gpu_float.random() - gpu_float.constant(0.5f); + gpu_res_float.device(sycl_device) = gpu_float.abs(); + gpu_res_half.device(sycl_device) = gpu_float.cast<Eigen::half>().abs().cast<float>(); + + Tensor<float, 1> half_prec(num_elem); + Tensor<float, 1> full_prec(num_elem); + sycl_device.memcpyDeviceToHost(half_prec.data(), d_res_half, num_elem*sizeof(float)); + sycl_device.memcpyDeviceToHost(full_prec.data(), d_res_float, num_elem*sizeof(float)); + sycl_device.synchronize(); + + for (int i = 0; i < num_elem; ++i) { + std::cout << "Checking unary " << i << std::endl; + VERIFY_IS_APPROX(full_prec(i), half_prec(i)); + } +} + +void test_gpu_elementwise(const Eigen::SyclDevice &sycl_device) { + int num_elem = 101; + + float* d_float1 = static_cast<float*>(sycl_device.allocate(num_elem * sizeof(float))); + float* d_float2 = static_cast<float*>(sycl_device.allocate(num_elem * sizeof(float))); + float* d_res_half = static_cast<float*>(sycl_device.allocate(num_elem * sizeof(float))); + float* d_res_float = static_cast<float*>(sycl_device.allocate(num_elem * sizeof(float))); + + Eigen::TensorMap<Tensor<float, 1>, Eigen::Aligned> gpu_float1(d_float1, num_elem); + Eigen::TensorMap<Tensor<float, 1>, Eigen::Aligned> gpu_float2(d_float2, num_elem); + Eigen::TensorMap<Tensor<float, 1>, Eigen::Aligned> gpu_res_half(d_res_half, num_elem); + Eigen::TensorMap<Tensor<float, 1>, Eigen::Aligned> gpu_res_float(d_res_float, num_elem); + + gpu_float1.device(sycl_device) = gpu_float1.random(); + gpu_float2.device(sycl_device) = gpu_float2.random(); + gpu_res_float.device(sycl_device) = (gpu_float1 + gpu_float2) * gpu_float1; + gpu_res_half.device(sycl_device) = ((gpu_float1.cast<Eigen::half>() + gpu_float2.cast<Eigen::half>()) * gpu_float1.cast<Eigen::half>()).cast<float>(); + + Tensor<float, 1> half_prec(num_elem); + Tensor<float, 1> full_prec(num_elem); + + sycl_device.memcpyDeviceToHost(half_prec.data(), d_res_half,num_elem * sizeof(float)); + sycl_device.memcpyDeviceToHost(full_prec.data(), d_res_float,num_elem * sizeof(float)); + + for (int i = 0; i < num_elem; ++i) { + std::cout << "Checking elemwise " << i << ": full prec = " << full_prec(i) << " vs half prec = " << half_prec(i) << std::endl; + VERIFY_IS_APPROX(static_cast<Eigen::half>(full_prec(i)), static_cast<Eigen::half>(half_prec(i))); + } +} + +void test_gpu_trancendental(const Eigen::SyclDevice &sycl_device) { + int num_elem = 101; + + float* d_float1 = (float*)sycl_device.allocate(num_elem * sizeof(float)); + float* d_float2 = (float*)sycl_device.allocate(num_elem * sizeof(float)); + float* d_float3 = (float*)sycl_device.allocate(num_elem * sizeof(float)); + Eigen::half* d_res1_half = (Eigen::half*)sycl_device.allocate(num_elem * sizeof(Eigen::half)); + Eigen::half* d_res1_float = (Eigen::half*)sycl_device.allocate(num_elem * sizeof(Eigen::half)); + Eigen::half* d_res2_half = (Eigen::half*)sycl_device.allocate(num_elem * sizeof(Eigen::half)); + Eigen::half* d_res2_float = (Eigen::half*)sycl_device.allocate(num_elem * sizeof(Eigen::half)); + Eigen::half* d_res3_half = (Eigen::half*)sycl_device.allocate(num_elem * sizeof(Eigen::half)); + Eigen::half* d_res3_float = (Eigen::half*)sycl_device.allocate(num_elem * sizeof(Eigen::half)); + + Eigen::TensorMap<Eigen::Tensor<float, 1>, Eigen::Aligned> gpu_float1(d_float1, num_elem); + Eigen::TensorMap<Eigen::Tensor<float, 1>, Eigen::Aligned> gpu_float2(d_float2, num_elem); + Eigen::TensorMap<Eigen::Tensor<float, 1>, Eigen::Aligned> gpu_float3(d_float3, num_elem); + Eigen::TensorMap<Eigen::Tensor<Eigen::half, 1>, Eigen::Aligned> gpu_res1_half(d_res1_half, num_elem); + Eigen::TensorMap<Eigen::Tensor<Eigen::half, 1>, Eigen::Aligned> gpu_res1_float(d_res1_float, num_elem); + Eigen::TensorMap<Eigen::Tensor<Eigen::half, 1>, Eigen::Aligned> gpu_res2_half(d_res2_half, num_elem); + Eigen::TensorMap<Eigen::Tensor<Eigen::half, 1>, Eigen::Aligned> gpu_res2_float(d_res2_float, num_elem); + Eigen::TensorMap<Eigen::Tensor<Eigen::half, 1>, Eigen::Aligned> gpu_res3_half(d_res3_half, num_elem); + Eigen::TensorMap<Eigen::Tensor<Eigen::half, 1>, Eigen::Aligned> gpu_res3_float(d_res3_float, num_elem); + Eigen::TensorMap<Eigen::Tensor<Eigen::half, 1>, Eigen::Aligned> gpu_res4_half(d_res3_half, num_elem); + Eigen::TensorMap<Eigen::Tensor<Eigen::half, 1>, Eigen::Aligned> gpu_res4_float(d_res3_float, num_elem); + + gpu_float1.device(sycl_device) = gpu_float1.random() - gpu_float1.constant(0.5f); + gpu_float2.device(sycl_device) = gpu_float2.random() + gpu_float1.constant(0.5f); + gpu_float3.device(sycl_device) = gpu_float3.random(); + gpu_res1_float.device(sycl_device) = gpu_float1.exp().cast<Eigen::half>(); + gpu_res2_float.device(sycl_device) = gpu_float2.log().cast<Eigen::half>(); + gpu_res3_float.device(sycl_device) = gpu_float3.log1p().cast<Eigen::half>(); + gpu_res4_float.device(sycl_device) = gpu_float3.expm1().cast<Eigen::half>(); + + gpu_res1_half.device(sycl_device) = gpu_float1.cast<Eigen::half>(); + gpu_res1_half.device(sycl_device) = gpu_res1_half.exp(); + + gpu_res2_half.device(sycl_device) = gpu_float2.cast<Eigen::half>(); + gpu_res2_half.device(sycl_device) = gpu_res2_half.log(); + + gpu_res3_half.device(sycl_device) = gpu_float3.cast<Eigen::half>(); + gpu_res3_half.device(sycl_device) = gpu_res3_half.log1p(); + + gpu_res3_half.device(sycl_device) = gpu_float3.cast<Eigen::half>(); + gpu_res3_half.device(sycl_device) = gpu_res3_half.expm1(); + + Tensor<float, 1> input1(num_elem); + Tensor<Eigen::half, 1> half_prec1(num_elem); + Tensor<Eigen::half, 1> full_prec1(num_elem); + Tensor<float, 1> input2(num_elem); + Tensor<Eigen::half, 1> half_prec2(num_elem); + Tensor<Eigen::half, 1> full_prec2(num_elem); + Tensor<float, 1> input3(num_elem); + Tensor<Eigen::half, 1> half_prec3(num_elem); + Tensor<Eigen::half, 1> full_prec3(num_elem); + sycl_device.memcpyDeviceToHost(input1.data(), d_float1, num_elem*sizeof(float)); + sycl_device.memcpyDeviceToHost(input2.data(), d_float2, num_elem*sizeof(float)); + sycl_device.memcpyDeviceToHost(input3.data(), d_float3, num_elem*sizeof(float)); + sycl_device.memcpyDeviceToHost(half_prec1.data(), d_res1_half, num_elem*sizeof(Eigen::half)); + sycl_device.memcpyDeviceToHost(full_prec1.data(), d_res1_float, num_elem*sizeof(Eigen::half)); + sycl_device.memcpyDeviceToHost(half_prec2.data(), d_res2_half, num_elem*sizeof(Eigen::half)); + sycl_device.memcpyDeviceToHost(full_prec2.data(), d_res2_float, num_elem*sizeof(Eigen::half)); + sycl_device.memcpyDeviceToHost(half_prec3.data(), d_res3_half, num_elem*sizeof(Eigen::half)); + sycl_device.memcpyDeviceToHost(full_prec3.data(), d_res3_float, num_elem*sizeof(Eigen::half)); + sycl_device.synchronize(); + + for (int i = 0; i < num_elem; ++i) { + std::cout << "Checking elemwise exp " << i << " input = " << input1(i) << " full = " << full_prec1(i) << " half = " << half_prec1(i) << std::endl; + VERIFY_IS_APPROX(full_prec1(i), half_prec1(i)); + } + for (int i = 0; i < num_elem; ++i) { + std::cout << "Checking elemwise log " << i << " input = " << input2(i) << " full = " << full_prec2(i) << " half = " << half_prec2(i) << std::endl; + if(std::abs(input2(i)-1.f)<0.05f) // log lacks accuracy nearby 1 + VERIFY_IS_APPROX(full_prec2(i)+Eigen::half(0.1f), half_prec2(i)+Eigen::half(0.1f)); + else + VERIFY_IS_APPROX(full_prec2(i), half_prec2(i)); + } + for (int i = 0; i < num_elem; ++i) { + std::cout << "Checking elemwise plog1 " << i << " input = " << input3(i) << " full = " << full_prec3(i) << " half = " << half_prec3(i) << std::endl; + VERIFY_IS_APPROX(full_prec3(i), half_prec3(i)); + } +} + +void test_gpu_contractions(const Eigen::SyclDevice &sycl_device) { + int rows = 23; + int cols = 23; + int num_elem = rows*cols; + + float* d_float1 = (float*)sycl_device.allocate(num_elem * sizeof(float)); + float* d_float2 = (float*)sycl_device.allocate(num_elem * sizeof(float)); + Eigen::half* d_res_half = (Eigen::half*)sycl_device.allocate(num_elem * sizeof(Eigen::half)); + Eigen::half* d_res_float = (Eigen::half*)sycl_device.allocate(num_elem * sizeof(Eigen::half)); + + Eigen::TensorMap<Eigen::Tensor<float, 2>, Eigen::Aligned> gpu_float1( + d_float1, rows, cols); + Eigen::TensorMap<Eigen::Tensor<float, 2>, Eigen::Aligned> gpu_float2( + d_float2, rows, cols); + Eigen::TensorMap<Eigen::Tensor<Eigen::half, 2>, Eigen::Aligned> gpu_res_half( + d_res_half, rows, cols); + Eigen::TensorMap<Eigen::Tensor<Eigen::half, 2>, Eigen::Aligned> gpu_res_float( + d_res_float, rows, cols); + + gpu_float1.device(sycl_device) = gpu_float1.random() - gpu_float1.constant(0.5f); + gpu_float2.device(sycl_device) = gpu_float2.random() - gpu_float2.constant(0.5f); + + typedef typename Tensor<float, 2>::DimensionPair DimPair; + Eigen::array<DimPair, 1> dims; + gpu_res_float.device(sycl_device) = gpu_float1.contract(gpu_float2, dims).cast<Eigen::half>(); + gpu_res_half.device(sycl_device) = gpu_float1.cast<Eigen::half>().contract(gpu_float2.cast<Eigen::half>(), dims); + + Tensor<Eigen::half, 2> half_prec(rows, cols); + Tensor<Eigen::half, 2> full_prec(rows, cols); + sycl_device.memcpyDeviceToHost(half_prec.data(), d_res_half, num_elem*sizeof(Eigen::half)); + sycl_device.memcpyDeviceToHost(full_prec.data(), d_res_float, num_elem*sizeof(Eigen::half)); + sycl_device.synchronize(); + + for (int i = 0; i < rows; ++i) { + for (int j = 0; j < cols; ++j) { + std::cout << "Checking contract " << i << " " << j << full_prec(i, j) << " " << half_prec(i, j) << std::endl; + if (numext::abs(full_prec(i, j) - half_prec(i, j)) > Eigen::half(1e-2f)) { + VERIFY_IS_APPROX(full_prec(i, j), half_prec(i, j)); + } + } + } +} + +void test_gpu_reductions(const Eigen::SyclDevice &sycl_device, int size1, int size2, int redux) { + std::cout << "Reducing " << size1 << " by " << size2 + << " tensor along dim " << redux << std::endl; + + int num_elem = size1*size2; + int result_size = (redux == 1 ? size1 : size2); + + float* d_float = (float*)sycl_device.allocate(num_elem * sizeof(float)); + Eigen::half* d_res_half = (Eigen::half*)sycl_device.allocate(result_size * sizeof(Eigen::half)); + Eigen::half* d_res_float = (Eigen::half*)sycl_device.allocate(result_size * sizeof(Eigen::half)); + + Eigen::TensorMap<Eigen::Tensor<float, 2>, Eigen::Aligned> gpu_float( + d_float, size1, size2); + Eigen::TensorMap<Eigen::Tensor<Eigen::half, 1>, Eigen::Aligned> gpu_res_half( + d_res_half, result_size); + Eigen::TensorMap<Eigen::Tensor<Eigen::half, 1>, Eigen::Aligned> gpu_res_float( + d_res_float, result_size); + + gpu_float.device(sycl_device) = gpu_float.random() * 2.0f; + + Eigen::array<int, 1> redux_dim = {redux}; + gpu_res_float.device(sycl_device) = gpu_float.sum(redux_dim).cast<Eigen::half>(); + gpu_res_half.device(sycl_device) = gpu_float.cast<Eigen::half>().sum(redux_dim); + + Tensor<Eigen::half, 1> half_prec(result_size); + Tensor<Eigen::half, 1> full_prec(result_size); + sycl_device.memcpyDeviceToHost(half_prec.data(), d_res_half, result_size*sizeof(Eigen::half)); + sycl_device.memcpyDeviceToHost(full_prec.data(), d_res_float, result_size*sizeof(Eigen::half)); + sycl_device.synchronize(); + + for (int i = 0; i < result_size; ++i) { + std::cout << "EXPECTED " << full_prec(i) << " GOT " << half_prec(i) << std::endl; + VERIFY_IS_APPROX(full_prec(i), half_prec(i)); + } +} + +void test_gpu_reductions(const Eigen::SyclDevice &sycl_device) { + test_gpu_reductions(sycl_device, 13, 13, 0); + test_gpu_reductions(sycl_device, 13, 13, 1); + + test_gpu_reductions(sycl_device, 35, 36, 0); + test_gpu_reductions(sycl_device, 35, 36, 1); + + test_gpu_reductions(sycl_device, 36, 35, 0); + test_gpu_reductions(sycl_device, 36, 35, 1); +} + +void test_gpu_full_reductions(const Eigen::SyclDevice &sycl_device) { + int size = 13; + int num_elem = size*size; + + float* d_float = (float*)sycl_device.allocate(num_elem * sizeof(float)); + Eigen::half* d_res_half = (Eigen::half*)sycl_device.allocate(1 * sizeof(Eigen::half)); + Eigen::half* d_res_float = (Eigen::half*)sycl_device.allocate(1 * sizeof(Eigen::half)); + + Eigen::TensorMap<Eigen::Tensor<float, 2>, Eigen::Aligned> gpu_float( + d_float, size, size); + Eigen::TensorMap<Eigen::Tensor<Eigen::half, 0>, Eigen::Aligned> gpu_res_half( + d_res_half); + Eigen::TensorMap<Eigen::Tensor<Eigen::half, 0>, Eigen::Aligned> gpu_res_float( + d_res_float); + + gpu_float.device(sycl_device) = gpu_float.random(); + + gpu_res_float.device(sycl_device) = gpu_float.sum().cast<Eigen::half>(); + gpu_res_half.device(sycl_device) = gpu_float.cast<Eigen::half>().sum(); + + Tensor<Eigen::half, 0> half_prec; + Tensor<Eigen::half, 0> full_prec; + sycl_device.memcpyDeviceToHost(half_prec.data(), d_res_half, sizeof(Eigen::half)); + sycl_device.memcpyDeviceToHost(full_prec.data(), d_res_float, sizeof(Eigen::half)); + sycl_device.synchronize(); + + VERIFY_IS_APPROX(full_prec(), half_prec()); + + gpu_res_float.device(sycl_device) = gpu_float.maximum().cast<Eigen::half>(); + gpu_res_half.device(sycl_device) = gpu_float.cast<Eigen::half>().maximum(); + sycl_device.memcpyDeviceToHost(half_prec.data(), d_res_half, sizeof(Eigen::half)); + sycl_device.memcpyDeviceToHost(full_prec.data(), d_res_float, sizeof(Eigen::half)); + sycl_device.synchronize(); + + VERIFY_IS_APPROX(full_prec(), half_prec()); +} + +void test_gpu_forced_evals(const Eigen::SyclDevice &sycl_device) { + int num_elem = 101; + + float* d_float = (float*)sycl_device.allocate(num_elem * sizeof(float)); + float* d_res_half1 = (float*)sycl_device.allocate(num_elem * sizeof(float)); + float* d_res_half2 = (float*)sycl_device.allocate(num_elem * sizeof(float)); + float* d_res_float = (float*)sycl_device.allocate(num_elem * sizeof(float)); + + Eigen::TensorMap<Eigen::Tensor<float, 1>, Eigen::Aligned> gpu_float( + d_float, num_elem); + Eigen::TensorMap<Eigen::Tensor<float, 1>, Eigen::Aligned> gpu_res_half1( + d_res_half1, num_elem); + Eigen::TensorMap<Eigen::Tensor<float, 1>, Eigen::Unaligned> gpu_res_half2( + d_res_half2, num_elem); + Eigen::TensorMap<Eigen::Tensor<float, 1>, Eigen::Aligned> gpu_res_float( + d_res_float, num_elem); + + Eigen::array<int, 1> no_bcast; + no_bcast[0] = 1; + + gpu_float.device(sycl_device) = gpu_float.random() - gpu_float.constant(0.5f); + gpu_res_float.device(sycl_device) = gpu_float.abs(); + gpu_res_half1.device(sycl_device) = gpu_float.cast<Eigen::half>().abs().eval().cast<float>(); + gpu_res_half2.device(sycl_device) = gpu_float.cast<Eigen::half>().abs().broadcast(no_bcast).eval().cast<float>(); + + Tensor<float, 1> half_prec1(num_elem); + Tensor<float, 1> half_prec2(num_elem); + Tensor<float, 1> full_prec(num_elem); + sycl_device.memcpyDeviceToHost(half_prec1.data(), d_res_half1, num_elem*sizeof(float)); + sycl_device.memcpyDeviceToHost(half_prec2.data(), d_res_half2, num_elem*sizeof(float)); + sycl_device.memcpyDeviceToHost(full_prec.data(), d_res_float, num_elem*sizeof(float)); + sycl_device.synchronize(); + + for (int i = 0; i < num_elem; ++i) { + std::cout << "Checking forced eval " << i << full_prec(i) << " vs " << half_prec1(i) << " vs " << half_prec2(i) << std::endl; + VERIFY_IS_APPROX(full_prec(i), half_prec1(i)); + VERIFY_IS_APPROX(full_prec(i), half_prec2(i)); + } +} + +EIGEN_DECLARE_TEST(cxx11_tensor_of_float16_sycl) +{ + for (const auto& s : Eigen::get_sycl_supported_devices()) { + QueueInterface queueInterface(s); + auto sycl_device = Eigen::SyclDevice(&queueInterface); + + CALL_SUBTEST_1(test_gpu_numext(sycl_device)); + CALL_SUBTEST_1(test_gpu_conversion(sycl_device)); + CALL_SUBTEST_1(test_gpu_unary(sycl_device)); + CALL_SUBTEST_1(test_gpu_elementwise(sycl_device)); + CALL_SUBTEST_1(test_gpu_trancendental(sycl_device)); + CALL_SUBTEST_2(test_gpu_contractions(sycl_device)); + CALL_SUBTEST_3(test_gpu_reductions(sycl_device)); + CALL_SUBTEST_4(test_gpu_full_reductions(sycl_device)); + CALL_SUBTEST_5(test_gpu_forced_evals(sycl_device)); + } +}
diff --git a/unsupported/test/cxx11_tensor_padding_sycl.cpp b/unsupported/test/cxx11_tensor_padding_sycl.cpp index 727a9ff..9b16bb1 100644 --- a/unsupported/test/cxx11_tensor_padding_sycl.cpp +++ b/unsupported/test/cxx11_tensor_padding_sycl.cpp
@@ -152,6 +152,7 @@ EIGEN_DECLARE_TEST(cxx11_tensor_padding_sycl) { for (const auto& device :Eigen::get_sycl_supported_devices()) { + CALL_SUBTEST(sycl_padding_test_per_device<half>(device)); CALL_SUBTEST(sycl_padding_test_per_device<float>(device)); } }
diff --git a/unsupported/test/cxx11_tensor_patch_sycl.cpp b/unsupported/test/cxx11_tensor_patch_sycl.cpp index 7f92bec..1216ad2 100644 --- a/unsupported/test/cxx11_tensor_patch_sycl.cpp +++ b/unsupported/test/cxx11_tensor_patch_sycl.cpp
@@ -244,6 +244,7 @@ EIGEN_DECLARE_TEST(cxx11_tensor_patch_sycl) { for (const auto& device :Eigen::get_sycl_supported_devices()) { + CALL_SUBTEST(sycl_tensor_patch_test_per_device<half>(device)); CALL_SUBTEST(sycl_tensor_patch_test_per_device<float>(device)); } }
diff --git a/unsupported/test/cxx11_tensor_random_sycl.cpp b/unsupported/test/cxx11_tensor_random_sycl.cpp index 14a7c48..dcdfdd3 100644 --- a/unsupported/test/cxx11_tensor_random_sycl.cpp +++ b/unsupported/test/cxx11_tensor_random_sycl.cpp
@@ -78,6 +78,7 @@ EIGEN_DECLARE_TEST(cxx11_tensor_random_sycl) { for (const auto& device :Eigen::get_sycl_supported_devices()) { + CALL_SUBTEST(sycl_random_test_per_device<half>(device)); CALL_SUBTEST(sycl_random_test_per_device<float>(device)); #ifdef EIGEN_SYCL_DOUBLE_SUPPORT CALL_SUBTEST(sycl_random_test_per_device<double>(device));
diff --git a/unsupported/test/cxx11_tensor_reverse_sycl.cpp b/unsupported/test/cxx11_tensor_reverse_sycl.cpp index dd30c23..5ed007f 100644 --- a/unsupported/test/cxx11_tensor_reverse_sycl.cpp +++ b/unsupported/test/cxx11_tensor_reverse_sycl.cpp
@@ -248,6 +248,7 @@ #ifdef EIGEN_SYCL_DOUBLE_SUPPORT CALL_SUBTEST_4(sycl_reverse_test_per_device<double>(device)); #endif - CALL_SUBTEST_5(sycl_reverse_test_per_device<float>(device)); + CALL_SUBTEST_5(sycl_reverse_test_per_device<half>(device)); + CALL_SUBTEST_6(sycl_reverse_test_per_device<float>(device)); } }
diff --git a/unsupported/test/cxx11_tensor_shuffling_sycl.cpp b/unsupported/test/cxx11_tensor_shuffling_sycl.cpp index ca4e8b5..70c4116 100644 --- a/unsupported/test/cxx11_tensor_shuffling_sycl.cpp +++ b/unsupported/test/cxx11_tensor_shuffling_sycl.cpp
@@ -112,6 +112,7 @@ } EIGEN_DECLARE_TEST(cxx11_tensor_shuffling_sycl) { for (const auto& device : Eigen::get_sycl_supported_devices()) { + CALL_SUBTEST(sycl_shuffling_test_per_device<half>(device)); CALL_SUBTEST(sycl_shuffling_test_per_device<float>(device)); } }
diff --git a/unsupported/test/cxx11_tensor_volume_patch_sycl.cpp b/unsupported/test/cxx11_tensor_volume_patch_sycl.cpp index 8d99a48..dc86dfe 100644 --- a/unsupported/test/cxx11_tensor_volume_patch_sycl.cpp +++ b/unsupported/test/cxx11_tensor_volume_patch_sycl.cpp
@@ -217,6 +217,7 @@ EIGEN_DECLARE_TEST(cxx11_tensor_volume_patch_sycl) { for (const auto& device :Eigen::get_sycl_supported_devices()) { + CALL_SUBTEST(sycl_tensor_volume_patch_test_per_device<half>(device)); CALL_SUBTEST(sycl_tensor_volume_patch_test_per_device<float>(device)); } }