Please, help us to better know about our user community by answering the following short survey: https://forms.gle/wpyrxWi18ox9Z5ae9
Eigen  3.4.0
InteropHeaders.h
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
Namespace containing all symbols from the Eigen library.
Definition: Core:141
EIGEN_DEFAULT_DENSE_INDEX_TYPE Index
The Index type as used for the API.
Definition: Meta.h:74
Definition: Eigen_Colamd.h:50