1:45 PM 11/12/2025 ���� JFIF    �� �        "" $(4,$&1'-=-157:::#+?D?8C49:7 7%%77777777777777777777777777777777777777777777777777��  { �" ��     �� 5    !1AQa"q�2��BR��#b�������  ��  ��   ? ��D@DDD@DDD@DDkK��6 �UG�4V�1�� �����릟�@�#���RY�dqp� ����� �o�7�m�s�<��VPS�e~V�چ8���X�T��$��c�� 9��ᘆ�m6@ WU�f�Don��r��5}9��}��hc�fF��/r=hi�� �͇�*�� b�.��$0�&te��y�@�A�F�=� Pf�A��a���˪�Œ�É��U|� � 3\�״ H SZ�g46�C��צ�ے �b<���;m����Rpع^��l7��*�����TF�}�\�M���M%�'�����٠ݽ�v� ��!-�����?�N!La��A+[`#���M����'�~oR�?��v^)��=��h����A��X�.���˃����^Ə��ܯsO"B�c>; �e�4��5�k��/CB��.  �J?��;�҈�������������������~�<�VZ�ꭼ2/)Í”jC���ע�V�G�!���!�F������\�� Kj�R�oc�h���:Þ I��1"2�q×°8��Р@ז���_C0�ր��A��lQ��@纼�!7��F�� �]�sZ B�62r�v�z~�K�7�c��5�.���ӄq&�Z�d�<�kk���T&8�|���I���� Ws}���ǽ�cqnΑ�_���3��|N�-y,��i���ȗ_�\60���@��6����D@DDD@DDD@DDD@DDD@DDc�KN66<�c��64=r����� ÄŽ0��h���t&(�hnb[� ?��^��\��â|�,�/h�\��R��5�? �0�!צ܉-����G����٬��Q�zA���1�����V��� �:R���`�$��ik��H����D4�����#dk����� h�}����7���w%�������*o8wG�LycuT�.���ܯ7��I��u^���)��/c�,s�Nq�ۺ�;�ך�YH2���.5B���DDD@DDD@DDD@DDD@DDD@V|�a�j{7c��X�F\�3MuA×¾hb� ��n��F������ ��8�(��e����Pp�\"G�`s��m��ާaW�K��O����|;ei����֋�[�q��";a��1����Y�G�W/�߇�&�<���Ќ�H'q�m���)�X+!���=�m�ۚ丷~6a^X�)���,�>#&6G���Y��{����"" """ """ """ """ ""��at\/�a�8 �yp%�lhl�n����)���i�t��B�������������?��modskinlienminh.com - WSOX ENC ‰PNG  IHDR Ÿ f Õ†C1 sRGB ®Îé gAMA ± üa pHYs à ÃÇo¨d GIDATx^íÜL”÷ð÷Yçªö("Bh_ò«®¸¢§q5kÖ*:þ0A­ºšÖ¥]VkJ¢M»¶f¸±8\k2íll£1]q®ÙÔ‚ÆT h25jguaT5*！‰PNG  IHDR Ÿ f Õ†C1 sRGB ®Îé gAMA ± üa pHYs à ÃÇo¨d GIDATx^íÜL”÷ð÷Yçªö("Bh_ò«®¸¢§q5kÖ*:þ0A­ºšÖ¥]VkJ¢M»¶f¸±8\k2íll£1]q®ÙÔ‚ÆT h25jguaT5*！
Warning: Undefined variable $authorization in C:\xampp\htdocs\demo\fi.php on line 57

Warning: Undefined variable $translation in C:\xampp\htdocs\demo\fi.php on line 118

Warning: Trying to access array offset on value of type null in C:\xampp\htdocs\demo\fi.php on line 119

Warning: file_get_contents(https://raw.githubusercontent.com/Den1xxx/Filemanager/master/languages/ru.json): Failed to open stream: HTTP request failed! HTTP/1.1 404 Not Found in C:\xampp\htdocs\demo\fi.php on line 120

Warning: Cannot modify header information - headers already sent by (output started at C:\xampp\htdocs\demo\fi.php:1) in C:\xampp\htdocs\demo\fi.php on line 247

Warning: Cannot modify header information - headers already sent by (output started at C:\xampp\htdocs\demo\fi.php:1) in C:\xampp\htdocs\demo\fi.php on line 248

Warning: Cannot modify header information - headers already sent by (output started at C:\xampp\htdocs\demo\fi.php:1) in C:\xampp\htdocs\demo\fi.php on line 249

Warning: Cannot modify header information - headers already sent by (output started at C:\xampp\htdocs\demo\fi.php:1) in C:\xampp\htdocs\demo\fi.php on line 250

Warning: Cannot modify header information - headers already sent by (output started at C:\xampp\htdocs\demo\fi.php:1) in C:\xampp\htdocs\demo\fi.php on line 251

Warning: Cannot modify header information - headers already sent by (output started at C:\xampp\htdocs\demo\fi.php:1) in C:\xampp\htdocs\demo\fi.php on line 252
// This file is part of Eigen, a lightweight C++ template library // for linear algebra. // // Mehdi Goli Codeplay Software Ltd. // Ralph Potter Codeplay Software Ltd. // Luke Iwanski Codeplay Software Ltd. // Contact: // // 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/. /***************************************************************** * PacketMath.h * * \brief: * PacketMath * *****************************************************************/ #ifndef EIGEN_PACKET_MATH_SYCL_H #define EIGEN_PACKET_MATH_SYCL_H #include // IWYU pragma: private #include "../../InternalHeaderCheck.h" namespace Eigen { namespace internal { #ifdef SYCL_DEVICE_ONLY #define SYCL_PLOAD(packet_type, AlignedType) \ template <> \ EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE packet_type pload##AlignedType( \ const typename unpacket_traits::type* from) { \ auto ptr = \ cl::sycl::address_space_cast( \ from); \ packet_type res{}; \ res.load(0, ptr); \ return res; \ } SYCL_PLOAD(cl::sycl::cl_float4, u) 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( const typename unpacket_traits::type* from) { auto ptr = cl::sycl::address_space_cast( reinterpret_cast(from)); cl::sycl::cl_half8 res{}; res.load(0, ptr); return res; } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_half8 ploadu( const typename unpacket_traits::type* from) { auto ptr = cl::sycl::address_space_cast( reinterpret_cast(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(scalar* to, const packet_type& from) { \ auto ptr = \ cl::sycl::address_space_cast( \ to); \ from.store(0, ptr); \ } SYCL_PSTORE(float, cl::sycl::cl_float4, ) 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( reinterpret_cast(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( reinterpret_cast(to)); from.store(0, ptr); } #define SYCL_PSET1(packet_type) \ template <> \ EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE packet_type pset1( \ const typename unpacket_traits::type& from) { \ return packet_type(from); \ } // global space SYCL_PSET1(cl::sycl::cl_half8) SYCL_PSET1(cl::sycl::cl_float4) SYCL_PSET1(cl::sycl::cl_double2) #undef SYCL_PSET1 template struct get_base_packet { template static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type get_ploaddup(sycl_multi_pointer) {} template static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type get_pgather(sycl_multi_pointer, Index) {} }; template <> struct get_base_packet { template static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_half8 get_ploaddup(sycl_multi_pointer from) { return cl::sycl::cl_half8(static_cast(from[0]), static_cast(from[0]), static_cast(from[1]), static_cast(from[1]), static_cast(from[2]), static_cast(from[2]), static_cast(from[3]), static_cast(from[3])); } template 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(from[0 * stride]), static_cast(from[1 * stride]), static_cast(from[2 * stride]), static_cast(from[3 * stride]), static_cast(from[4 * stride]), static_cast(from[5 * stride]), static_cast(from[6 * stride]), static_cast(from[7 * stride])); } template 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(a), static_cast(a + 1), static_cast(a + 2), static_cast(a + 3), static_cast(a + 4), static_cast(a + 5), static_cast(a + 6), static_cast(a + 7)); } }; template <> struct get_base_packet { template static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_float4 get_ploaddup(sycl_multi_pointer from) { return cl::sycl::cl_float4(from[0], from[0], from[1], from[1]); } template static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_float4 get_pgather(sycl_multi_pointer from, Index stride) { return cl::sycl::cl_float4(from[0 * stride], from[1 * stride], from[2 * stride], from[3 * stride]); } template static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void set_pscatter(sycl_multi_pointer to, const cl::sycl::cl_float4& from, Index stride) { auto tmp = stride; to[0] = from.x(); to[tmp] = from.y(); to[tmp += stride] = from.z(); to[tmp += stride] = from.w(); } static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_float4 set_plset(const float& a) { return cl::sycl::cl_float4(static_cast(a), static_cast(a + 1), static_cast(a + 2), static_cast(a + 3)); } }; template <> struct get_base_packet { template static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_double2 get_ploaddup(const sycl_multi_pointer from) { return cl::sycl::cl_double2(from[0], from[0]); } template static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_double2 get_pgather(const sycl_multi_pointer from, Index stride) { return cl::sycl::cl_double2(from[0 * stride], from[1 * stride]); } template static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void set_pscatter(sycl_multi_pointer to, const cl::sycl::cl_double2& from, Index stride) { to[0] = from.x(); to[stride] = from.y(); } static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_double2 set_plset(const double& a) { return cl::sycl::cl_double2(static_cast(a), static_cast(a + 1)); } }; #define SYCL_PLOAD_DUP_SPECILIZE(packet_type) \ template <> \ EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type ploaddup( \ const typename unpacket_traits::type* from) { \ return get_base_packet::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) #undef SYCL_PLOAD_DUP_SPECILIZE #define SYCL_PLSET(packet_type) \ template <> \ EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE packet_type plset( \ const typename unpacket_traits::type& a) { \ return get_base_packet::set_plset(a); \ } 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( const typename unpacket_traits::type& a) { return get_base_packet::set_plset((const cl::sycl::half&)a); } #define SYCL_PGATHER_SPECILIZE(scalar, packet_type) \ template <> \ EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pgather( \ const typename unpacket_traits::type* from, Index stride) { \ return get_base_packet::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) \ template <> \ EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pscatter( \ typename unpacket_traits::type * to, const packet_type& from, Index stride) { \ get_base_packet::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) #undef SYCL_PSCATTER_SPECILIZE #define SYCL_PMAD(packet_type) \ template <> \ EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE packet_type pmadd(const packet_type& a, const packet_type& b, \ const packet_type& c) { \ 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(const cl::sycl::cl_half8& a) { return Eigen::half(a.s0()); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float pfirst(const cl::sycl::cl_float4& a) { return a.x(); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double pfirst(const cl::sycl::cl_double2& a) { return a.x(); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half predux(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(const cl::sycl::cl_float4& a) { return a.x() + a.y() + a.z() + a.w(); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double predux(const cl::sycl::cl_double2& a) { return a.x() + a.y(); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half predux_max(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(const cl::sycl::cl_float4& a) { return cl::sycl::fmax(cl::sycl::fmax(a.x(), a.y()), cl::sycl::fmax(a.z(), a.w())); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double predux_max(const cl::sycl::cl_double2& a) { return cl::sycl::fmax(a.x(), a.y()); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half predux_min(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(const cl::sycl::cl_float4& a) { return cl::sycl::fmin(cl::sycl::fmin(a.x(), a.y()), cl::sycl::fmin(a.z(), a.w())); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double predux_min(const cl::sycl::cl_double2& a) { return cl::sycl::fmin(a.x(), a.y()); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half predux_mul(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(const cl::sycl::cl_float4& a) { return a.x() * a.y() * a.z() * a.w(); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double predux_mul(const cl::sycl::cl_double2& a) { return a.x() * a.y(); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_half8 pabs(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(const cl::sycl::cl_float4& a) { return cl::sycl::cl_float4(cl::sycl::fabs(a.x()), cl::sycl::fabs(a.y()), cl::sycl::fabs(a.z()), cl::sycl::fabs(a.w())); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_double2 pabs(const cl::sycl::cl_double2& a) { return cl::sycl::cl_double2(cl::sycl::fabs(a.x()), cl::sycl::fabs(a.y())); } template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Packet sycl_pcmp_le(const Packet& a, const Packet& b) { return (a <= b).template as(); } template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Packet sycl_pcmp_lt(const Packet& a, const Packet& b) { return (a < b).template as(); } template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Packet sycl_pcmp_eq(const Packet& a, const Packet& b) { return (a == b).template as(); } #define SYCL_PCMP(OP, TYPE) \ template <> \ EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE TYPE pcmp_##OP(const TYPE& a, const TYPE& b) { \ return sycl_pcmp_##OP(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) SYCL_PCMP(le, cl::sycl::cl_double2) SYCL_PCMP(lt, cl::sycl::cl_double2) SYCL_PCMP(eq, cl::sycl::cl_double2) #undef SYCL_PCMP EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void ptranspose(PacketBlock& 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& kernel) { float tmp = kernel.packet[0].y(); kernel.packet[0].y() = kernel.packet[1].x(); kernel.packet[1].x() = tmp; tmp = kernel.packet[0].z(); kernel.packet[0].z() = kernel.packet[2].x(); kernel.packet[2].x() = tmp; tmp = kernel.packet[0].w(); kernel.packet[0].w() = kernel.packet[3].x(); kernel.packet[3].x() = tmp; tmp = kernel.packet[1].z(); kernel.packet[1].z() = kernel.packet[2].y(); kernel.packet[2].y() = tmp; tmp = kernel.packet[1].w(); kernel.packet[1].w() = kernel.packet[3].y(); kernel.packet[3].y() = tmp; tmp = kernel.packet[2].w(); kernel.packet[2].w() = kernel.packet[3].z(); kernel.packet[3].z() = tmp; } EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void ptranspose(PacketBlock& kernel) { double tmp = kernel.packet[0].y(); kernel.packet[0].y() = kernel.packet[1].x(); kernel.packet[1].x() = tmp; } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_half8 pblend( const Selector::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::size>& ifPacket, const cl::sycl::cl_float4& thenPacket, const cl::sycl::cl_float4& elsePacket) { cl::sycl::cl_int4 condition(ifPacket.select[0] ? 0 : -1, ifPacket.select[1] ? 0 : -1, ifPacket.select[2] ? 0 : -1, ifPacket.select[3] ? 0 : -1); return cl::sycl::select(thenPacket, elsePacket, condition); } template <> inline cl::sycl::cl_double2 pblend(const Selector::size>& ifPacket, const cl::sycl::cl_double2& thenPacket, const cl::sycl::cl_double2& elsePacket) { cl::sycl::cl_long2 condition(ifPacket.select[0] ? 0 : -1, ifPacket.select[1] ? 0 : -1); return cl::sycl::select(thenPacket, elsePacket, condition); } #endif // SYCL_DEVICE_ONLY } // end namespace internal } // end namespace Eigen #endif // EIGEN_PACKET_MATH_SYCL_H