InteropHeaders.h (7428B)
1 // This file is part of Eigen, a lightweight C++ template library 2 // for linear algebra. 3 // 4 // Mehdi Goli Codeplay Software Ltd. 5 // Ralph Potter Codeplay Software Ltd. 6 // Luke Iwanski Codeplay Software Ltd. 7 // Contact: <eigen@codeplay.com> 8 // 9 // This Source Code Form is subject to the terms of the Mozilla 10 // Public License v. 2.0. If a copy of the MPL was not distributed 11 // with this file, You can obtain one at http://mozilla.org/MPL/2.0/. 12 13 /***************************************************************** 14 * InteropHeaders.h 15 * 16 * \brief: 17 * InteropHeaders 18 * 19 *****************************************************************/ 20 21 #ifndef EIGEN_INTEROP_HEADERS_SYCL_H 22 #define EIGEN_INTEROP_HEADERS_SYCL_H 23 24 namespace Eigen { 25 26 #if !defined(EIGEN_DONT_VECTORIZE_SYCL) 27 28 namespace internal { 29 30 template <int has_blend, int lengths> 31 struct sycl_packet_traits : default_packet_traits { 32 enum { 33 Vectorizable = 1, 34 AlignedOnScalar = 1, 35 size = lengths, 36 HasHalfPacket = 0, 37 HasDiv = 1, 38 HasLog = 1, 39 HasExp = 1, 40 HasSqrt = 1, 41 HasRsqrt = 1, 42 HasSin = 1, 43 HasCos = 1, 44 HasTan = 1, 45 HasASin = 1, 46 HasACos = 1, 47 HasATan = 1, 48 HasSinh = 1, 49 HasCosh = 1, 50 HasTanh = 1, 51 HasLGamma = 0, 52 HasDiGamma = 0, 53 HasZeta = 0, 54 HasPolygamma = 0, 55 HasErf = 0, 56 HasErfc = 0, 57 HasNdtri = 0, 58 HasIGamma = 0, 59 HasIGammac = 0, 60 HasBetaInc = 0, 61 HasBlend = has_blend, 62 // This flag is used to indicate whether packet comparison is supported. 63 // pcmp_eq, pcmp_lt and pcmp_le should be defined for it to be true. 64 HasCmp = 1, 65 HasMax = 1, 66 HasMin = 1, 67 HasMul = 1, 68 HasAdd = 1, 69 HasFloor = 1, 70 HasRound = 1, 71 HasRint = 1, 72 HasLog1p = 1, 73 HasExpm1 = 1, 74 HasCeil = 1, 75 }; 76 }; 77 78 #ifdef SYCL_DEVICE_ONLY 79 #define SYCL_PACKET_TRAITS(packet_type, has_blend, unpacket_type, lengths) \ 80 template <> \ 81 struct packet_traits<unpacket_type> \ 82 : sycl_packet_traits<has_blend, lengths> { \ 83 typedef packet_type type; \ 84 typedef packet_type half; \ 85 }; 86 87 SYCL_PACKET_TRAITS(cl::sycl::cl_float4, 1, float, 4) 88 SYCL_PACKET_TRAITS(cl::sycl::cl_float4, 1, const float, 4) 89 SYCL_PACKET_TRAITS(cl::sycl::cl_double2, 0, double, 2) 90 SYCL_PACKET_TRAITS(cl::sycl::cl_double2, 0, const double, 2) 91 #undef SYCL_PACKET_TRAITS 92 93 // Make sure this is only available when targeting a GPU: we don't want to 94 // introduce conflicts between these packet_traits definitions and the ones 95 // we'll use on the host side (SSE, AVX, ...) 96 #define SYCL_ARITHMETIC(packet_type) \ 97 template <> \ 98 struct is_arithmetic<packet_type> { \ 99 enum { value = true }; \ 100 }; 101 SYCL_ARITHMETIC(cl::sycl::cl_float4) 102 SYCL_ARITHMETIC(cl::sycl::cl_double2) 103 #undef SYCL_ARITHMETIC 104 105 #define SYCL_UNPACKET_TRAITS(packet_type, unpacket_type, lengths) \ 106 template <> \ 107 struct unpacket_traits<packet_type> { \ 108 typedef unpacket_type type; \ 109 enum { size = lengths, vectorizable = true, alignment = Aligned16 }; \ 110 typedef packet_type half; \ 111 }; 112 SYCL_UNPACKET_TRAITS(cl::sycl::cl_float4, float, 4) 113 SYCL_UNPACKET_TRAITS(cl::sycl::cl_double2, double, 2) 114 115 #undef SYCL_UNPACKET_TRAITS 116 #endif 117 118 } // end namespace internal 119 120 #endif 121 122 namespace TensorSycl { 123 namespace internal { 124 125 template <typename PacketReturnType, int PacketSize> 126 struct PacketWrapper; 127 // This function should never get called on the device 128 #ifndef SYCL_DEVICE_ONLY 129 template <typename PacketReturnType, int PacketSize> 130 struct PacketWrapper { 131 typedef typename ::Eigen::internal::unpacket_traits<PacketReturnType>::type 132 Scalar; 133 template <typename Index> 134 EIGEN_DEVICE_FUNC static Scalar scalarize(Index, PacketReturnType &) { 135 eigen_assert(false && "THERE IS NO PACKETIZE VERSION FOR THE CHOSEN TYPE"); 136 abort(); 137 } 138 EIGEN_DEVICE_FUNC static PacketReturnType convert_to_packet_type(Scalar in, 139 Scalar) { 140 return ::Eigen::internal::template plset<PacketReturnType>(in); 141 } 142 EIGEN_DEVICE_FUNC static void set_packet(PacketReturnType, Scalar *) { 143 eigen_assert(false && "THERE IS NO PACKETIZE VERSION FOR THE CHOSEN TYPE"); 144 abort(); 145 } 146 }; 147 148 #elif defined(SYCL_DEVICE_ONLY) 149 template <typename PacketReturnType> 150 struct PacketWrapper<PacketReturnType, 4> { 151 typedef typename ::Eigen::internal::unpacket_traits<PacketReturnType>::type 152 Scalar; 153 template <typename Index> 154 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static Scalar scalarize(Index index, PacketReturnType &in) { 155 switch (index) { 156 case 0: 157 return in.x(); 158 case 1: 159 return in.y(); 160 case 2: 161 return in.z(); 162 case 3: 163 return in.w(); 164 default: 165 //INDEX MUST BE BETWEEN 0 and 3.There is no abort function in SYCL kernel. so we cannot use abort here. 166 // The code will never reach here 167 __builtin_unreachable(); 168 } 169 __builtin_unreachable(); 170 } 171 172 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static PacketReturnType convert_to_packet_type( 173 Scalar in, Scalar other) { 174 return PacketReturnType(in, other, other, other); 175 } 176 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static void set_packet(PacketReturnType &lhs, Scalar *rhs) { 177 lhs = PacketReturnType(rhs[0], rhs[1], rhs[2], rhs[3]); 178 } 179 }; 180 181 template <typename PacketReturnType> 182 struct PacketWrapper<PacketReturnType, 1> { 183 typedef typename ::Eigen::internal::unpacket_traits<PacketReturnType>::type 184 Scalar; 185 template <typename Index> 186 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static Scalar scalarize(Index, PacketReturnType &in) { 187 return in; 188 } 189 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static PacketReturnType convert_to_packet_type(Scalar in, 190 Scalar) { 191 return PacketReturnType(in); 192 } 193 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static void set_packet(PacketReturnType &lhs, Scalar *rhs) { 194 lhs = rhs[0]; 195 } 196 }; 197 198 template <typename PacketReturnType> 199 struct PacketWrapper<PacketReturnType, 2> { 200 typedef typename ::Eigen::internal::unpacket_traits<PacketReturnType>::type 201 Scalar; 202 template <typename Index> 203 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static Scalar scalarize(Index index, PacketReturnType &in) { 204 switch (index) { 205 case 0: 206 return in.x(); 207 case 1: 208 return in.y(); 209 default: 210 //INDEX MUST BE BETWEEN 0 and 1.There is no abort function in SYCL kernel. so we cannot use abort here. 211 // The code will never reach here 212 __builtin_unreachable(); 213 } 214 __builtin_unreachable(); 215 } 216 217 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static PacketReturnType convert_to_packet_type( 218 Scalar in, Scalar other) { 219 return PacketReturnType(in, other); 220 } 221 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static void set_packet(PacketReturnType &lhs, Scalar *rhs) { 222 lhs = PacketReturnType(rhs[0], rhs[1]); 223 } 224 }; 225 226 #endif 227 228 } // end namespace internal 229 } // end namespace TensorSycl 230 } // end namespace Eigen 231 232 #endif // EIGEN_INTEROP_HEADERS_SYCL_H