Please, help us to better know about our user community by answering the following short survey: https://forms.gle/wpyrxWi18ox9Z5ae9
Eigen  3.4.0
 
Loading...
Searching...
No Matches
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
24namespace Eigen {
25
26#if !defined(EIGEN_DONT_VECTORIZE_SYCL)
27
28namespace internal {
29
30template <int has_blend, int lengths>
31struct 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
87SYCL_PACKET_TRAITS(cl::sycl::cl_float4, 1, float, 4)
88SYCL_PACKET_TRAITS(cl::sycl::cl_float4, 1, const float, 4)
89SYCL_PACKET_TRAITS(cl::sycl::cl_double2, 0, double, 2)
90SYCL_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 };
101SYCL_ARITHMETIC(cl::sycl::cl_float4)
102SYCL_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 };
112SYCL_UNPACKET_TRAITS(cl::sycl::cl_float4, float, 4)
113SYCL_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
122namespace TensorSycl {
123namespace internal {
124
125template <typename PacketReturnType, int PacketSize>
126struct PacketWrapper;
127// This function should never get called on the device
128#ifndef SYCL_DEVICE_ONLY
129template <typename PacketReturnType, int PacketSize>
130struct 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)
149template <typename PacketReturnType>
150struct 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
181template <typename PacketReturnType>
182struct 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
198template <typename PacketReturnType>
199struct 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