Eigen  5.0.1
 
Loading...
Searching...
No Matches
PacketMath.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// SPDX-FileCopyrightText: The Eigen Authors
13// SPDX-License-Identifier: MPL-2.0
14
15/*****************************************************************
16 * PacketMath.h
17 *
18 * \brief:
19 * PacketMath
20 *
21 *****************************************************************/
22
23#ifndef EIGEN_PACKET_MATH_SYCL_H
24#define EIGEN_PACKET_MATH_SYCL_H
25#include <type_traits>
26
27// IWYU pragma: private
28#include "../../InternalHeaderCheck.h"
29
30namespace Eigen {
31
32namespace internal {
33#ifdef SYCL_DEVICE_ONLY
34#define SYCL_PLOAD(packet_type, AlignedType) \
35 template <> \
36 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE packet_type pload##AlignedType<packet_type>( \
37 const typename unpacket_traits<packet_type>::type* from) { \
38 auto ptr = \
39 cl::sycl::address_space_cast<cl::sycl::access::address_space::generic_space, cl::sycl::access::decorated::no>( \
40 from); \
41 packet_type res{}; \
42 res.load(0, ptr); \
43 return res; \
44 }
45
46SYCL_PLOAD(cl::sycl::cl_float4, u)
47SYCL_PLOAD(cl::sycl::cl_float4, )
48SYCL_PLOAD(cl::sycl::cl_double2, u)
49SYCL_PLOAD(cl::sycl::cl_double2, )
50#undef SYCL_PLOAD
51
52template <>
53EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_half8 pload<cl::sycl::cl_half8>(
54 const typename unpacket_traits<cl::sycl::cl_half8>::type* from) {
55 auto ptr =
56 cl::sycl::address_space_cast<cl::sycl::access::address_space::generic_space, cl::sycl::access::decorated::no>(
57 reinterpret_cast<const cl::sycl::cl_half*>(from));
58 cl::sycl::cl_half8 res{};
59 res.load(0, ptr);
60 return res;
61}
62
63template <>
64EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_half8 ploadu<cl::sycl::cl_half8>(
65 const typename unpacket_traits<cl::sycl::cl_half8>::type* from) {
66 auto ptr =
67 cl::sycl::address_space_cast<cl::sycl::access::address_space::generic_space, cl::sycl::access::decorated::no>(
68 reinterpret_cast<const cl::sycl::cl_half*>(from));
69 cl::sycl::cl_half8 res{};
70 res.load(0, ptr);
71 return res;
72}
73
74#define SYCL_PSTORE(scalar, packet_type, alignment) \
75 template <> \
76 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void pstore##alignment(scalar* to, const packet_type& from) { \
77 auto ptr = \
78 cl::sycl::address_space_cast<cl::sycl::access::address_space::generic_space, cl::sycl::access::decorated::no>( \
79 to); \
80 from.store(0, ptr); \
81 }
82
83SYCL_PSTORE(float, cl::sycl::cl_float4, )
84SYCL_PSTORE(float, cl::sycl::cl_float4, u)
85SYCL_PSTORE(double, cl::sycl::cl_double2, )
86SYCL_PSTORE(double, cl::sycl::cl_double2, u)
87#undef SYCL_PSTORE
88
89template <>
90EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void pstoreu(Eigen::half* to, const cl::sycl::cl_half8& from) {
91 auto ptr =
92 cl::sycl::address_space_cast<cl::sycl::access::address_space::generic_space, cl::sycl::access::decorated::no>(
93 reinterpret_cast<cl::sycl::cl_half*>(to));
94 from.store(0, ptr);
95}
96
97template <>
98EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void pstore(Eigen::half* to, const cl::sycl::cl_half8& from) {
99 auto ptr =
100 cl::sycl::address_space_cast<cl::sycl::access::address_space::generic_space, cl::sycl::access::decorated::no>(
101 reinterpret_cast<cl::sycl::cl_half*>(to));
102 from.store(0, ptr);
103}
104
105#define SYCL_PSET1(packet_type) \
106 template <> \
107 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE packet_type pset1<packet_type>( \
108 const typename unpacket_traits<packet_type>::type& from) { \
109 return packet_type(from); \
110 }
111
112// global space
113SYCL_PSET1(cl::sycl::cl_half8)
114SYCL_PSET1(cl::sycl::cl_float4)
115SYCL_PSET1(cl::sycl::cl_double2)
116
117#undef SYCL_PSET1
118
119template <typename packet_type>
120struct get_base_packet {
121 template <typename sycl_multi_pointer>
122 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type get_ploaddup(sycl_multi_pointer) {}
123
124 template <typename sycl_multi_pointer>
125 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type get_pgather(sycl_multi_pointer, Index) {}
126};
127
128template <>
129struct get_base_packet<cl::sycl::cl_half8> {
130 template <typename sycl_multi_pointer>
131 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_half8 get_ploaddup(sycl_multi_pointer from) {
132 return cl::sycl::cl_half8(static_cast<cl::sycl::half>(from[0]), static_cast<cl::sycl::half>(from[0]),
133 static_cast<cl::sycl::half>(from[1]), static_cast<cl::sycl::half>(from[1]),
134 static_cast<cl::sycl::half>(from[2]), static_cast<cl::sycl::half>(from[2]),
135 static_cast<cl::sycl::half>(from[3]), static_cast<cl::sycl::half>(from[3]));
136 }
137 template <typename sycl_multi_pointer>
138 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_half8 get_pgather(sycl_multi_pointer from, Index stride) {
139 return cl::sycl::cl_half8(
140 static_cast<cl::sycl::half>(from[0 * stride]), static_cast<cl::sycl::half>(from[1 * stride]),
141 static_cast<cl::sycl::half>(from[2 * stride]), static_cast<cl::sycl::half>(from[3 * stride]),
142 static_cast<cl::sycl::half>(from[4 * stride]), static_cast<cl::sycl::half>(from[5 * stride]),
143 static_cast<cl::sycl::half>(from[6 * stride]), static_cast<cl::sycl::half>(from[7 * stride]));
144 }
145
146 template <typename sycl_multi_pointer>
147 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void set_pscatter(sycl_multi_pointer to, const cl::sycl::cl_half8& from,
148 Index stride) {
149 auto tmp = stride;
150 to[0] = Eigen::half(from.s0());
151 to[tmp] = Eigen::half(from.s1());
152 to[tmp += stride] = Eigen::half(from.s2());
153 to[tmp += stride] = Eigen::half(from.s3());
154 to[tmp += stride] = Eigen::half(from.s4());
155 to[tmp += stride] = Eigen::half(from.s5());
156 to[tmp += stride] = Eigen::half(from.s6());
157 to[tmp += stride] = Eigen::half(from.s7());
158 }
159 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_half8 set_plset(const cl::sycl::half& a) {
160 return cl::sycl::cl_half8(static_cast<cl::sycl::half>(a), static_cast<cl::sycl::half>(a + 1),
161 static_cast<cl::sycl::half>(a + 2), static_cast<cl::sycl::half>(a + 3),
162 static_cast<cl::sycl::half>(a + 4), static_cast<cl::sycl::half>(a + 5),
163 static_cast<cl::sycl::half>(a + 6), static_cast<cl::sycl::half>(a + 7));
164 }
165};
166
167template <>
168struct get_base_packet<cl::sycl::cl_float4> {
169 template <typename sycl_multi_pointer>
170 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_float4 get_ploaddup(sycl_multi_pointer from) {
171 return cl::sycl::cl_float4(from[0], from[0], from[1], from[1]);
172 }
173 template <typename sycl_multi_pointer>
174 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_float4 get_pgather(sycl_multi_pointer from, Index stride) {
175 return cl::sycl::cl_float4(from[0 * stride], from[1 * stride], from[2 * stride], from[3 * stride]);
176 }
177
178 template <typename sycl_multi_pointer>
179 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void set_pscatter(sycl_multi_pointer to, const cl::sycl::cl_float4& from,
180 Index stride) {
181 auto tmp = stride;
182 to[0] = from.x();
183 to[tmp] = from.y();
184 to[tmp += stride] = from.z();
185 to[tmp += stride] = from.w();
186 }
187 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_float4 set_plset(const float& a) {
188 return cl::sycl::cl_float4(static_cast<float>(a), static_cast<float>(a + 1), static_cast<float>(a + 2),
189 static_cast<float>(a + 3));
190 }
191};
192
193template <>
194struct get_base_packet<cl::sycl::cl_double2> {
195 template <typename sycl_multi_pointer>
196 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_double2 get_ploaddup(const sycl_multi_pointer from) {
197 return cl::sycl::cl_double2(from[0], from[0]);
198 }
199
200 template <typename sycl_multi_pointer, typename Index>
201 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_double2 get_pgather(const sycl_multi_pointer from,
202 Index stride) {
203 return cl::sycl::cl_double2(from[0 * stride], from[1 * stride]);
204 }
205
206 template <typename sycl_multi_pointer>
207 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void set_pscatter(sycl_multi_pointer to,
208 const cl::sycl::cl_double2& from, Index stride) {
209 to[0] = from.x();
210 to[stride] = from.y();
211 }
212
213 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE cl::sycl::cl_double2 set_plset(const double& a) {
214 return cl::sycl::cl_double2(static_cast<double>(a), static_cast<double>(a + 1));
215 }
216};
217
218#define SYCL_PLOAD_DUP_SPECILIZE(packet_type) \
219 template <> \
220 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type ploaddup<packet_type>( \
221 const typename unpacket_traits<packet_type>::type* from) { \
222 return get_base_packet<packet_type>::get_ploaddup(from); \
223 }
224
225SYCL_PLOAD_DUP_SPECILIZE(cl::sycl::cl_half8)
226SYCL_PLOAD_DUP_SPECILIZE(cl::sycl::cl_float4)
227SYCL_PLOAD_DUP_SPECILIZE(cl::sycl::cl_double2)
228
229#undef SYCL_PLOAD_DUP_SPECILIZE
230
231#define SYCL_PLSET(packet_type) \
232 template <> \
233 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE packet_type plset<packet_type>( \
234 const typename unpacket_traits<packet_type>::type& a) { \
235 return get_base_packet<packet_type>::set_plset(a); \
236 }
237SYCL_PLSET(cl::sycl::cl_float4)
238SYCL_PLSET(cl::sycl::cl_double2)
239#undef SYCL_PLSET
240
241template <>
242EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_half8 plset<cl::sycl::cl_half8>(
243 const typename unpacket_traits<cl::sycl::cl_half8>::type& a) {
244 return get_base_packet<cl::sycl::cl_half8>::set_plset((const cl::sycl::half&)a);
245}
246
247#define SYCL_PGATHER_SPECILIZE(scalar, packet_type) \
248 template <> \
249 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE packet_type pgather<scalar, packet_type>( \
250 const typename unpacket_traits<packet_type>::type* from, Index stride) { \
251 return get_base_packet<packet_type>::get_pgather(from, stride); \
252 }
253
254SYCL_PGATHER_SPECILIZE(Eigen::half, cl::sycl::cl_half8)
255SYCL_PGATHER_SPECILIZE(float, cl::sycl::cl_float4)
256SYCL_PGATHER_SPECILIZE(double, cl::sycl::cl_double2)
257#undef SYCL_PGATHER_SPECILIZE
258
259#define SYCL_PSCATTER_SPECILIZE(scalar, packet_type) \
260 template <> \
261 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE void pscatter<scalar, packet_type>( \
262 typename unpacket_traits<packet_type>::type * to, const packet_type& from, Index stride) { \
263 get_base_packet<packet_type>::set_pscatter(to, from, stride); \
264 }
265
266SYCL_PSCATTER_SPECILIZE(Eigen::half, cl::sycl::cl_half8)
267SYCL_PSCATTER_SPECILIZE(float, cl::sycl::cl_float4)
268SYCL_PSCATTER_SPECILIZE(double, cl::sycl::cl_double2)
269
270#undef SYCL_PSCATTER_SPECILIZE
271
272#define SYCL_PMAD(packet_type) \
273 template <> \
274 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE packet_type pmadd(const packet_type& a, const packet_type& b, \
275 const packet_type& c) { \
276 return cl::sycl::mad(a, b, c); \
277 }
278
279SYCL_PMAD(cl::sycl::cl_half8)
280SYCL_PMAD(cl::sycl::cl_float4)
281SYCL_PMAD(cl::sycl::cl_double2)
282#undef SYCL_PMAD
283
284template <>
285EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half pfirst<cl::sycl::cl_half8>(const cl::sycl::cl_half8& a) {
286 return Eigen::half(a.s0());
287}
288template <>
289EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float pfirst<cl::sycl::cl_float4>(const cl::sycl::cl_float4& a) {
290 return a.x();
291}
292template <>
293EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double pfirst<cl::sycl::cl_double2>(const cl::sycl::cl_double2& a) {
294 return a.x();
295}
296
297template <>
298EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half predux<cl::sycl::cl_half8>(const cl::sycl::cl_half8& a) {
299 return Eigen::half(a.s0() + a.s1() + a.s2() + a.s3() + a.s4() + a.s5() + a.s6() + a.s7());
300}
301
302template <>
303EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float predux<cl::sycl::cl_float4>(const cl::sycl::cl_float4& a) {
304 return a.x() + a.y() + a.z() + a.w();
305}
306
307template <>
308EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double predux<cl::sycl::cl_double2>(const cl::sycl::cl_double2& a) {
309 return a.x() + a.y();
310}
311
312template <>
313EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE bool predux_any(const cl::sycl::cl_half8& a) {
314 return cl::sycl::any(a.template as<cl::sycl::cl_short8>() != cl::sycl::cl_short8(0));
315}
316
317template <>
318EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE bool predux_any(const cl::sycl::cl_float4& a) {
319 return cl::sycl::any(a.template as<cl::sycl::cl_int4>() != cl::sycl::cl_int4(0));
320}
321
322template <>
323EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE bool predux_any(const cl::sycl::cl_double2& a) {
324 return cl::sycl::any(a.template as<cl::sycl::cl_long2>() != cl::sycl::cl_long2(0));
325}
326
327template <>
328EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half predux_max<cl::sycl::cl_half8>(const cl::sycl::cl_half8& a) {
329 return Eigen::half(cl::sycl::fmax(cl::sycl::fmax(cl::sycl::fmax(a.s0(), a.s1()), cl::sycl::fmax(a.s2(), a.s3())),
330 cl::sycl::fmax(cl::sycl::fmax(a.s4(), a.s5()), cl::sycl::fmax(a.s6(), a.s7()))));
331}
332template <>
333EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float predux_max<cl::sycl::cl_float4>(const cl::sycl::cl_float4& a) {
334 return cl::sycl::fmax(cl::sycl::fmax(a.x(), a.y()), cl::sycl::fmax(a.z(), a.w()));
335}
336template <>
337EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double predux_max<cl::sycl::cl_double2>(const cl::sycl::cl_double2& a) {
338 return cl::sycl::fmax(a.x(), a.y());
339}
340
341template <>
342EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half predux_min<cl::sycl::cl_half8>(const cl::sycl::cl_half8& a) {
343 return Eigen::half(cl::sycl::fmin(cl::sycl::fmin(cl::sycl::fmin(a.s0(), a.s1()), cl::sycl::fmin(a.s2(), a.s3())),
344 cl::sycl::fmin(cl::sycl::fmin(a.s4(), a.s5()), cl::sycl::fmin(a.s6(), a.s7()))));
345}
346template <>
347EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float predux_min<cl::sycl::cl_float4>(const cl::sycl::cl_float4& a) {
348 return cl::sycl::fmin(cl::sycl::fmin(a.x(), a.y()), cl::sycl::fmin(a.z(), a.w()));
349}
350template <>
351EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double predux_min<cl::sycl::cl_double2>(const cl::sycl::cl_double2& a) {
352 return cl::sycl::fmin(a.x(), a.y());
353}
354
355template <>
356EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Eigen::half predux_mul<cl::sycl::cl_half8>(const cl::sycl::cl_half8& a) {
357 return Eigen::half(a.s0() * a.s1() * a.s2() * a.s3() * a.s4() * a.s5() * a.s6() * a.s7());
358}
359template <>
360EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float predux_mul<cl::sycl::cl_float4>(const cl::sycl::cl_float4& a) {
361 return a.x() * a.y() * a.z() * a.w();
362}
363template <>
364EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double predux_mul<cl::sycl::cl_double2>(const cl::sycl::cl_double2& a) {
365 return a.x() * a.y();
366}
367
368template <>
369EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_half8 pabs<cl::sycl::cl_half8>(const cl::sycl::cl_half8& a) {
370 return cl::sycl::cl_half8(cl::sycl::fabs(a.s0()), cl::sycl::fabs(a.s1()), cl::sycl::fabs(a.s2()),
371 cl::sycl::fabs(a.s3()), cl::sycl::fabs(a.s4()), cl::sycl::fabs(a.s5()),
372 cl::sycl::fabs(a.s6()), cl::sycl::fabs(a.s7()));
373}
374template <>
375EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_float4 pabs<cl::sycl::cl_float4>(const cl::sycl::cl_float4& a) {
376 return cl::sycl::cl_float4(cl::sycl::fabs(a.x()), cl::sycl::fabs(a.y()), cl::sycl::fabs(a.z()),
377 cl::sycl::fabs(a.w()));
378}
379template <>
380EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE cl::sycl::cl_double2 pabs<cl::sycl::cl_double2>(const cl::sycl::cl_double2& a) {
381 return cl::sycl::cl_double2(cl::sycl::fabs(a.x()), cl::sycl::fabs(a.y()));
382}
383
384template <typename Packet>
385EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Packet sycl_pcmp_le(const Packet& a, const Packet& b) {
386 return (a <= b).template as<Packet>();
387}
388
389template <typename Packet>
390EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Packet sycl_pcmp_lt(const Packet& a, const Packet& b) {
391 return (a < b).template as<Packet>();
392}
393
394template <typename Packet>
395EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Packet sycl_pcmp_eq(const Packet& a, const Packet& b) {
396 return (a == b).template as<Packet>();
397}
398
399#define SYCL_PCMP(OP, TYPE) \
400 template <> \
401 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE TYPE pcmp_##OP<TYPE>(const TYPE& a, const TYPE& b) { \
402 return sycl_pcmp_##OP<TYPE>(a, b); \
403 }
404
405SYCL_PCMP(le, cl::sycl::cl_half8)
406SYCL_PCMP(lt, cl::sycl::cl_half8)
407SYCL_PCMP(eq, cl::sycl::cl_half8)
408SYCL_PCMP(le, cl::sycl::cl_float4)
409SYCL_PCMP(lt, cl::sycl::cl_float4)
410SYCL_PCMP(eq, cl::sycl::cl_float4)
411SYCL_PCMP(le, cl::sycl::cl_double2)
412SYCL_PCMP(lt, cl::sycl::cl_double2)
413SYCL_PCMP(eq, cl::sycl::cl_double2)
414#undef SYCL_PCMP
415
416EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void ptranspose(PacketBlock<cl::sycl::cl_half8, 8>& kernel) {
417 cl::sycl::cl_half tmp = kernel.packet[0].s1();
418 kernel.packet[0].s1() = kernel.packet[1].s0();
419 kernel.packet[1].s0() = tmp;
420
421 tmp = kernel.packet[0].s2();
422 kernel.packet[0].s2() = kernel.packet[2].s0();
423 kernel.packet[2].s0() = tmp;
424
425 tmp = kernel.packet[0].s3();
426 kernel.packet[0].s3() = kernel.packet[3].s0();
427 kernel.packet[3].s0() = tmp;
428
429 tmp = kernel.packet[0].s4();
430 kernel.packet[0].s4() = kernel.packet[4].s0();
431 kernel.packet[4].s0() = tmp;
432
433 tmp = kernel.packet[0].s5();
434 kernel.packet[0].s5() = kernel.packet[5].s0();
435 kernel.packet[5].s0() = tmp;
436
437 tmp = kernel.packet[0].s6();
438 kernel.packet[0].s6() = kernel.packet[6].s0();
439 kernel.packet[6].s0() = tmp;
440
441 tmp = kernel.packet[0].s7();
442 kernel.packet[0].s7() = kernel.packet[7].s0();
443 kernel.packet[7].s0() = tmp;
444
445 tmp = kernel.packet[1].s2();
446 kernel.packet[1].s2() = kernel.packet[2].s1();
447 kernel.packet[2].s1() = tmp;
448
449 tmp = kernel.packet[1].s3();
450 kernel.packet[1].s3() = kernel.packet[3].s1();
451 kernel.packet[3].s1() = tmp;
452
453 tmp = kernel.packet[1].s4();
454 kernel.packet[1].s4() = kernel.packet[4].s1();
455 kernel.packet[4].s1() = tmp;
456
457 tmp = kernel.packet[1].s5();
458 kernel.packet[1].s5() = kernel.packet[5].s1();
459 kernel.packet[5].s1() = tmp;
460
461 tmp = kernel.packet[1].s6();
462 kernel.packet[1].s6() = kernel.packet[6].s1();
463 kernel.packet[6].s1() = tmp;
464
465 tmp = kernel.packet[1].s7();
466 kernel.packet[1].s7() = kernel.packet[7].s1();
467 kernel.packet[7].s1() = tmp;
468
469 tmp = kernel.packet[2].s3();
470 kernel.packet[2].s3() = kernel.packet[3].s2();
471 kernel.packet[3].s2() = tmp;
472
473 tmp = kernel.packet[2].s4();
474 kernel.packet[2].s4() = kernel.packet[4].s2();
475 kernel.packet[4].s2() = tmp;
476
477 tmp = kernel.packet[2].s5();
478 kernel.packet[2].s5() = kernel.packet[5].s2();
479 kernel.packet[5].s2() = tmp;
480
481 tmp = kernel.packet[2].s6();
482 kernel.packet[2].s6() = kernel.packet[6].s2();
483 kernel.packet[6].s2() = tmp;
484
485 tmp = kernel.packet[2].s7();
486 kernel.packet[2].s7() = kernel.packet[7].s2();
487 kernel.packet[7].s2() = tmp;
488
489 tmp = kernel.packet[3].s4();
490 kernel.packet[3].s4() = kernel.packet[4].s3();
491 kernel.packet[4].s3() = tmp;
492
493 tmp = kernel.packet[3].s5();
494 kernel.packet[3].s5() = kernel.packet[5].s3();
495 kernel.packet[5].s3() = tmp;
496
497 tmp = kernel.packet[3].s6();
498 kernel.packet[3].s6() = kernel.packet[6].s3();
499 kernel.packet[6].s3() = tmp;
500
501 tmp = kernel.packet[3].s7();
502 kernel.packet[3].s7() = kernel.packet[7].s3();
503 kernel.packet[7].s3() = tmp;
504
505 tmp = kernel.packet[4].s5();
506 kernel.packet[4].s5() = kernel.packet[5].s4();
507 kernel.packet[5].s4() = tmp;
508
509 tmp = kernel.packet[4].s6();
510 kernel.packet[4].s6() = kernel.packet[6].s4();
511 kernel.packet[6].s4() = tmp;
512
513 tmp = kernel.packet[4].s7();
514 kernel.packet[4].s7() = kernel.packet[7].s4();
515 kernel.packet[7].s4() = tmp;
516
517 tmp = kernel.packet[5].s6();
518 kernel.packet[5].s6() = kernel.packet[6].s5();
519 kernel.packet[6].s5() = tmp;
520
521 tmp = kernel.packet[5].s7();
522 kernel.packet[5].s7() = kernel.packet[7].s5();
523 kernel.packet[7].s5() = tmp;
524
525 tmp = kernel.packet[6].s7();
526 kernel.packet[6].s7() = kernel.packet[7].s6();
527 kernel.packet[7].s6() = tmp;
528}
529
530EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void ptranspose(PacketBlock<cl::sycl::cl_float4, 4>& kernel) {
531 float tmp = kernel.packet[0].y();
532 kernel.packet[0].y() = kernel.packet[1].x();
533 kernel.packet[1].x() = tmp;
534
535 tmp = kernel.packet[0].z();
536 kernel.packet[0].z() = kernel.packet[2].x();
537 kernel.packet[2].x() = tmp;
538
539 tmp = kernel.packet[0].w();
540 kernel.packet[0].w() = kernel.packet[3].x();
541 kernel.packet[3].x() = tmp;
542
543 tmp = kernel.packet[1].z();
544 kernel.packet[1].z() = kernel.packet[2].y();
545 kernel.packet[2].y() = tmp;
546
547 tmp = kernel.packet[1].w();
548 kernel.packet[1].w() = kernel.packet[3].y();
549 kernel.packet[3].y() = tmp;
550
551 tmp = kernel.packet[2].w();
552 kernel.packet[2].w() = kernel.packet[3].z();
553 kernel.packet[3].z() = tmp;
554}
555
556EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE void ptranspose(PacketBlock<cl::sycl::cl_double2, 2>& kernel) {
557 double tmp = kernel.packet[0].y();
558 kernel.packet[0].y() = kernel.packet[1].x();
559 kernel.packet[1].x() = tmp;
560}
561
562#endif // SYCL_DEVICE_ONLY
563
564} // end namespace internal
565
566} // end namespace Eigen
567
568#endif // EIGEN_PACKET_MATH_SYCL_H