DPC++ Runtime
Runtime libraries for oneAPI DPC++
util.hpp
Go to the documentation of this file.
1 //==----------------- util.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 // Utility functions used for implementing Explicit SIMD APIs.
9 //===----------------------------------------------------------------------===//
10 
11 #pragma once
12 
14 
17 
18 #include <type_traits>
19 
20 #ifdef __SYCL_DEVICE_ONLY__
21 #define __ESIMD_INTRIN __DPCPP_SYCL_EXTERNAL SYCL_ESIMD_FUNCTION
22 #else
23 #define __ESIMD_INTRIN inline
24 #endif // __SYCL_DEVICE_ONLY__
25 
26 namespace sycl {
28 namespace ext::intel::esimd::detail {
29 
31 struct OperandSize {
32  enum { BYTE = 1, WORD = 2, DWORD = 4, QWORD = 8, OWORD = 16, GRF = 32 };
33 };
34 
37 template <unsigned int N, unsigned int K, bool K_gt_eq_N> struct NextPowerOf2;
38 
39 template <unsigned int N, unsigned int K> struct NextPowerOf2<N, K, true> {
40  static constexpr unsigned int get() { return K; }
41 };
42 
43 template <unsigned int N, unsigned int K> struct NextPowerOf2<N, K, false> {
44  static constexpr unsigned int get() {
45  return NextPowerOf2<N, K * 2, K * 2 >= N>::get();
46  }
47 };
48 
49 template <unsigned int N> constexpr unsigned int getNextPowerOf2() {
50  return NextPowerOf2<N, 1, (1 >= N)>::get();
51 }
52 
53 template <> constexpr unsigned int getNextPowerOf2<0>() { return 0; }
54 
57 template <unsigned int N, bool N_gt_1> struct Log2;
58 
59 template <unsigned int N> struct Log2<N, false> {
60  static constexpr unsigned int get() { return 0; }
61 };
62 
63 template <unsigned int N> struct Log2<N, true> {
64  static constexpr unsigned int get() {
65  return 1 + Log2<(N >> 1), ((N >> 1) > 1)>::get();
66  }
67 };
68 
69 template <unsigned int N> constexpr unsigned int log2() {
70  return Log2<N, (N > 1)>::get();
71 }
72 
74 template <typename T> struct is_esimd_vector : public std::false_type {};
75 
76 template <typename T, int N>
77 struct is_esimd_vector<simd<T, N>> : public std::true_type {};
78 
79 template <typename T, int N>
80 using is_hw_int_type =
81  typename std::bool_constant<std::is_integral_v<T> && (sizeof(T) == N)>;
82 
83 template <typename T> using is_qword_type = is_hw_int_type<T, 8>;
84 template <typename T> using is_dword_type = is_hw_int_type<T, 4>;
85 template <typename T> using is_word_type = is_hw_int_type<T, 2>;
86 template <typename T> using is_byte_type = is_hw_int_type<T, 1>;
87 
88 template <typename T, int N>
89 using is_hw_fp_type = typename std::bool_constant<std::is_floating_point_v<T> &&
90  (sizeof(T) == N)>;
91 
92 template <typename T> using is_fp_type = is_hw_fp_type<T, 4>;
93 template <typename T> using is_df_type = is_hw_fp_type<T, 8>;
94 
95 template <typename T>
96 using is_fp_or_dword_type =
97  typename std::bool_constant<is_fp_type<T>::value ||
98  is_dword_type<T>::value>;
99 
101 template <typename T> struct simd_type {
102  using type = simd<T, 1>;
103 };
104 template <typename T, int N> struct simd_type<raw_vector_type<T, N>> {
105  using type = simd<T, N>;
106 };
107 
108 template <typename T> struct simd_type<T &> {
109  using type = typename simd_type<T>::type;
110 };
111 template <typename T> struct simd_type<T &&> {
112  using type = typename simd_type<T>::type;
113 };
114 template <typename T> struct simd_type<const T> {
115  using type = typename simd_type<T>::type;
116 };
117 
118 template <typename T> struct dword_type {
119  using type = T;
120 };
121 template <> struct dword_type<char> {
122  using type = int;
123 };
124 template <> struct dword_type<short> {
125  using type = int;
126 };
127 template <> struct dword_type<uchar> {
128  using type = uint;
129 };
130 template <> struct dword_type<ushort> {
131  using type = uint;
132 };
133 
134 template <typename T> struct byte_type {
135  using type = T;
136 };
137 template <> struct byte_type<short> {
138  using type = char;
139 };
140 template <> struct byte_type<int> {
141  using type = char;
142 };
143 template <> struct byte_type<ushort> {
144  using type = uchar;
145 };
146 template <> struct byte_type<uint> {
147  using type = uchar;
148 };
149 
150 template <typename T> struct word_type {
151  using type = T;
152 };
153 template <> struct word_type<char> {
154  using type = short;
155 };
156 template <> struct word_type<int> {
157  using type = short;
158 };
159 template <> struct word_type<uchar> {
160  using type = ushort;
161 };
162 template <> struct word_type<uint> {
163  using type = ushort;
164 };
165 
166 // Utility for compile time loop unrolling.
167 template <unsigned N> class ForHelper {
168  template <unsigned I, typename Action> static inline void repeat(Action A) {
169  if constexpr (I < N)
170  A(I);
171  if constexpr (I + 1 < N)
172  repeat<I + 1, Action>(A);
173  }
174 
175 public:
176  template <typename Action> static inline void unroll(Action A) {
177  ForHelper::template repeat<0, Action>(A);
178  }
179 };
180 
181 #ifdef __ESIMD_FORCE_STATELESS_MEM
182 template <typename T, typename AccessorTy, typename OffsetTy = uint32_t>
185 auto accessorToPointer(AccessorTy Acc, OffsetTy Offset = 0) {
186  using QualCharPtrType =
187  std::conditional_t<std::is_const_v<typename AccessorTy::value_type>,
188  const char *, char *>;
189  using QualTPtrType =
190  std::conditional_t<std::is_const_v<typename AccessorTy::value_type>,
191  const T *, T *>;
192  auto BytePtr = reinterpret_cast<QualCharPtrType>(Acc.get_pointer()) + Offset;
193  return reinterpret_cast<QualTPtrType>(BytePtr);
194 }
195 #endif // __ESIMD_FORCE_STATELESS_MEM
196 
197 } // namespace ext::intel::esimd::detail
198 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
199 } // namespace sycl
200 
T
type_traits.hpp
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:11
sycl::_V1::addressing_mode::repeat
@ repeat
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
sycl::_V1::ext::intel::esimd::uchar
unsigned char uchar
Definition: common.hpp:41
std::get
constexpr tuple_element< I, tuple< Types... > >::type & get(sycl::detail::tuple< Types... > &Arg) noexcept
Definition: tuple.hpp:199
types.hpp
sycl::_V1::ext::intel::esimd::rgba_channel::A
@ A
simd
Definition: simd.hpp:1027
sycl::_V1::ext::intel::esimd::ushort
unsigned short ushort
Definition: common.hpp:42
sycl::_V1::ext::intel::esimd::log2
__ESIMD_API simd< T, N > log2(simd< T, N > src, Sat sat={})
Logarithm base 2.
Definition: math.hpp:381
sycl::_V1::ext::intel::esimd::uint
unsigned int uint
Definition: common.hpp:43