arch/SYCL/InteropHeaders.h

0001 // This file is part of Eigen, a lightweight C++ template library
0002 // for linear algebra.
0003 //
0004 // Mehdi Goli    Codeplay Software Ltd.
0005 // Ralph Potter  Codeplay Software Ltd.
0006 // Luke Iwanski  Codeplay Software Ltd.
0007 // Contact: <eigen@codeplay.com>
0008 //
0009 // This Source Code Form is subject to the terms of the Mozilla
0010 // Public License v. 2.0. If a copy of the MPL was not distributed
0011 // with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
0012
0013 /*****************************************************************
0014  * InteropHeaders.h
0015  *
0016  * \brief:
0017  *  InteropHeaders
0018  *
0019  *****************************************************************/
0020
0021 #ifndef EIGEN_INTEROP_HEADERS_SYCL_H
0022 #define EIGEN_INTEROP_HEADERS_SYCL_H
0023
0024 namespace Eigen {
0025
0026 #if !defined(EIGEN_DONT_VECTORIZE_SYCL)
0027
0028 namespace internal {
0029
0030 template <int has_blend, int lengths>
0031 struct sycl_packet_traits : default_packet_traits {
0032   enum {
0033     Vectorizable = 1,
0034     AlignedOnScalar = 1,
0035     size = lengths,
0036     HasHalfPacket = 0,
0037     HasDiv = 1,
0038     HasLog = 1,
0039     HasExp = 1,
0040     HasSqrt = 1,
0041     HasRsqrt = 1,
0042     HasSin = 1,
0043     HasCos = 1,
0044     HasTan = 1,
0045     HasASin = 1,
0046     HasACos = 1,
0047     HasATan = 1,
0048     HasSinh = 1,
0049     HasCosh = 1,
0050     HasTanh = 1,
0051     HasLGamma = 0,
0052     HasDiGamma = 0,
0053     HasZeta = 0,
0054     HasPolygamma = 0,
0055     HasErf = 0,
0056     HasErfc = 0,
0057     HasNdtri = 0,
0058     HasIGamma = 0,
0059     HasIGammac = 0,
0060     HasBetaInc = 0,
0061     HasBlend = has_blend,
0062     // This flag is used to indicate whether packet comparison is supported.
0063     // pcmp_eq, pcmp_lt and pcmp_le should be defined for it to be true.
0064     HasCmp = 1,
0065     HasMax = 1,
0066     HasMin = 1,
0067     HasMul = 1,
0068     HasAdd = 1,
0069     HasFloor = 1,
0070     HasRound = 1,
0071     HasRint = 1,
0072     HasLog1p = 1,
0073     HasExpm1 = 1,
0074     HasCeil = 1,
0075   };
0076 };
0077
0078 #ifdef SYCL_DEVICE_ONLY
0079 #define SYCL_PACKET_TRAITS(packet_type, has_blend, unpacket_type, lengths) \
0080   template <>                                                              \
0081   struct packet_traits<unpacket_type>                                      \
0082       : sycl_packet_traits<has_blend, lengths> {                           \
0083     typedef packet_type type;                                              \
0084     typedef packet_type half;                                              \
0085   };
0086
0087 SYCL_PACKET_TRAITS(cl::sycl::cl_float4, 1, float, 4)
0088 SYCL_PACKET_TRAITS(cl::sycl::cl_float4, 1, const float, 4)
0089 SYCL_PACKET_TRAITS(cl::sycl::cl_double2, 0, double, 2)
0090 SYCL_PACKET_TRAITS(cl::sycl::cl_double2, 0, const double, 2)
0091 #undef SYCL_PACKET_TRAITS
0092
0093 // Make sure this is only available when targeting a GPU: we don't want to
0094 // introduce conflicts between these packet_traits definitions and the ones
0095 // we'll use on the host side (SSE, AVX, ...)
0096 #define SYCL_ARITHMETIC(packet_type)  \
0097   template <>                         \
0098   struct is_arithmetic<packet_type> { \
0099     enum { value = true };            \
0100   };
0101 SYCL_ARITHMETIC(cl::sycl::cl_float4)
0102 SYCL_ARITHMETIC(cl::sycl::cl_double2)
0103 #undef SYCL_ARITHMETIC
0104
0105 #define SYCL_UNPACKET_TRAITS(packet_type, unpacket_type, lengths)        \
0106   template <>                                                            \
0107   struct unpacket_traits<packet_type> {                                  \
0108     typedef unpacket_type type;                                          \
0109     enum { size = lengths, vectorizable = true, alignment = Aligned16 }; \
0110     typedef packet_type half;                                            \
0111   };
0112 SYCL_UNPACKET_TRAITS(cl::sycl::cl_float4, float, 4)
0113 SYCL_UNPACKET_TRAITS(cl::sycl::cl_double2, double, 2)
0114
0115 #undef SYCL_UNPACKET_TRAITS
0116 #endif
0117
0118 }  // end namespace internal
0119
0120 #endif
0121
0122 namespace TensorSycl {
0123 namespace internal {
0124
0125 template <typename PacketReturnType, int PacketSize>
0126 struct PacketWrapper;
0127 // This function should never get called on the device
0128 #ifndef SYCL_DEVICE_ONLY
0129 template <typename PacketReturnType, int PacketSize>
0130 struct PacketWrapper {
0131   typedef typename ::Eigen::internal::unpacket_traits<PacketReturnType>::type
0132       Scalar;
0133   template <typename Index>
0134   EIGEN_DEVICE_FUNC static Scalar scalarize(Index, PacketReturnType &) {
0135     eigen_assert(false && "THERE IS NO PACKETIZE VERSION FOR  THE CHOSEN TYPE");
0136     abort();
0137   }
0138   EIGEN_DEVICE_FUNC static PacketReturnType convert_to_packet_type(Scalar in,
0139                                                                    Scalar) {
0140     return ::Eigen::internal::template plset<PacketReturnType>(in);
0141   }
0142   EIGEN_DEVICE_FUNC static void set_packet(PacketReturnType, Scalar *) {
0143     eigen_assert(false && "THERE IS NO PACKETIZE VERSION FOR  THE CHOSEN TYPE");
0144     abort();
0145   }
0146 };
0147
0148 #elif defined(SYCL_DEVICE_ONLY)
0149 template <typename PacketReturnType>
0150 struct PacketWrapper<PacketReturnType, 4> {
0151   typedef typename ::Eigen::internal::unpacket_traits<PacketReturnType>::type
0152       Scalar;
0153   template <typename Index>
0154   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static Scalar scalarize(Index index, PacketReturnType &in) {
0155     switch (index) {
0156       case 0:
0157         return in.x();
0158       case 1:
0159         return in.y();
0160       case 2:
0161         return in.z();
0162       case 3:
0163         return in.w();
0164       default:
0165       //INDEX MUST BE BETWEEN 0 and 3.There is no abort function in SYCL kernel. so we cannot use abort here.
0166       // The code will never reach here
0167       __builtin_unreachable();
0168     }
0169     __builtin_unreachable();
0170   }
0171
0172   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static PacketReturnType convert_to_packet_type(
0173       Scalar in, Scalar other) {
0174     return PacketReturnType(in, other, other, other);
0175   }
0176   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static void set_packet(PacketReturnType &lhs, Scalar *rhs) {
0177     lhs = PacketReturnType(rhs[0], rhs[1], rhs[2], rhs[3]);
0178   }
0179 };
0180
0181 template <typename PacketReturnType>
0182 struct PacketWrapper<PacketReturnType, 1> {
0183   typedef typename ::Eigen::internal::unpacket_traits<PacketReturnType>::type
0184       Scalar;
0185   template <typename Index>
0186   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static Scalar scalarize(Index, PacketReturnType &in) {
0187     return in;
0188   }
0189   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static PacketReturnType convert_to_packet_type(Scalar in,
0190                                                                    Scalar) {
0191     return PacketReturnType(in);
0192   }
0193   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static void set_packet(PacketReturnType &lhs, Scalar *rhs) {
0194     lhs = rhs[0];
0195   }
0196 };
0197
0198 template <typename PacketReturnType>
0199 struct PacketWrapper<PacketReturnType, 2> {
0200   typedef typename ::Eigen::internal::unpacket_traits<PacketReturnType>::type
0201       Scalar;
0202   template <typename Index>
0203   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static Scalar scalarize(Index index, PacketReturnType &in) {
0204     switch (index) {
0205       case 0:
0206         return in.x();
0207       case 1:
0208         return in.y();
0209       default:
0210         //INDEX MUST BE BETWEEN 0 and 1.There is no abort function in SYCL kernel. so we cannot use abort here.
0211       // The code will never reach here
0212         __builtin_unreachable();
0213     }
0214     __builtin_unreachable();
0215   }
0216
0217   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static PacketReturnType convert_to_packet_type(
0218       Scalar in, Scalar other) {
0219     return PacketReturnType(in, other);
0220   }
0221   EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE static void set_packet(PacketReturnType &lhs, Scalar *rhs) {
0222     lhs = PacketReturnType(rhs[0], rhs[1]);
0223   }
0224 };
0225
0226 #endif
0227
0228 }  // end namespace internal
0229 }  // end namespace TensorSycl
0230 }  // end namespace Eigen
0231
0232 #endif  // EIGEN_INTEROP_HEADERS_SYCL_H