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