DPC++ Runtime
Runtime libraries for oneAPI DPC++
builtins_helper.hpp
Go to the documentation of this file.
1 //==----------- builtins_helper.hpp - SYCL built-in helper ----------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #pragma once
10 
11 #include <sycl/detail/export.hpp>
13 #include <sycl/exception.hpp>
14 #include <sycl/pointers.hpp>
15 #include <sycl/types.hpp>
16 
17 // TODO Decide whether to mark functions with this attribute.
18 #define __NOEXC /*noexcept*/
19 
20 #define __MAKE_1V(Fun, Call, N, Ret, Arg1) \
21  __SYCL_EXPORT sycl::vec<Ret, N> Fun __NOEXC(sycl::vec<Arg1, N> x) { \
22  sycl::vec<Ret, N> r; \
23  detail::helper<N - 1>().run_1v( \
24  r, [](Arg1 x) { return __host_std::Call(x); }, x); \
25  return r; \
26  }
27 
28 #define __MAKE_1V_2V(Fun, Call, N, Ret, Arg1, Arg2) \
29  __SYCL_EXPORT sycl::vec<Ret, N> Fun __NOEXC(sycl::vec<Arg1, N> x, \
30  sycl::vec<Arg2, N> y) { \
31  sycl::vec<Ret, N> r; \
32  detail::helper<N - 1>().run_1v_2v( \
33  r, [](Arg1 x, Arg2 y) { return __host_std::Call(x, y); }, x, y); \
34  return r; \
35  }
36 
37 #define __MAKE_1V_2V_RS(Fun, Call, N, Ret, Arg1, Arg2) \
38  __SYCL_EXPORT Ret Fun __NOEXC(sycl::vec<Arg1, N> x, sycl::vec<Arg2, N> y) { \
39  Ret r = Ret(); \
40  detail::helper<N - 1>().run_1v_2v_rs( \
41  r, [](Ret &r, Arg1 x, Arg2 y) { return __host_std::Call(r, x, y); }, \
42  x, y); \
43  return r; \
44  }
45 
46 #define __MAKE_1V_RS(Fun, Call, N, Ret, Arg1) \
47  __SYCL_EXPORT Ret Fun __NOEXC(sycl::vec<Arg1, N> x) { \
48  Ret r = Ret(); \
49  detail::helper<N - 1>().run_1v_rs( \
50  r, [](Ret &r, Arg1 x) { return __host_std::Call(r, x); }, x); \
51  return r; \
52  }
53 
54 #define __MAKE_1V_2V_3V(Fun, Call, N, Ret, Arg1, Arg2, Arg3) \
55  __SYCL_EXPORT sycl::vec<Ret, N> Fun __NOEXC( \
56  sycl::vec<Arg1, N> x, sycl::vec<Arg2, N> y, sycl::vec<Arg3, N> z) { \
57  sycl::vec<Ret, N> r; \
58  detail::helper<N - 1>().run_1v_2v_3v( \
59  r, [](Arg1 x, Arg2 y, Arg3 z) { return __host_std::Call(x, y, z); }, \
60  x, y, z); \
61  return r; \
62  }
63 
64 #define __MAKE_1V_2S_3S(Fun, N, Ret, Arg1, Arg2, Arg3) \
65  __SYCL_EXPORT sycl::vec<Ret, N> Fun __NOEXC(sycl::vec<Arg1, N> x, Arg2 y, \
66  Arg3 z) { \
67  sycl::vec<Ret, N> r; \
68  detail::helper<N - 1>().run_1v_2s_3s( \
69  r, [](Arg1 x, Arg2 y, Arg3 z) { return __host_std::Fun(x, y, z); }, x, \
70  y, z); \
71  return r; \
72  }
73 
74 #define __MAKE_1V_2S(Fun, N, Ret, Arg1, Arg2) \
75  __SYCL_EXPORT sycl::vec<Ret, N> Fun __NOEXC(sycl::vec<Arg1, N> x, Arg2 y) { \
76  sycl::vec<Ret, N> r; \
77  detail::helper<N - 1>().run_1v_2s( \
78  r, [](Arg1 x, Arg2 y) { return __host_std::Fun(x, y); }, x, y); \
79  return r; \
80  }
81 
82 #define __MAKE_SR_1V_AND(Fun, Call, N, Ret, Arg1) \
83  __SYCL_EXPORT Ret Fun __NOEXC(sycl::vec<Arg1, N> x) { \
84  Ret r; \
85  detail::helper<N - 1>().run_1v_sr_and( \
86  r, [](Arg1 x) { return __host_std::Call(x); }, x); \
87  return r; \
88  }
89 
90 #define __MAKE_SR_1V_OR(Fun, Call, N, Ret, Arg1) \
91  __SYCL_EXPORT Ret Fun __NOEXC(sycl::vec<Arg1, N> x) { \
92  Ret r; \
93  detail::helper<N - 1>().run_1v_sr_or( \
94  r, [](Arg1 x) { return __host_std::Call(x); }, x); \
95  return r; \
96  }
97 
98 #define __MAKE_1V_2P(Fun, N, Ret, Arg1, Arg2) \
99  __SYCL_EXPORT sycl::vec<Ret, N> Fun __NOEXC(sycl::vec<Arg1, N> x, \
100  sycl::vec<Arg2, N> *y) { \
101  sycl::vec<Ret, N> r; \
102  detail::helper<N - 1>().run_1v_2p( \
103  r, [](Arg1 x, Arg2 *y) { return __host_std::Fun(x, y); }, x, y); \
104  return r; \
105  }
106 
107 #define __MAKE_1V_2V_3P(Fun, N, Ret, Arg1, Arg2, Arg3) \
108  __SYCL_EXPORT sycl::vec<Ret, N> Fun __NOEXC( \
109  sycl::vec<Arg1, N> x, sycl::vec<Arg2, N> y, sycl::vec<Arg3, N> *z) { \
110  sycl::vec<Ret, N> r; \
111  detail::helper<N - 1>().run_1v_2v_3p( \
112  r, [](Arg1 x, Arg2 y, Arg3 *z) { return __host_std::Fun(x, y, z); }, \
113  x, y, z); \
114  return r; \
115  }
116 
117 #define MAKE_1V(Fun, Ret, Arg1) MAKE_1V_FUNC(Fun, Fun, Ret, Arg1)
118 
119 #define MAKE_1V_FUNC(Fun, Call, Ret, Arg1) \
120  __MAKE_1V(Fun, Call, 1, Ret, Arg1) \
121  __MAKE_1V(Fun, Call, 2, Ret, Arg1) \
122  __MAKE_1V(Fun, Call, 3, Ret, Arg1) \
123  __MAKE_1V(Fun, Call, 4, Ret, Arg1) \
124  __MAKE_1V(Fun, Call, 8, Ret, Arg1) \
125  __MAKE_1V(Fun, Call, 16, Ret, Arg1)
126 
127 #define MAKE_1V_2V(Fun, Ret, Arg1, Arg2) \
128  MAKE_1V_2V_FUNC(Fun, Fun, Ret, Arg1, Arg2)
129 
130 #define MAKE_1V_2V_FUNC(Fun, Call, Ret, Arg1, Arg2) \
131  __MAKE_1V_2V(Fun, Call, 1, Ret, Arg1, Arg2) \
132  __MAKE_1V_2V(Fun, Call, 2, Ret, Arg1, Arg2) \
133  __MAKE_1V_2V(Fun, Call, 3, Ret, Arg1, Arg2) \
134  __MAKE_1V_2V(Fun, Call, 4, Ret, Arg1, Arg2) \
135  __MAKE_1V_2V(Fun, Call, 8, Ret, Arg1, Arg2) \
136  __MAKE_1V_2V(Fun, Call, 16, Ret, Arg1, Arg2)
137 
138 #define MAKE_1V_2V_3V(Fun, Ret, Arg1, Arg2, Arg3) \
139  MAKE_1V_2V_3V_FUNC(Fun, Fun, Ret, Arg1, Arg2, Arg3)
140 
141 #define MAKE_1V_2V_3V_FUNC(Fun, Call, Ret, Arg1, Arg2, Arg3) \
142  __MAKE_1V_2V_3V(Fun, Call, 1, Ret, Arg1, Arg2, Arg3) \
143  __MAKE_1V_2V_3V(Fun, Call, 2, Ret, Arg1, Arg2, Arg3) \
144  __MAKE_1V_2V_3V(Fun, Call, 3, Ret, Arg1, Arg2, Arg3) \
145  __MAKE_1V_2V_3V(Fun, Call, 4, Ret, Arg1, Arg2, Arg3) \
146  __MAKE_1V_2V_3V(Fun, Call, 8, Ret, Arg1, Arg2, Arg3) \
147  __MAKE_1V_2V_3V(Fun, Call, 16, Ret, Arg1, Arg2, Arg3)
148 
149 #define MAKE_SC_1V_2V_3V(Fun, Ret, Arg1, Arg2, Arg3) \
150  MAKE_SC_3ARG(Fun, Ret, Arg1, Arg2, Arg3) \
151  MAKE_1V_2V_3V_FUNC(Fun, Fun, Ret, Arg1, Arg2, Arg3)
152 
153 #define MAKE_SC_FSC_1V_2V_3V_FV(FunSc, FunV, Ret, Arg1, Arg2, Arg3) \
154  MAKE_SC_3ARG(FunSc, Ret, Arg1, Arg2, Arg3) \
155  MAKE_1V_2V_3V_FUNC(FunSc, FunV, Ret, Arg1, Arg2, Arg3)
156 
157 #define MAKE_SC_3ARG(Fun, Ret, Arg1, Arg2, Arg3) \
158  __SYCL_EXPORT Ret Fun __NOEXC(Arg1 x, Arg2 y, Arg3 z) { \
159  return (Ret)__##Fun(x, y, z); \
160  }
161 
162 #define MAKE_1V_2S(Fun, Ret, Arg1, Arg2) \
163  __MAKE_1V_2S(Fun, 1, Ret, Arg1, Arg2) \
164  __MAKE_1V_2S(Fun, 2, Ret, Arg1, Arg2) \
165  __MAKE_1V_2S(Fun, 3, Ret, Arg1, Arg2) \
166  __MAKE_1V_2S(Fun, 4, Ret, Arg1, Arg2) \
167  __MAKE_1V_2S(Fun, 8, Ret, Arg1, Arg2) \
168  __MAKE_1V_2S(Fun, 16, Ret, Arg1, Arg2)
169 
170 #define MAKE_1V_2S_3S(Fun, Ret, Arg1, Arg2, Arg3) \
171  __MAKE_1V_2S_3S(Fun, 1, Ret, Arg1, Arg2, Arg3) \
172  __MAKE_1V_2S_3S(Fun, 2, Ret, Arg1, Arg2, Arg3) \
173  __MAKE_1V_2S_3S(Fun, 3, Ret, Arg1, Arg2, Arg3) \
174  __MAKE_1V_2S_3S(Fun, 4, Ret, Arg1, Arg2, Arg3) \
175  __MAKE_1V_2S_3S(Fun, 8, Ret, Arg1, Arg2, Arg3) \
176  __MAKE_1V_2S_3S(Fun, 16, Ret, Arg1, Arg2, Arg3)
177 
178 #define MAKE_SR_1V_AND(Fun, Call, Ret, Arg1) \
179  __MAKE_SR_1V_AND(Fun, Call, 1, Ret, Arg1) \
180  __MAKE_SR_1V_AND(Fun, Call, 2, Ret, Arg1) \
181  __MAKE_SR_1V_AND(Fun, Call, 3, Ret, Arg1) \
182  __MAKE_SR_1V_AND(Fun, Call, 4, Ret, Arg1) \
183  __MAKE_SR_1V_AND(Fun, Call, 8, Ret, Arg1) \
184  __MAKE_SR_1V_AND(Fun, Call, 16, Ret, Arg1)
185 
186 #define MAKE_SR_1V_OR(Fun, Call, Ret, Arg1) \
187  __MAKE_SR_1V_OR(Fun, Call, 1, Ret, Arg1) \
188  __MAKE_SR_1V_OR(Fun, Call, 2, Ret, Arg1) \
189  __MAKE_SR_1V_OR(Fun, Call, 3, Ret, Arg1) \
190  __MAKE_SR_1V_OR(Fun, Call, 4, Ret, Arg1) \
191  __MAKE_SR_1V_OR(Fun, Call, 8, Ret, Arg1) \
192  __MAKE_SR_1V_OR(Fun, Call, 16, Ret, Arg1)
193 
194 #define MAKE_1V_2P(Fun, Ret, Arg1, Arg2) \
195  __MAKE_1V_2P(Fun, 1, Ret, Arg1, Arg2) \
196  __MAKE_1V_2P(Fun, 2, Ret, Arg1, Arg2) \
197  __MAKE_1V_2P(Fun, 3, Ret, Arg1, Arg2) \
198  __MAKE_1V_2P(Fun, 4, Ret, Arg1, Arg2) \
199  __MAKE_1V_2P(Fun, 8, Ret, Arg1, Arg2) \
200  __MAKE_1V_2P(Fun, 16, Ret, Arg1, Arg2)
201 
202 #define MAKE_GEO_1V_2V_RS(Fun, Call, Ret, Arg1, Arg2) \
203  __MAKE_1V_2V_RS(Fun, Call, 1, Ret, Arg1, Arg2) \
204  __MAKE_1V_2V_RS(Fun, Call, 2, Ret, Arg1, Arg2) \
205  __MAKE_1V_2V_RS(Fun, Call, 3, Ret, Arg1, Arg2) \
206  __MAKE_1V_2V_RS(Fun, Call, 4, Ret, Arg1, Arg2) \
207  __MAKE_1V_2V_RS(Fun, Call, 8, Ret, Arg1, Arg2) \
208  __MAKE_1V_2V_RS(Fun, Call, 16, Ret, Arg1, Arg2)
209 
210 #define MAKE_1V_2V_3P(Fun, Ret, Arg1, Arg2, Arg3) \
211  __MAKE_1V_2V_3P(Fun, 1, Ret, Arg1, Arg2, Arg3) \
212  __MAKE_1V_2V_3P(Fun, 2, Ret, Arg1, Arg2, Arg3) \
213  __MAKE_1V_2V_3P(Fun, 3, Ret, Arg1, Arg2, Arg3) \
214  __MAKE_1V_2V_3P(Fun, 4, Ret, Arg1, Arg2, Arg3) \
215  __MAKE_1V_2V_3P(Fun, 8, Ret, Arg1, Arg2, Arg3) \
216  __MAKE_1V_2V_3P(Fun, 16, Ret, Arg1, Arg2, Arg3)
217 
218 namespace __host_std {
219 namespace detail {
220 
221 template <int N> struct helper {
222  template <typename Res, typename Op, typename T1>
223  inline void run_1v(Res &r, Op op, T1 x) {
224  helper<N - 1>().run_1v(r, op, x);
225  r.template swizzle<N>() = op(x.template swizzle<N>());
226  }
227 
228  template <typename Res, typename Op, typename T1, typename T2>
229  inline void run_1v_2v(Res &r, Op op, T1 x, T2 y) {
230  helper<N - 1>().run_1v_2v(r, op, x, y);
231  r.template swizzle<N>() =
232  op(x.template swizzle<N>(), y.template swizzle<N>());
233  }
234 
235  template <typename Res, typename Op, typename T1, typename T2>
236  inline void run_1v_2s(Res &r, Op op, T1 x, T2 y) {
237  helper<N - 1>().run_1v_2s(r, op, x, y);
238  r.template swizzle<N>() = op(x.template swizzle<N>(), y);
239  }
240 
241  template <typename Res, typename Op, typename T1, typename T2, typename T3>
242  inline void run_1v_2s_3s(Res &r, Op op, T1 x, T2 y, T3 z) {
243  helper<N - 1>().run_1v_2s_3s(r, op, x, y, z);
244  r.template swizzle<N>() = op(x.template swizzle<N>(), y, z);
245  }
246 
247  template <typename Res, typename Op, typename T1, typename T2>
248  inline void run_1v_2v_rs(Res &r, Op op, T1 x, T2 y) {
249  helper<N - 1>().run_1v_2v_rs(r, op, x, y);
250  op(r, x.template swizzle<N>(), y.template swizzle<N>());
251  }
252 
253  template <typename Res, typename Op, typename T1>
254  inline void run_1v_rs(Res &r, Op op, T1 x) {
255  helper<N - 1>().run_1v_rs(r, op, x);
256  op(r, x.template swizzle<N>());
257  }
258 
259  template <typename Res, typename Op, typename T1, typename T2>
260  inline void run_1v_2p(Res &r, Op op, T1 x, T2 y) {
261  helper<N - 1>().run_1v_2p(r, op, x, y);
262  // TODO avoid creating a temporary variable
264  r.template swizzle<N>() = op(x.template swizzle<N>(), &temp);
265  y->template swizzle<N>() = temp;
266  }
267 
268  template <typename Res, typename Op, typename T1, typename T2, typename T3>
269  inline void run_1v_2v_3p(Res &r, Op op, T1 x, T2 y, T3 z) {
270  helper<N - 1>().run_1v_2v_3p(r, op, x, y, z);
271  // TODO avoid creating a temporary variable
273  r.template swizzle<N>() =
274  op(x.template swizzle<N>(), y.template swizzle<N>(), &temp);
275  z->template swizzle<N>() = temp;
276  }
277 
278  template <typename Res, typename Op, typename T1, typename T2, typename T3>
279  inline void run_1v_2v_3v(Res &r, Op op, T1 x, T2 y, T3 z) {
280  helper<N - 1>().run_1v_2v_3v(r, op, x, y, z);
281  r.template swizzle<N>() =
282  op(x.template swizzle<N>(), y.template swizzle<N>(),
283  z.template swizzle<N>());
284  }
285 
286  template <typename Res, typename Op, typename T1>
287  inline void run_1v_sr_or(Res &r, Op op, T1 x) {
288  helper<N - 1>().run_1v_sr_or(r, op, x);
289  r = (op(x.template swizzle<N>()) || r);
290  }
291 
292  template <typename Res, typename Op, typename T1>
293  inline void run_1v_sr_and(Res &r, Op op, T1 x) {
294  helper<N - 1>().run_1v_sr_and(r, op, x);
295  r = (op(x.template swizzle<N>()) && r);
296  }
297 };
298 
299 template <> struct helper<0> {
300  template <typename Res, typename Op, typename T1>
301  inline void run_1v(Res &r, Op op, T1 x) {
302  r.template swizzle<0>() = op(x.template swizzle<0>());
303  }
304 
305  template <typename Res, typename Op, typename T1, typename T2>
306  inline void run_1v_2v(Res &r, Op op, T1 x, T2 y) {
307  r.template swizzle<0>() =
308  op(x.template swizzle<0>(), y.template swizzle<0>());
309  }
310 
311  template <typename Res, typename Op, typename T1, typename T2>
312  inline void run_1v_2s(Res &r, Op op, T1 x, T2 y) {
313  r.template swizzle<0>() = op(x.template swizzle<0>(), y);
314  }
315 
316  template <typename Res, typename Op, typename T1, typename T2, typename T3>
317  inline void run_1v_2s_3s(Res &r, Op op, T1 x, T2 y, T3 z) {
318  r.template swizzle<0>() = op(x.template swizzle<0>(), y, z);
319  }
320 
321  template <typename Res, typename Op, typename T1, typename T2>
322  inline void run_1v_2v_rs(Res &r, Op op, T1 x, T2 y) {
323  op(r, x.template swizzle<0>(), y.template swizzle<0>());
324  }
325 
326  template <typename Res, typename Op, typename T1>
327  inline void run_1v_rs(Res &r, Op op, T1 x) {
328  op(r, x.template swizzle<0>());
329  }
330 
331  template <typename Res, typename Op, typename T1, typename T2>
332  inline void run_1v_2p(Res &r, Op op, T1 x, T2 y) {
333  // TODO avoid creating a temporary variable
335  r.template swizzle<0>() = op(x.template swizzle<0>(), &temp);
336  y->template swizzle<0>() = temp;
337  }
338 
339  template <typename Res, typename Op, typename T1, typename T2, typename T3>
340  inline void run_1v_2v_3p(Res &r, Op op, T1 x, T2 y, T3 z) {
341  // TODO avoid creating a temporary variable
343  r.template swizzle<0>() =
344  op(x.template swizzle<0>(), y.template swizzle<0>(), &temp);
345  z->template swizzle<0>() = temp;
346  }
347 
348  template <typename Res, typename Op, typename T1, typename T2, typename T3>
349  inline void run_1v_2v_3v(Res &r, Op op, T1 x, T2 y, T3 z) {
350  r.template swizzle<0>() =
351  op(x.template swizzle<0>(), y.template swizzle<0>(),
352  z.template swizzle<0>());
353  }
354 
355  template <typename Res, typename Op, typename T1>
356  inline void run_1v_sr_or(Res &r, Op op, T1 x) {
357  r = op(x.template swizzle<0>());
358  }
359 
360  template <typename Res, typename Op, typename T1>
361  inline void run_1v_sr_and(Res &r, Op op, T1 x) {
362  r = op(x.template swizzle<0>());
363  }
364 };
365 
366 } // namespace detail
367 } // namespace __host_std
std::conditional_t< std::is_same_v< ElementType, half >, sycl::detail::half_impl::BIsRepresentationT, ElementType > element_type
Definition: multi_ptr.hpp:752
void run_1v_2s(Res &r, Op op, T1 x, T2 y)
void run_1v_2v_rs(Res &r, Op op, T1 x, T2 y)
void run_1v_2v_3v(Res &r, Op op, T1 x, T2 y, T3 z)
void run_1v_rs(Res &r, Op op, T1 x)
void run_1v_2v(Res &r, Op op, T1 x, T2 y)
void run_1v_2p(Res &r, Op op, T1 x, T2 y)
void run_1v(Res &r, Op op, T1 x)
void run_1v_2v_3p(Res &r, Op op, T1 x, T2 y, T3 z)
void run_1v_2s_3s(Res &r, Op op, T1 x, T2 y, T3 z)
void run_1v_sr_and(Res &r, Op op, T1 x)
void run_1v_sr_or(Res &r, Op op, T1 x)
void run_1v_rs(Res &r, Op op, T1 x)
void run_1v_2s(Res &r, Op op, T1 x, T2 y)
void run_1v_sr_or(Res &r, Op op, T1 x)
void run_1v_sr_and(Res &r, Op op, T1 x)
void run_1v_2p(Res &r, Op op, T1 x, T2 y)
void run_1v_2v(Res &r, Op op, T1 x, T2 y)
void run_1v_2s_3s(Res &r, Op op, T1 x, T2 y, T3 z)
void run_1v_2v_3v(Res &r, Op op, T1 x, T2 y, T3 z)
void run_1v(Res &r, Op op, T1 x)
void run_1v_2v_3p(Res &r, Op op, T1 x, T2 y, T3 z)
void run_1v_2v_rs(Res &r, Op op, T1 x, T2 y)