DPC++ Runtime
Runtime libraries for oneAPI DPC++
host_util.hpp
Go to the documentation of this file.
1 //==-------------------------- host_util.hpp -------------------------------==//
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 // Utility functions and definitions for implementing ESIMD intrinsics on host
9 //===----------------------------------------------------------------------===//
10 
11 #pragma once
12 
14 
15 #ifndef __SYCL_DEVICE_ONLY__
16 
17 #include <assert.h>
18 #include <limits>
19 
21 
22 #define SIMDCF_ELEMENT_SKIP(i)
23 
25 namespace sycl {
26 namespace detail {
27 namespace half_impl {
28 class half;
29 } // namespace half_impl
30 } // namespace detail
31 } // namespace sycl
32 } // __SYCL_INLINE_NAMESPACE(cl)
33 
35 namespace __ESIMD_EMU_DNS {
36 
38 constexpr int sat_is_on = 1;
39 
40 static long long abs(long long a) {
41  if (a < 0) {
42  return -a;
43  } else {
44  return a;
45  }
46 }
47 
48 template <typename RT, class SFINAE = void> struct satur;
49 
50 template <typename RT>
51 struct satur<RT, std::enable_if_t<std::is_integral_v<RT>>> {
52  static_assert(!__ESIMD_DNS::is_wrapper_elem_type_v<RT>);
53 
54  template <typename T> static RT saturate(const T val, const int flags) {
55  if ((flags & sat_is_on) == 0) {
56  return (RT)val;
57  }
58 
59  // min/max can be macros on Windows, so wrap them into parens to avoid their
60  // expansion
61  const RT t_max = (std::numeric_limits<RT>::max)();
62  const RT t_min = (std::numeric_limits<RT>::min)();
63 
64  if (val > t_max) {
65  return t_max;
66  } else if ((val >= 0) && (t_min < 0)) {
67  // RT is "signed" if t_min < 0
68  // when comparing a signed and a unsigned variable, the signed one cast to
69  // unsigned first.
70  return (RT)val;
71  } else if (val < t_min) {
72  return t_min;
73  } else {
74  return (RT)val;
75  }
76  }
77 };
78 
79 // Host implemenation of saturation for FP types, including non-standarad
80 // wrapper types such as sycl::half. Template parameters are defined in terms
81 // of user-level types (sycl::half), function parameter and return types -
82 // in terms of raw bit representation type(_Float16 for half on device).
83 template <class Tdst>
84 struct satur<Tdst,
85  std::enable_if_t<__ESIMD_DNS::is_generic_floating_point_v<Tdst>>> {
86  template <typename Tsrc>
87  static __ESIMD_DNS::__raw_t<Tdst>
88  saturate(const __ESIMD_DNS::__raw_t<Tsrc> raw_src, const int flags) {
89  Tsrc src = __ESIMD_DNS::bitcast_to_wrapper_type<Tsrc>(raw_src);
90 
91  // perform comparison on user type!
92  if ((flags & sat_is_on) == 0 || (src >= 0 && src <= 1)) {
93  // convert_scalar accepts/returns user types - need to bitcast
94  Tdst dst = __ESIMD_DNS::convert_scalar<Tdst, Tsrc>(src);
95  return __ESIMD_DNS::bitcast_to_raw_type<Tdst>(dst);
96  }
97  if (src < 0) {
98  return __ESIMD_DNS::bitcast_to_raw_type<Tdst>(Tdst{0});
99  }
100  assert(src > 1);
101  return __ESIMD_DNS::bitcast_to_raw_type<Tdst>(Tdst{1});
102  }
103 };
104 
105 template <typename T1, bool B> struct SetSatur {
106  static unsigned int set() { return 0; }
107 };
108 
109 template <> struct SetSatur<float, true> {
110  static unsigned int set() { return sat_is_on; }
111 };
112 
113 template <> struct SetSatur<double, true> {
114  static unsigned int set() { return sat_is_on; }
115 };
116 
117 // TODO replace restype_ex with detail::computation_type_t and represent half
118 // as sycl::half rather than 'using half = sycl::detail::half_impl::half;'
119 // above
120 
121 // used for intermediate type in dp4a emulation
122 template <typename T1, typename T2> struct restype_ex {
123 private:
124  restype_ex();
125 };
126 
127 template <> struct restype_ex<char, char> {
128  using type = int;
129 };
130 template <> struct restype_ex<char, unsigned char> {
131  using type = int;
132 };
133 template <> struct restype_ex<char, short> {
134  using type = int;
135 };
136 template <> struct restype_ex<char, unsigned short> {
137  using type = int;
138 };
139 template <> struct restype_ex<char, int> {
140  using type = long long;
141 };
142 template <> struct restype_ex<char, unsigned int> {
143  using type = long long;
144 };
145 template <> struct restype_ex<char, half> {
146  using type = half;
147 };
148 template <> struct restype_ex<char, float> {
149  using type = float;
150 };
151 template <> struct restype_ex<char, double> {
152  using type = double;
153 };
154 
155 template <> struct restype_ex<unsigned char, char> {
156  using type = int;
157 };
158 template <> struct restype_ex<unsigned char, unsigned char> {
159  using type = int;
160 };
161 template <> struct restype_ex<unsigned char, short> {
162  using type = int;
163 };
164 template <> struct restype_ex<unsigned char, unsigned short> {
165  using type = int;
166 };
167 template <> struct restype_ex<unsigned char, int> {
168  using type = long long;
169 };
170 template <> struct restype_ex<unsigned char, unsigned int> {
171  using type = long long;
172 };
173 template <> struct restype_ex<unsigned char, half> {
174  using type = half;
175 };
176 template <> struct restype_ex<unsigned char, float> {
177  using type = float;
178 };
179 template <> struct restype_ex<unsigned char, double> {
180  using type = double;
181 };
182 template <> struct restype_ex<unsigned char, long long> {
183  using type = long long;
184 };
185 template <> struct restype_ex<unsigned char, unsigned long long> {
186  using type = long long;
187 };
188 
189 template <> struct restype_ex<short, char> {
190  using type = int;
191 };
192 template <> struct restype_ex<short, unsigned char> {
193  using type = int;
194 };
195 template <> struct restype_ex<short, short> {
196  using type = int;
197 };
198 template <> struct restype_ex<short, unsigned short> {
199  using type = int;
200 };
201 template <> struct restype_ex<short, int> {
202  using type = long long;
203 };
204 template <> struct restype_ex<short, unsigned int> {
205  using type = long long;
206 };
207 template <> struct restype_ex<short, half> {
208  using type = half;
209 };
210 template <> struct restype_ex<short, float> {
211  using type = float;
212 };
213 template <> struct restype_ex<short, double> {
214  using type = double;
215 };
216 template <> struct restype_ex<short, long long> {
217  using type = long long;
218 };
219 template <> struct restype_ex<short, unsigned long long> {
220  using type = long long;
221 };
222 
223 template <> struct restype_ex<unsigned short, char> {
224  using type = int;
225 };
226 template <> struct restype_ex<unsigned short, unsigned char> {
227  using type = int;
228 };
229 template <> struct restype_ex<unsigned short, short> {
230  using type = int;
231 };
232 template <> struct restype_ex<unsigned short, unsigned short> {
233  using type = int;
234 };
235 template <> struct restype_ex<unsigned short, int> {
236  using type = long long;
237 };
238 template <> struct restype_ex<unsigned short, unsigned int> {
239  using type = long long;
240 };
241 template <> struct restype_ex<unsigned short, half> {
242  using type = half;
243 };
244 template <> struct restype_ex<unsigned short, float> {
245  using type = float;
246 };
247 template <> struct restype_ex<unsigned short, double> {
248  using type = double;
249 };
250 template <> struct restype_ex<unsigned short, long long> {
251  using type = long long;
252 };
253 template <> struct restype_ex<unsigned short, unsigned long long> {
254  using type = long long;
255 };
256 
257 template <> struct restype_ex<int, char> {
258  using type = long long;
259 };
260 template <> struct restype_ex<int, unsigned char> {
261  using type = long long;
262 };
263 template <> struct restype_ex<int, short> {
264  using type = long long;
265 };
266 template <> struct restype_ex<int, unsigned short> {
267  using type = long long;
268 };
269 template <> struct restype_ex<int, int> {
270  using type = long long;
271 };
272 template <> struct restype_ex<int, unsigned int> {
273  using type = long long;
274 };
275 template <> struct restype_ex<int, half> {
276  using type = half;
277 };
278 template <> struct restype_ex<int, float> {
279  using type = float;
280 };
281 template <> struct restype_ex<int, double> {
282  using type = double;
283 };
284 template <> struct restype_ex<int, long long> {
285  using type = long long;
286 };
287 template <> struct restype_ex<int, unsigned long long> {
288  using type = long long;
289 };
290 
291 template <> struct restype_ex<unsigned int, char> {
292  using type = long long;
293 };
294 template <> struct restype_ex<unsigned int, unsigned char> {
295  using type = long long;
296 };
297 template <> struct restype_ex<unsigned int, short> {
298  using type = long long;
299 };
300 template <> struct restype_ex<unsigned int, unsigned short> {
301  using type = long long;
302 };
303 template <> struct restype_ex<unsigned int, int> {
304  using type = long long;
305 };
306 template <> struct restype_ex<unsigned int, unsigned int> {
307  using type = long long;
308 };
309 template <> struct restype_ex<unsigned int, half> {
310  using type = half;
311 };
312 template <> struct restype_ex<unsigned int, float> {
313  using type = float;
314 };
315 template <> struct restype_ex<unsigned int, double> {
316  using type = double;
317 };
318 template <> struct restype_ex<unsigned int, long long> {
319  using type = long long;
320 };
321 template <> struct restype_ex<unsigned int, unsigned long long> {
322  using type = long long;
323 };
324 
325 template <> struct restype_ex<half, char> {
326  using type = half;
327 };
328 template <> struct restype_ex<half, unsigned char> {
329  using type = half;
330 };
331 template <> struct restype_ex<half, short> {
332  using type = half;
333 };
334 template <> struct restype_ex<half, unsigned short> {
335  using type = half;
336 };
337 template <> struct restype_ex<half, int> {
338  using type = half;
339 };
340 template <> struct restype_ex<half, unsigned int> {
341  using type = half;
342 };
343 template <> struct restype_ex<half, half> {
344  using type = half;
345 };
346 template <> struct restype_ex<half, float> {
347  using type = float;
348 };
349 template <> struct restype_ex<half, double> {
350  using type = double;
351 };
352 template <> struct restype_ex<half, long long> {
353  using type = half;
354 };
355 template <> struct restype_ex<half, unsigned long long> {
356  using type = half;
357 };
358 
359 template <> struct restype_ex<float, char> {
360  using type = float;
361 };
362 template <> struct restype_ex<float, unsigned char> {
363  using type = float;
364 };
365 template <> struct restype_ex<float, short> {
366  using type = float;
367 };
368 template <> struct restype_ex<float, unsigned short> {
369  using type = float;
370 };
371 template <> struct restype_ex<float, int> {
372  using type = float;
373 };
374 template <> struct restype_ex<float, unsigned int> {
375  using type = float;
376 };
377 template <> struct restype_ex<float, half> {
378  using type = float;
379 };
380 template <> struct restype_ex<float, float> {
381  using type = float;
382 };
383 template <> struct restype_ex<float, double> {
384  using type = double;
385 };
386 template <> struct restype_ex<float, long long> {
387  using type = float;
388 };
389 template <> struct restype_ex<float, unsigned long long> {
390  using type = float;
391 };
392 
393 template <> struct restype_ex<double, char> {
394  using type = double;
395 };
396 template <> struct restype_ex<double, unsigned char> {
397  using type = double;
398 };
399 template <> struct restype_ex<double, short> {
400  using type = double;
401 };
402 template <> struct restype_ex<double, unsigned short> {
403  using type = double;
404 };
405 template <> struct restype_ex<double, int> {
406  using type = double;
407 };
408 template <> struct restype_ex<double, unsigned int> {
409  using type = double;
410 };
411 template <> struct restype_ex<double, half> {
412  using type = double;
413 };
414 template <> struct restype_ex<double, float> {
415  using type = double;
416 };
417 template <> struct restype_ex<double, double> {
418  using type = double;
419 };
420 template <> struct restype_ex<double, long long> {
421  using type = double;
422 };
423 template <> struct restype_ex<double, unsigned long long> {
424  using type = double;
425 };
426 
427 template <> struct restype_ex<long long, char> {
428  using type = long long;
429 };
430 template <> struct restype_ex<long long, unsigned char> {
431  using type = long long;
432 };
433 template <> struct restype_ex<long long, short> {
434  using type = long long;
435 };
436 template <> struct restype_ex<long long, unsigned short> {
437  using type = long long;
438 };
439 template <> struct restype_ex<long long, int> {
440  using type = long long;
441 };
442 template <> struct restype_ex<long long, unsigned int> {
443  using type = long long;
444 };
445 template <> struct restype_ex<long long, half> {
446  using type = half;
447 };
448 template <> struct restype_ex<long long, float> {
449  using type = float;
450 };
451 template <> struct restype_ex<long long, double> {
452  using type = double;
453 };
454 template <> struct restype_ex<long long, long long> {
455  using type = long long;
456 };
457 template <> struct restype_ex<long long, unsigned long long> {
458  using type = long long;
459 };
460 
461 template <> struct restype_ex<unsigned long long, char> {
462  using type = long long;
463 };
464 template <> struct restype_ex<unsigned long long, unsigned char> {
465  using type = long long;
466 };
467 template <> struct restype_ex<unsigned long long, short> {
468  using type = long long;
469 };
470 template <> struct restype_ex<unsigned long long, unsigned short> {
471  using type = long long;
472 };
473 template <> struct restype_ex<unsigned long long, int> {
474  using type = long long;
475 };
476 template <> struct restype_ex<unsigned long long, unsigned int> {
477  using type = long long;
478 };
479 template <> struct restype_ex<unsigned long long, half> {
480  using type = half;
481 };
482 template <> struct restype_ex<unsigned long long, float> {
483  using type = float;
484 };
485 template <> struct restype_ex<unsigned long long, double> {
486  using type = double;
487 };
488 template <> struct restype_ex<unsigned long long, long long> {
489  using type = long long;
490 };
491 template <> struct restype_ex<unsigned long long, unsigned long long> {
492  using type = long long;
493 };
494 
495 // used in emulation of shl etc operations
496 template <typename T> struct maxtype {
497  using type = T;
498 };
499 template <> struct maxtype<char> {
500  using type = int;
501 };
502 template <> struct maxtype<short> {
503  using type = int;
504 };
505 template <> struct maxtype<unsigned char> {
506  using type = unsigned int;
507 };
508 template <> struct maxtype<unsigned short> {
509  using type = unsigned int;
510 };
511 
512 // used in emulation of abs
513 template <typename T> struct abstype {
514  using type = T;
515 };
516 template <> struct abstype<char> {
517  using type = unsigned char;
518 };
519 template <> struct abstype<short> {
520  using type = unsigned short;
521 };
522 template <> struct abstype<long long> {
523  using type = unsigned long long;
524 };
525 
526 template <bool VALUE> struct check_true {
527  static const bool value = false;
528 };
529 template <> struct check_true<true> {
530  static const bool value = true;
531 };
532 
533 template <typename T> struct is_inttype {
534  static const bool value = false;
535 };
536 template <> struct is_inttype<char> {
537  static const bool value = true;
538 };
539 template <> struct is_inttype<unsigned char> {
540  static const bool value = true;
541 };
542 template <> struct is_inttype<short> {
543  static const bool value = true;
544 };
545 template <> struct is_inttype<unsigned short> {
546  static const bool value = true;
547 };
548 template <> struct is_inttype<int> {
549  static const bool value = true;
550 };
551 template <> struct is_inttype<unsigned int> {
552  static const bool value = true;
553 };
554 template <> struct is_inttype<long long> {
555  static const bool value = true;
556 };
557 template <> struct is_inttype<unsigned long long> {
558  static const bool value = true;
559 };
560 
561 template <typename T> struct is_byte_type {
562  static const bool value = false;
563 };
564 template <> struct is_byte_type<char> {
565  static const bool value = true;
566 };
567 template <> struct is_byte_type<unsigned char> {
568  static const bool value = true;
569 };
570 
571 template <typename T> struct is_word_type {
572  static const bool value = false;
573 };
574 template <> struct is_word_type<short> {
575  static const bool value = true;
576 };
577 template <> struct is_word_type<unsigned short> {
578  static const bool value = true;
579 };
580 
581 template <typename T> struct is_dword_type {
582  static const bool value = false;
583 };
584 template <> struct is_dword_type<int> {
585  static const bool value = true;
586 };
587 template <> struct is_dword_type<unsigned int> {
588  static const bool value = true;
589 };
590 
591 template <typename T> struct is_qf_type {
592  static const bool value = false;
593 };
594 template <> struct is_qf_type<unsigned char> {
595  static const bool value = true;
596 };
597 
598 template <typename T> struct is_hf_type {
599  static const bool value = false;
600 };
601 template <> struct is_hf_type<half> {
602  static const bool value = true;
603 };
604 
605 template <typename T> struct is_fp_type {
606  static const bool value = false;
607 };
608 template <> struct is_fp_type<float> {
609  static const bool value = true;
610 };
611 
612 template <typename T> struct is_df_type {
613  static const bool value = false;
614 };
615 template <> struct is_df_type<double> {
616  static const bool value = true;
617 };
618 
619 template <typename T> struct is_fp_or_dword_type {
620  static const bool value = false;
621 };
622 template <> struct is_fp_or_dword_type<int> {
623  static const bool value = true;
624 };
625 template <> struct is_fp_or_dword_type<unsigned int> {
626  static const bool value = true;
627 };
628 template <> struct is_fp_or_dword_type<float> {
629  static const bool value = true;
630 };
631 // The check is only used for dataport APIs,
632 // which also support df data type.
633 template <> struct is_fp_or_dword_type<double> {
634  static const bool value = true;
635 };
636 
637 template <typename T> struct is_ushort_type {
638  static const bool value = false;
639 };
640 template <> struct is_ushort_type<unsigned short> {
641  static const bool value = true;
642 };
643 
644 template <typename T1, typename T2> struct is_float_dword {
645  static const bool value = false;
646 };
647 template <> struct is_float_dword<float, int> {
648  static const bool value = true;
649 };
650 template <> struct is_float_dword<float, unsigned int> {
651  static const bool value = true;
652 };
653 template <> struct is_float_dword<int, float> {
654  static const bool value = true;
655 };
656 template <> struct is_float_dword<unsigned int, float> {
657  static const bool value = true;
658 };
659 
660 template <typename T> struct hftype {
661  static const bool value = false;
662 };
663 template <> struct hftype<half> {
664  static const bool value = true;
665 };
666 
667 template <typename T> struct fptype {
668  static const bool value = false;
669 };
670 template <> struct fptype<float> {
671  static const bool value = true;
672 };
673 
674 template <typename T> struct dftype {
675  static const bool value = false;
676 };
677 template <> struct dftype<double> {
678  static const bool value = true;
679 };
680 
681 template <typename T> struct bytetype;
682 template <> struct bytetype<char> {
683  static const bool value = true;
684 };
685 template <> struct bytetype<unsigned char> {
686  static const bool value = true;
687 };
688 
689 template <typename T> struct wordtype;
690 template <> struct wordtype<short> {
691  static const bool value = true;
692 };
693 template <> struct wordtype<unsigned short> {
694  static const bool value = true;
695 };
696 
697 template <typename T> struct dwordtype;
698 template <> struct dwordtype<int> {
699  static const bool value = true;
700 };
701 template <> struct dwordtype<unsigned int> {
702  static const bool value = true;
703 };
704 
705 } // namespace __ESIMD_EMU_DNS
706 } // __SYCL_INLINE_NAMESPACE(cl)
707 
708 #endif // #ifndef __SYCL_DEVICE_ONLY__
709 
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
T
sycl
Definition: invoke_simd.hpp:68
max
simd< _Tp, _Abi > max(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
cl::sycl::half
cl::sycl::detail::half_impl::half half
Definition: aliases.hpp:77
char
cl::sycl::abs
detail::enable_if_t< detail::is_genfloat< T >::value, T > abs(T x) __NOEXC
Definition: builtins.hpp:530
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
elem_type_traits.hpp
cl::sycl::image_channel_order::a
@ a
std
Definition: accessor.hpp:2617
cl::sycl::detail::pi
Definition: backend_traits_opencl.hpp:193
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
min
simd< _Tp, _Abi > min(const simd< _Tp, _Abi > &, const simd< _Tp, _Abi > &) noexcept
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12