DPC++ Runtime
Runtime libraries for oneAPI DPC++
math_intrin.hpp
Go to the documentation of this file.
1 //==------------ math_intrin.hpp - DPC++ Explicit SIMD API -----------------==//
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 // Declares Explicit SIMD math intrinsics used to implement working with
9 // the SIMD classes objects.
10 //===----------------------------------------------------------------------===//
11 
12 #pragma once
13 
15 
16 #include <sycl/builtins.hpp>
22 
23 #include <cstdint>
24 
25 #define __ESIMD_raw_vec_t(T, SZ) \
26  __ESIMD_DNS::vector_type_t<__ESIMD_DNS::__raw_t<T>, SZ>
27 #define __ESIMD_cpp_vec_t(T, SZ) \
28  __ESIMD_DNS::vector_type_t<__ESIMD_DNS::__cpp_t<T>, SZ>
29 
30 // saturation intrinsics
31 template <typename T0, typename T1, int SZ>
32 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
33  __esimd_sat(__ESIMD_raw_vec_t(T1, SZ) src);
34 
35 template <typename T0, typename T1, int SZ>
36 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
37  __esimd_fptoui_sat(__ESIMD_raw_vec_t(T1, SZ) src);
38 
39 template <typename T0, typename T1, int SZ>
40 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
41  __esimd_fptosi_sat(__ESIMD_raw_vec_t(T1, SZ) src);
42 
43 template <typename T0, typename T1, int SZ>
44 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
45  __esimd_uutrunc_sat(__ESIMD_raw_vec_t(T1, SZ) src);
46 
47 template <typename T0, typename T1, int SZ>
48 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
49  __esimd_ustrunc_sat(__ESIMD_raw_vec_t(T1, SZ) src);
50 
51 template <typename T0, typename T1, int SZ>
52 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
53  __esimd_sutrunc_sat(__ESIMD_raw_vec_t(T1, SZ) src);
54 
55 template <typename T0, typename T1, int SZ>
56 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
57  __esimd_sstrunc_sat(__ESIMD_raw_vec_t(T1, SZ) src);
58 
59 template <typename T, int SZ>
60 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
61  __esimd_abs(__ESIMD_raw_vec_t(T, SZ) src0);
62 
64 template <typename T, int SZ>
65 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
66  __esimd_fmax(__ESIMD_raw_vec_t(T, SZ) src0, __ESIMD_raw_vec_t(T, SZ) src1);
67 template <typename T, int SZ>
68 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
69  __esimd_umax(__ESIMD_raw_vec_t(T, SZ) src0, __ESIMD_raw_vec_t(T, SZ) src1);
70 template <typename T, int SZ>
71 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
72  __esimd_smax(__ESIMD_raw_vec_t(T, SZ) src0, __ESIMD_raw_vec_t(T, SZ) src1);
73 
75 template <typename T, int SZ>
76 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
77  __esimd_fmin(__ESIMD_raw_vec_t(T, SZ) src0, __ESIMD_raw_vec_t(T, SZ) src1);
78 template <typename T, int SZ>
79 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
80  __esimd_umin(__ESIMD_raw_vec_t(T, SZ) src0, __ESIMD_raw_vec_t(T, SZ) src1);
81 template <typename T, int SZ>
82 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
83  __esimd_smin(__ESIMD_raw_vec_t(T, SZ) src0, __ESIMD_raw_vec_t(T, SZ) src1);
84 
85 template <typename T, int SZ>
86 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<unsigned int, SZ>
87  __esimd_cbit(__ESIMD_raw_vec_t(T, SZ) src0);
88 
89 template <typename T0, int SZ>
90 __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ)
91  __esimd_fbl(__ESIMD_raw_vec_t(T0, SZ) src0);
92 
93 template <typename T0, int SZ>
94 __ESIMD_INTRIN __ESIMD_raw_vec_t(int, SZ)
95  __esimd_sfbh(__ESIMD_raw_vec_t(T0, SZ) src0);
96 
97 template <typename T0, int SZ>
98 __ESIMD_INTRIN __ESIMD_raw_vec_t(uint32_t, SZ)
99  __esimd_ufbh(__ESIMD_raw_vec_t(T0, SZ) src0);
100 
101 #define __ESIMD_UNARY_EXT_MATH_INTRIN(name) \
102  template <class T, int SZ> \
103  __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ) \
104  __esimd_##name(__ESIMD_raw_vec_t(T, SZ) src)
105 
106 __ESIMD_UNARY_EXT_MATH_INTRIN(inv);
107 __ESIMD_UNARY_EXT_MATH_INTRIN(log);
108 __ESIMD_UNARY_EXT_MATH_INTRIN(exp);
109 __ESIMD_UNARY_EXT_MATH_INTRIN(sqrt);
110 __ESIMD_UNARY_EXT_MATH_INTRIN(ieee_sqrt);
111 __ESIMD_UNARY_EXT_MATH_INTRIN(rsqrt);
112 __ESIMD_UNARY_EXT_MATH_INTRIN(sin);
113 __ESIMD_UNARY_EXT_MATH_INTRIN(cos);
114 
115 #undef __ESIMD_UNARY_EXT_MATH_INTRIN
116 
117 template <class T, int SZ>
118 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
119  __esimd_pow(__ESIMD_raw_vec_t(T, SZ) src0, __ESIMD_raw_vec_t(T, SZ) src1);
120 
121 template <class T, int SZ>
122 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
123  __esimd_ieee_div(__ESIMD_raw_vec_t(T, SZ) src0,
124  __ESIMD_raw_vec_t(T, SZ) src1);
125 
126 template <int SZ>
127 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<float, SZ>
128 __esimd_rndd(__ESIMD_DNS::vector_type_t<float, SZ> src0);
129 template <int SZ>
130 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<float, SZ>
131 __esimd_rndu(__ESIMD_DNS::vector_type_t<float, SZ> src0);
132 template <int SZ>
133 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<float, SZ>
134 __esimd_rnde(__ESIMD_DNS::vector_type_t<float, SZ> src0);
135 template <int SZ>
136 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<float, SZ>
137 __esimd_rndz(__ESIMD_DNS::vector_type_t<float, SZ> src0);
138 
139 template <int N>
140 __ESIMD_INTRIN uint32_t
141 __esimd_pack_mask(__ESIMD_DNS::vector_type_t<uint16_t, N> src0);
142 
143 template <int N>
144 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<uint16_t, N>
145 __esimd_unpack_mask(uint32_t src0);
146 
147 template <typename T1, typename T2, typename T3, typename T4, int N>
148 __ESIMD_INTRIN __ESIMD_raw_vec_t(T1, N)
149  __esimd_uudp4a(__ESIMD_raw_vec_t(T2, N) src0, __ESIMD_raw_vec_t(T3, N) src1,
150  __ESIMD_raw_vec_t(T4, N) src2);
151 
152 template <typename T1, typename T2, typename T3, typename T4, int N>
153 __ESIMD_INTRIN __ESIMD_raw_vec_t(T1, N)
154  __esimd_usdp4a(__ESIMD_raw_vec_t(T2, N) src0, __ESIMD_raw_vec_t(T3, N) src1,
155  __ESIMD_raw_vec_t(T4, N) src2);
156 
157 template <typename T1, typename T2, typename T3, typename T4, int N>
158 __ESIMD_INTRIN __ESIMD_raw_vec_t(T1, N)
159  __esimd_sudp4a(__ESIMD_raw_vec_t(T2, N) src0, __ESIMD_raw_vec_t(T3, N) src1,
160  __ESIMD_raw_vec_t(T4, N) src2);
161 
162 template <typename T1, typename T2, typename T3, typename T4, int N>
163 __ESIMD_INTRIN __ESIMD_raw_vec_t(T1, N)
164  __esimd_ssdp4a(__ESIMD_raw_vec_t(T2, N) src0, __ESIMD_raw_vec_t(T3, N) src1,
165  __ESIMD_raw_vec_t(T4, N) src2);
166 
167 template <typename T1, typename T2, typename T3, typename T4, int N>
168 __ESIMD_INTRIN __ESIMD_raw_vec_t(T1, N)
169  __esimd_uudp4a_sat(__ESIMD_raw_vec_t(T2, N) src0,
170  __ESIMD_raw_vec_t(T3, N) src1,
171  __ESIMD_raw_vec_t(T4, N) src2);
172 
173 template <typename T1, typename T2, typename T3, typename T4, int N>
174 __ESIMD_INTRIN __ESIMD_raw_vec_t(T1, N)
175  __esimd_usdp4a_sat(__ESIMD_raw_vec_t(T2, N) src0,
176  __ESIMD_raw_vec_t(T3, N) src1,
177  __ESIMD_raw_vec_t(T4, N) src2);
178 
179 template <typename T1, typename T2, typename T3, typename T4, int N>
180 __ESIMD_INTRIN __ESIMD_raw_vec_t(T1, N)
181  __esimd_sudp4a_sat(__ESIMD_raw_vec_t(T2, N) src0,
182  __ESIMD_raw_vec_t(T3, N) src1,
183  __ESIMD_raw_vec_t(T4, N) src2);
184 
185 template <typename T1, typename T2, typename T3, typename T4, int N>
186 __ESIMD_INTRIN __ESIMD_raw_vec_t(T1, N)
187  __esimd_ssdp4a_sat(__ESIMD_raw_vec_t(T2, N) src0,
188  __ESIMD_raw_vec_t(T3, N) src1,
189  __ESIMD_raw_vec_t(T4, N) src2);
190 
191 #ifdef __SYCL_DEVICE_ONLY__
192 
193 // lane-id for reusing scalar math functions.
194 // Depending upon the SIMT mode(8/16/32), the return value is
195 // in the range of 0-7, 0-15, or 0-31.
196 __ESIMD_INTRIN int __esimd_lane_id();
197 
198 // Wrapper for designating a scalar region of code that will be
199 // vectorized by the backend compiler.
200 #define __ESIMD_SIMT_BEGIN(N, lane) \
201  [&]() SYCL_ESIMD_FUNCTION ESIMD_NOINLINE [[intel::sycl_esimd_vectorize(N)]] { \
202  int lane = __esimd_lane_id();
203 #define __ESIMD_SIMT_END \
204  } \
205  ();
206 
207 #define ESIMD_MATH_INTRINSIC_IMPL(type, func) \
208  template <int SZ> \
209  __ESIMD_INTRIN __ESIMD_raw_vec_t(type, SZ) \
210  ocl_##func(__ESIMD_raw_vec_t(type, SZ) src0) { \
211  __ESIMD_raw_vec_t(type, SZ) retv; \
212  __ESIMD_SIMT_BEGIN(SZ, lane) \
213  retv[lane] = sycl::func(src0[lane]); \
214  __ESIMD_SIMT_END \
215  return retv; \
216  }
217 
219 namespace __ESIMD_DNS {
220 // TODO support half vectors in std sycl math functions.
221 ESIMD_MATH_INTRINSIC_IMPL(float, sin)
222 ESIMD_MATH_INTRINSIC_IMPL(float, cos)
223 ESIMD_MATH_INTRINSIC_IMPL(float, exp)
224 ESIMD_MATH_INTRINSIC_IMPL(float, log)
225 } // namespace __ESIMD_DNS
226 } // __SYCL_INLINE_NAMESPACE(cl)
227 
228 #undef __ESIMD_SIMT_BEGIN
229 #undef __ESIMD_SIMT_END
230 #undef ESIMD_MATH_INTRINSIC_IMPL
231 
232 #else // __SYCL_DEVICE_ONLY__
233 
234 // Typical implementation of a generic intrinsic supporting non-standard
235 // types (half, bfloat*,...) should be like this:
236 // - user type information is encoded in template parameters, but function
237 // parameters and return type are raw types
238 // - before use, parameters are converted to EnclosingCppT
239 // - return value is calculated using the converted parameters,
240 // but before return it is converted back to the user type and is bitcast
241 // (that's what .data() basically does) to the raw type
242 //
243 // template <class T, int SZ>
244 // __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ) __esimd_intrin(
245 // __ESIMD_raw_vec_t(T, SZ) raw_src0, __ESIMD_raw_vec_t(T, SZ) raw_src1) {
246 //
247 // simd<T, SZ> ret;
248 // simd<T, SZ> src0{raw_src0};
249 // simd<T, SZ> src1{raw_src1};
250 // ret = function_of(src0, src1);
251 // return ret.data();
252 //
253 // TODO Not following this approach in some of the intrinsics, and performing
254 // calculations on the raw type will lead to runtime compuation error. A guard
255 // if (__ESIMD_DNS::is_wrapper_elem_type_v<T>) __ESIMD_UNSUPPORTED_ON_HOST;
256 // is temporarily used for now, until wrapper types are supported by these
257 // intrinsics.
258 
259 template <typename T>
260 inline T extract(const uint32_t &width, const uint32_t &offset, uint32_t src,
261  const uint32_t &sign_extend) {
262  uint32_t mask = ((1 << width) - 1) << offset;
263  T ret = (src & mask) >> offset;
264  if (sign_extend) {
265  if ((src >> (offset + width - 1)) & 0x1) {
266  uint32_t sign_extend = ((1 << (32 - width)) - 1) << width;
267  ret = ret | sign_extend;
268  }
269  }
270 
271  return ret;
272 }
273 
274 #define __ESIMD_DEFAULT_HOST_SATURATE_INTRIN(name) \
275  template <typename T0, typename T1, int SZ> \
276  __ESIMD_INTRIN __ESIMD_raw_vec_t(T0, SZ) \
277  __esimd_##name(__ESIMD_raw_vec_t(T1, SZ) src) { \
278  __ESIMD_raw_vec_t(T0, SZ) retv; \
279  for (int i = 0; i < SZ; i++) { \
280  SIMDCF_ELEMENT_SKIP(i); \
281  retv[i] = __ESIMD_EMU_DNS::satur<T0>::template saturate<T1>(src[i], 1); \
282  } \
283  return retv; \
284  }
285 
286 __ESIMD_DEFAULT_HOST_SATURATE_INTRIN(sat)
287 __ESIMD_DEFAULT_HOST_SATURATE_INTRIN(fptoui_sat)
288 __ESIMD_DEFAULT_HOST_SATURATE_INTRIN(fptosi_sat)
289 __ESIMD_DEFAULT_HOST_SATURATE_INTRIN(uutrunc_sat)
290 __ESIMD_DEFAULT_HOST_SATURATE_INTRIN(ustrunc_sat)
291 __ESIMD_DEFAULT_HOST_SATURATE_INTRIN(sutrunc_sat)
292 __ESIMD_DEFAULT_HOST_SATURATE_INTRIN(sstrunc_sat)
293 
294 template <typename T, int SZ>
295 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
296  __esimd_abs(__ESIMD_raw_vec_t(T, SZ) src0) {
297  if (__ESIMD_DNS::is_wrapper_elem_type_v<T>)
298  __ESIMD_UNSUPPORTED_ON_HOST;
299  int i;
300  typename __ESIMD_EMU_DNS::abstype<T>::type ret;
301  __ESIMD_raw_vec_t(T, SZ) retv;
302 
303  for (i = 0; i < SZ; i++) {
304  SIMDCF_ELEMENT_SKIP(i);
305  if (src0[i] < 0) {
306  ret = -(src0[i]);
307  } else {
308  ret = (src0[i]);
309  }
310  retv[i] = ret;
311  }
312  return retv;
313 }
314 
316 template <typename T, int SZ>
317 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
318  __esimd_fmax(__ESIMD_raw_vec_t(T, SZ) src0, __ESIMD_raw_vec_t(T, SZ) src1) {
319  if (__ESIMD_DNS::is_wrapper_elem_type_v<T>)
320  __ESIMD_UNSUPPORTED_ON_HOST;
321  int i;
322  __ESIMD_raw_vec_t(T, SZ) retv;
323 
324  for (i = 0; i < SZ; i++) {
325  SIMDCF_ELEMENT_SKIP(i);
326  if (src0[i] >= src1[i]) {
327  retv[i] = src0[i];
328  } else {
329  retv[i] = src1[i];
330  }
331  }
332 
333  return retv;
334 }
335 
336 template <typename T, int SZ>
337 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
338  __esimd_umax(__ESIMD_raw_vec_t(T, SZ) src0, __ESIMD_raw_vec_t(T, SZ) src1) {
339  if (__ESIMD_DNS::is_wrapper_elem_type_v<T>)
340  __ESIMD_UNSUPPORTED_ON_HOST;
341  int i;
342  __ESIMD_raw_vec_t(T, SZ) retv;
343 
344  for (i = 0; i < SZ; i++) {
345  SIMDCF_ELEMENT_SKIP(i);
346  if (src0[i] >= src1[i]) {
347  retv[i] = src0[i];
348  } else {
349  retv[i] = src1[i];
350  }
351  }
352 
353  return retv;
354 }
355 
356 template <typename T, int SZ>
357 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
358  __esimd_smax(__ESIMD_raw_vec_t(T, SZ) src0, __ESIMD_raw_vec_t(T, SZ) src1) {
359  if (__ESIMD_DNS::is_wrapper_elem_type_v<T>)
360  __ESIMD_UNSUPPORTED_ON_HOST;
361  int i;
362  __ESIMD_raw_vec_t(T, SZ) retv;
363 
364  for (i = 0; i < SZ; i++) {
365  SIMDCF_ELEMENT_SKIP(i);
366  if (src0[i] >= src1[i]) {
367  retv[i] = src0[i];
368  } else {
369  retv[i] = src1[i];
370  }
371  }
372 
373  return retv;
374 }
375 
377 template <typename T, int SZ>
378 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
379  __esimd_fmin(__ESIMD_raw_vec_t(T, SZ) src0, __ESIMD_raw_vec_t(T, SZ) src1) {
380  if (__ESIMD_DNS::is_wrapper_elem_type_v<T>)
381  __ESIMD_UNSUPPORTED_ON_HOST;
382  int i;
383  __ESIMD_raw_vec_t(T, SZ) retv;
384 
385  for (i = 0; i < SZ; i++) {
386  SIMDCF_ELEMENT_SKIP(i);
387  if (src0[i] < src1[i]) {
388  retv[i] = src0[i];
389  } else {
390  retv[i] = src1[i];
391  }
392  }
393 
394  return retv;
395 };
396 
397 template <typename T, int SZ>
398 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
399  __esimd_umin(__ESIMD_raw_vec_t(T, SZ) src0, __ESIMD_raw_vec_t(T, SZ) src1) {
400  if (__ESIMD_DNS::is_wrapper_elem_type_v<T>)
401  __ESIMD_UNSUPPORTED_ON_HOST;
402  int i;
403  __ESIMD_raw_vec_t(T, SZ) retv;
404 
405  for (i = 0; i < SZ; i++) {
406  SIMDCF_ELEMENT_SKIP(i);
407  if (src0[i] < src1[i]) {
408  retv[i] = src0[i];
409  } else {
410  retv[i] = src1[i];
411  }
412  }
413 
414  return retv;
415 }
416 
417 template <typename T, int SZ>
418 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
419  __esimd_smin(__ESIMD_raw_vec_t(T, SZ) src0, __ESIMD_raw_vec_t(T, SZ) src1) {
420  if (__ESIMD_DNS::is_wrapper_elem_type_v<T>)
421  __ESIMD_UNSUPPORTED_ON_HOST;
422  int i;
423  __ESIMD_raw_vec_t(T, SZ) retv;
424 
425  for (i = 0; i < SZ; i++) {
426  SIMDCF_ELEMENT_SKIP(i);
427  if (src0[i] < src1[i]) {
428  retv[i] = src0[i];
429  } else {
430  retv[i] = src1[i];
431  }
432  }
433 
434  return retv;
435 }
436 
437 template <typename T, int SZ>
438 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<unsigned int, SZ>
439 __esimd_cbit(__ESIMD_raw_vec_t(T, SZ) src0) {
440  if (__ESIMD_DNS::is_wrapper_elem_type_v<T>)
441  __ESIMD_UNSUPPORTED_ON_HOST;
442  int i;
443  uint32_t ret;
444  __ESIMD_raw_vec_t(uint32_t, SZ) retv;
445 
446  for (i = 0; i < SZ; i++) {
447  SIMDCF_ELEMENT_SKIP(i);
448  ret = src0[i];
449  uint32_t cnt = 0;
450  for (int j = 0; j < sizeof(T) * 8; j++) {
451  if ((ret & 1u) == 1) {
452  cnt++;
453  }
454  ret = ret >> 1;
455  }
456  retv[i] = cnt;
457  }
458 
459  return retv;
460 }
461 
462 template <typename T, int SZ>
463 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
464  __esimd_fbl(__ESIMD_raw_vec_t(T, SZ) src0) {
465  if (__ESIMD_DNS::is_wrapper_elem_type_v<T>)
466  __ESIMD_UNSUPPORTED_ON_HOST;
467  int i;
468  T ret;
469  __ESIMD_raw_vec_t(T, SZ) retv;
470 
471  for (i = 0; i < SZ; i++) {
472  SIMDCF_ELEMENT_SKIP(i);
473  ret = src0[i];
474  uint32_t cnt = 0;
475  while ((ret & 1u) == 0 && cnt != 32) {
476  cnt++;
477  ret = ret >> 1;
478  }
479  if (src0[i] == 0x0) {
480  retv[i] = 0xFFFFFFFF;
481  } else {
482  retv[i] = cnt;
483  }
484  }
485 
486  return retv;
487 }
488 
489 template <typename T, int SZ>
490 __ESIMD_INTRIN __ESIMD_raw_vec_t(int, SZ)
491  __esimd_sfbh(__ESIMD_raw_vec_t(T, SZ) src0) {
492  if (__ESIMD_DNS::is_wrapper_elem_type_v<T>)
493  __ESIMD_UNSUPPORTED_ON_HOST;
494  int i, cval;
495  int ret;
496  __ESIMD_raw_vec_t(int, SZ) retv;
497 
498  for (i = 0; i < SZ; i++) {
499  SIMDCF_ELEMENT_SKIP(i);
500  ret = src0[i];
501  uint32_t cnt = 0;
502  if (((ret >> 31u) & 1u) == 1) {
503  cval = 1;
504  } else {
505  cval = 0;
506  }
507  while (((ret >> 31u) & 1u) == cval && cnt != 32) {
508  cnt++;
509  ret = ret << 1;
510  }
511 
512  if ((src0[i] == 0xFFFFFFFF) || (src0[i] == 0x00000000)) {
513  retv[i] = 0xFFFFFFFF;
514  } else {
515  retv[i] = cnt;
516  }
517  }
518 
519  return retv;
520 }
521 
522 template <typename T, int SZ>
523 __ESIMD_INTRIN __ESIMD_raw_vec_t(uint32_t, SZ)
524  __esimd_ufbh(__ESIMD_raw_vec_t(T, SZ) src0) {
525  if (__ESIMD_DNS::is_wrapper_elem_type_v<T>)
526  __ESIMD_UNSUPPORTED_ON_HOST;
527  uint32_t ret;
528  __ESIMD_raw_vec_t(uint32_t, SZ) retv;
529 
530  for (int i = 0; i < SZ; i++) {
531  SIMDCF_ELEMENT_SKIP(i);
532  ret = src0[i];
533  uint32_t cnt = 0;
534  while ((ret & (1u << 31u)) == 0 && cnt != 32) {
535  cnt++;
536  ret = ret << 1;
537  }
538  if (src0[i] == 0x00000000) {
539  retv[i] = 0xFFFFFFFF;
540  } else {
541  retv[i] = cnt;
542  }
543  }
544 
545  return retv;
546 }
547 
548 // Host intrinsics are implemented via converting elements to enclosing Cpp
549 // type (always 'float' except ieee_sqrt, which can be 'double'), applying
550 // standard C++ library math function and converting back to the element type.
551 //
552 #define __ESIMD_UNARY_EXT_MATH_HOST_INTRIN(name, formula) \
553  template <class T, int SZ> \
554  __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ) \
555  __esimd_##name(__ESIMD_raw_vec_t(T, SZ) src) { \
556  using CppT = __ESIMD_DNS::__cpp_t<T>; \
557  using CppVecT = __ESIMD_cpp_vec_t(T, SZ); \
558  CppVecT ret_cpp{0}; \
559  CppVecT src_cpp = __ESIMD_DNS::convert_vector<CppT, T, SZ>(src); \
560  \
561  for (int i = 0; i < SZ; i++) { \
562  SIMDCF_ELEMENT_SKIP(i); \
563  ret_cpp[i] = formula; \
564  } \
565  __ESIMD_raw_vec_t(T, SZ) ret = \
566  __ESIMD_DNS::convert_vector<T, CppT, SZ>(ret_cpp); \
567  return ret; \
568  }
569 
570 __ESIMD_UNARY_EXT_MATH_HOST_INTRIN(inv, 1.f / src_cpp[i])
571 __ESIMD_UNARY_EXT_MATH_HOST_INTRIN(log, logf(src_cpp[i]) / logf(2.f))
572 __ESIMD_UNARY_EXT_MATH_HOST_INTRIN(exp, powf(2.f, src_cpp[i]))
573 __ESIMD_UNARY_EXT_MATH_HOST_INTRIN(sqrt, sqrt(src_cpp[i]))
574 __ESIMD_UNARY_EXT_MATH_HOST_INTRIN(ieee_sqrt, sqrt(src_cpp[i]))
575 __ESIMD_UNARY_EXT_MATH_HOST_INTRIN(rsqrt, 1.f / sqrt(src_cpp[i]))
576 __ESIMD_UNARY_EXT_MATH_HOST_INTRIN(sin, sin(src_cpp[i]))
577 __ESIMD_UNARY_EXT_MATH_HOST_INTRIN(cos, cos(src_cpp[i]))
578 
579 #undef __ESIMD_UNARY_EXT_MATH_HOST_INTRIN
580 
581 template <class T, int SZ>
582 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
583  __esimd_pow(__ESIMD_raw_vec_t(T, SZ) src0, __ESIMD_raw_vec_t(T, SZ) src1) {
584  using CppT = __ESIMD_DNS::__cpp_t<T>;
585  using CppVecT = __ESIMD_cpp_vec_t(T, SZ);
586 
587  CppVecT cpp_src0 = __ESIMD_DNS::convert_vector<CppT, T, SZ>(src0);
588  CppVecT cpp_src1 = __ESIMD_DNS::convert_vector<CppT, T, SZ>(src1);
589  CppVecT cpp_res;
590 
591  for (int i = 0; i < SZ; i++) {
592  SIMDCF_ELEMENT_SKIP(i);
593  cpp_res[i] = std::pow(std::fabs(cpp_src0[i]), cpp_src1[i]);
594  }
595  return __ESIMD_DNS::convert_vector<T, CppT, SZ>(cpp_res);
596 }
597 
598 template <class T, int SZ>
599 __ESIMD_INTRIN __ESIMD_raw_vec_t(T, SZ)
600  __esimd_ieee_div(__ESIMD_raw_vec_t(T, SZ) src0,
601  __ESIMD_raw_vec_t(T, SZ) src1) {
602  using CppT = __ESIMD_DNS::__cpp_t<T>;
603  using CppVecT = __ESIMD_cpp_vec_t(T, SZ);
604 
605  CppVecT cpp_src0 = __ESIMD_DNS::convert_vector<CppT, T, SZ>(src0);
606  CppVecT cpp_src1 = __ESIMD_DNS::convert_vector<CppT, T, SZ>(src1);
607  CppVecT cpp_res;
608 
609  for (int i = 0; i < SZ; i += 1) {
610  SIMDCF_ELEMENT_SKIP(i);
611  if (cpp_src1[i] == 0) {
613  cpp_res[i] = (cpp_src0[i] < 0) ? (-INFINITY) : INFINITY;
614  } else {
615  cpp_res[i] = cpp_src0[i] / cpp_src1[i];
616  }
617  }
618  return __ESIMD_DNS::convert_vector<T, CppT, SZ>(cpp_res);
619 }
620 
621 template <int SZ>
622 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<float, SZ>
623 __esimd_rndd(__ESIMD_DNS::vector_type_t<float, SZ> src0) {
624  __ESIMD_DNS::vector_type_t<float, SZ> retv;
625 
626  for (int i = 0; i < SZ; i++) {
627  SIMDCF_ELEMENT_SKIP(i);
628  retv[i] = floor(src0[i]);
629  }
630  return retv;
631 }
632 
633 template <int SZ>
634 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<float, SZ>
635 __esimd_rndu(__ESIMD_DNS::vector_type_t<float, SZ> src0) {
636  __ESIMD_DNS::vector_type_t<float, SZ> retv;
637  int increment;
638 
639  for (int i = 0; i < SZ; i++) {
640  SIMDCF_ELEMENT_SKIP(i);
641  if (src0[i] - floor(src0[i]) > 0.0f) {
642  increment = 1;
643  } else {
644  increment = 0;
645  }
646 
647  retv[i] = floor(src0[i]) + increment;
648  }
649 
650  return retv;
651 }
652 
653 template <int SZ>
654 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<float, SZ>
655 __esimd_rnde(__ESIMD_DNS::vector_type_t<float, SZ> src0) {
656  __ESIMD_DNS::vector_type_t<float, SZ> retv;
657  int increment;
658 
659  for (int i = 0; i < SZ; i++) {
660  SIMDCF_ELEMENT_SKIP(i);
661  if (src0[i] - floor(src0[i]) > 0.5f) {
662  increment = 1;
663  } else if (src0[i] - floor(src0[i]) < 0.5f) {
664  increment = 0;
665  } else {
666  increment = (int(floor(src0[i])) % 2 == 1);
667  }
668 
669  retv[i] = floor(src0[i]) + increment;
670  }
671 
672  return retv;
673 }
674 
675 template <int SZ>
676 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<float, SZ>
677 __esimd_rndz(__ESIMD_DNS::vector_type_t<float, SZ> src0) {
678  __ESIMD_DNS::vector_type_t<float, SZ> retv;
679  int increment;
680 
681  for (int i = 0; i < SZ; i++) {
682  SIMDCF_ELEMENT_SKIP(i);
683  if (fabs(src0[i]) < fabs(floor(src0[i]))) {
684  increment = 1;
685  } else {
686  increment = 0;
687  }
688  retv[i] = floor(src0[i]) + increment;
689  }
690 
691  return retv;
692 }
693 
694 template <int N>
695 __ESIMD_INTRIN uint32_t
696 __esimd_pack_mask(__ESIMD_DNS::vector_type_t<uint16_t, N> src0) {
697  // We don't check the arguments here as this function is only invoked by
698  // wrapper code (which does the checks already)
699  uint32_t retv = 0;
700  for (int i = 0; i < N; i++) {
701  if (src0[i] != 0) {
702  retv |= 0x1 << i;
703  }
704  }
705 
706  return retv;
707 }
708 
709 template <int N>
710 __ESIMD_INTRIN __ESIMD_DNS::vector_type_t<uint16_t, N>
711 __esimd_unpack_mask(uint32_t src0) {
712  __ESIMD_DNS::vector_type_t<uint16_t, N> retv = 0;
713  for (int i = 0; i < N; i++) {
714  if ((src0 >> i) & 0x1) {
715  retv[i] = 1;
716  }
717  }
718  return retv;
719 }
720 
721 template <typename T1, typename T2, typename T3, typename T4, int N>
722 __ESIMD_INTRIN __ESIMD_raw_vec_t(T1, N)
723  __esimd_dp4a(__ESIMD_raw_vec_t(T2, N) src0, __ESIMD_raw_vec_t(T3, N) src1,
724  __ESIMD_raw_vec_t(T4, N) src2) {
725 #define __ESIMD_WR(T) __ESIMD_DNS::is_wrapper_elem_type_v<T>
726  if (__ESIMD_WR(T1) || __ESIMD_WR(T2) || __ESIMD_WR(T3) || __ESIMD_WR(T4))
727  __ESIMD_UNSUPPORTED_ON_HOST;
728 #undef __ESIMD_IS_WR
729  using __ESIMD_EMU_DNS::restype_ex;
730  typename restype_ex<T2, typename restype_ex<T3, T4>::type>::type reta;
731  __ESIMD_raw_vec_t(T1, N) retv;
732 
733  int src1_a, src1_b, src1_c, src1_d, src2_a, src2_b, src2_c, src2_d, ret;
734 
735  uint32_t sat1 =
736  __ESIMD_EMU_DNS::SetSatur<
737  T2, __ESIMD_EMU_DNS::is_inttype<T1>::value>::set() ||
738  __ESIMD_EMU_DNS::SetSatur<
739  T3, __ESIMD_EMU_DNS::is_inttype<T1>::value>::set() ||
740  __ESIMD_EMU_DNS::SetSatur<T4,
741  __ESIMD_EMU_DNS::is_inttype<T1>::value>::set();
742 
743  for (uint32_t i = 0; i < N; i++) {
744 
745  SIMDCF_ELEMENT_SKIP(i);
746 
747  src1_a = extract<short>(8, 0, src1[i], 0);
748  src1_b = extract<short>(8, 8, src1[i], 0);
749  src1_c = extract<short>(8, 16, src1[i], 0);
750  src1_d = extract<short>(8, 24, src1[i], 0);
751  src2_a = extract<short>(8, 0, src2[i], 0);
752  src2_b = extract<short>(8, 8, src2[i], 0);
753  src2_c = extract<short>(8, 16, src2[i], 0);
754  src2_d = extract<short>(8, 24, src2[i], 0);
755 
756  ret = src1_a * src2_a + src1_b * src2_b + src1_c * src2_c + src1_d * src2_d;
757  reta = ret + src0[i];
758  retv[i] = __ESIMD_EMU_DNS::satur<T1>::template saturate(reta, sat1);
759  }
760 
761  return retv;
762 }
763 
764 #endif // #ifdef __SYCL_DEVICE_ONLY__
765 
766 #undef __ESIMD_raw_vec_t
767 #undef __ESIMD_cpp_vec_t
768 
cl::sycl::ext::intel::esimd::saturate
__ESIMD_API std::enable_if_t<!detail::is_generic_floating_point_v< T0 >||std::is_same_v< T1, T0 >, simd< T0, SZ > > saturate(simd< T1, SZ > src)
Conversion of input vector elements of type T1 into vector of elements of type T0 with saturation.
Definition: math.hpp:71
host_util.hpp
builtins.hpp
T
cl::sycl::log
detail::enable_if_t< __FAST_MATH_GENFLOAT(T), T > log(T x) __NOEXC
Definition: builtins.hpp:306
cl::sycl::sqrt
detail::enable_if_t< __FAST_MATH_GENFLOAT(T), T > sqrt(T x) __NOEXC
Definition: builtins.hpp:469
common.hpp
cl::sycl::sin
detail::enable_if_t< __FAST_MATH_GENFLOAT(T), T > sin(T x) __NOEXC
Definition: builtins.hpp:442
cl::sycl::floor
detail::enable_if_t< detail::is_genfloat< T >::value, T > floor(T x) __NOEXC
Definition: builtins.hpp:190
cl::sycl::rsqrt
detail::enable_if_t< __FAST_MATH_GENFLOAT(T), T > rsqrt(T x) __NOEXC
Definition: builtins.hpp:436
cl::sycl::fabs
detail::enable_if_t< detail::is_genfloat< T >::value, T > fabs(T x) __NOEXC
Definition: builtins.hpp:178
types.hpp
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
elem_type_traits.hpp
util.hpp
cl::sycl::cos
detail::enable_if_t< __FAST_MATH_GENFLOAT(T), T > cos(T x) __NOEXC
Definition: builtins.hpp:124
cl::sycl::pow
detail::enable_if_t< detail::is_genfloat< T >::value, T > pow(T x, T y) __NOEXC
Definition: builtins.hpp:378
cl::sycl::exp
detail::enable_if_t< __FAST_MATH_GENFLOAT(T), T > exp(T x) __NOEXC
Definition: builtins.hpp:154
cl::sycl::ext::intel::esimd::inv
__ESIMD_API simd< T, N > inv(simd< T, N > src, Sat sat={})
Inversion - calculates (1/x).
Definition: math.hpp:374
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:11