DPC++ Runtime
Runtime libraries for oneAPI DPC++
invoke_simd.hpp
Go to the documentation of this file.
1 //==------ invoke_simd.hpp - SYCL invoke_simd extension --*- C++ -*---------==//
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 // Implemenation of the sycl_ext_oneapi_invoke_simd extension.
9 // https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_invoke_simd.asciidoc
10 // ===--------------------------------------------------------------------=== //
11 
12 #pragma once
13 
14 // SYCL extension macro definition as required by the SYCL specification.
15 // 1 - Initial extension version. Base features are supported.
16 #define SYCL_EXT_ONEAPI_INVOKE_SIMD 1
17 
19 
21 #include <sycl/detail/boost/mp11.hpp>
22 #include <sycl/sub_group.hpp>
23 
24 #include <functional>
25 
26 // TODOs:
27 // * (a) TODO bool translation in spmd2simd.
28 // * (b) TODO enforce constness of a functor/lambda's () operator
29 // * (c) TODO support lambdas and functors in BE
30 
42 template <bool IsFunc, class SpmdRet, class HelperFunc,
43  class... UserSimdFuncAndSpmdArgs, class = std::enable_if_t<!IsFunc>>
44 SYCL_EXTERNAL __regcall SpmdRet
45 __builtin_invoke_simd(HelperFunc helper, const void *obj,
46  UserSimdFuncAndSpmdArgs... args)
47 #ifdef __SYCL_DEVICE_ONLY__
48  ;
49 #else
50 {
51  // __builtin_invoke_simd is not supported on the host device yet
52  throw sycl::exception(sycl::errc::feature_not_supported,
53  "__builtin_invoke_simd is not supported on host");
54 }
55 #endif // __SYCL_DEVICE_ONLY__
56 
57 template <bool IsFunc, class SpmdRet, class HelperFunc,
58  class... UserSimdFuncAndSpmdArgs, class = std::enable_if_t<IsFunc>>
59 SYCL_EXTERNAL __regcall SpmdRet
60 __builtin_invoke_simd(HelperFunc helper, UserSimdFuncAndSpmdArgs... args)
61 #ifdef __SYCL_DEVICE_ONLY__
62  ;
63 #else
64 {
65  // __builtin_invoke_simd is not supported on the host device yet
66  throw sycl::exception(sycl::errc::feature_not_supported,
67  "__builtin_invoke_simd is not supported on host");
68 }
69 #endif // __SYCL_DEVICE_ONLY__
70 
71 namespace sycl {
73 
74 namespace ext {
75 namespace oneapi {
76 namespace experimental {
77 
78 // --- Basic definitions prescribed by the spec.
79 namespace simd_abi {
80 // "Fixed-size simd width of N" ABI based on clang vectors - used as the ABI for
81 // SIMD objects this implementation of invoke_simd spec is based on.
82 template <class T, int N>
83 using native_fixed_size = typename std::experimental::__simd_abi<
84  std::experimental::_StorageKind::_VecExt, N>;
85 } // namespace simd_abi
86 
87 // The SIMD object type, which is the generic std::experimental::simd type with
88 // the native fixed size ABI.
89 template <class T, int N>
90 using simd = std::experimental::simd<T, simd_abi::native_fixed_size<T, N>>;
91 
92 // The SIMD mask object type.
93 template <class T, int N>
94 using simd_mask =
95  std::experimental::simd_mask<T, simd_abi::native_fixed_size<T, N>>;
96 
97 // --- Helpers
98 namespace detail {
99 
100 namespace __MP11_NS = sycl::detail::boost::mp11;
101 
102 // This structure performs the SPMD-to-SIMD parameter type conversion as defined
103 // by the spec.
104 template <class T, int N, class = void> struct spmd2simd;
105 // * `uniform<T>` converts to `T`
106 template <class T, int N> struct spmd2simd<uniform<T>, N> {
107  using type = T;
108 };
109 // * tuple of types converts to tuple of converted tuple element types.
110 template <class... T, int N> struct spmd2simd<std::tuple<T...>, N> {
111  using type = std::tuple<typename spmd2simd<T, N>::type...>;
112 };
113 // * arithmetic type converts to a simd vector with this element type and the
114 // width equal to caller's subgroup size and passed as the `N` template
115 // argument.
116 template <class T, int N>
117 struct spmd2simd<T, N, std::enable_if_t<std::is_arithmetic_v<T>>> {
118  using type = simd<T, N>;
119 };
120 
121 // This structure performs the SIMD-to-SPMD return type conversion as defined
122 // by the spec.
123 template <class, class = void> struct simd2spmd;
124 // * `uniform<T>` stays the same
125 template <class T> struct simd2spmd<uniform<T>> {
126  using type = uniform<T>;
127 };
128 // * `simd<T, N>` converts to T
129 template <class T, int N> struct simd2spmd<simd<T, N>> {
130  using type = T;
131 };
132 // * tuple of types converts to tuple of converted tuple element types.
133 template <class... T> struct simd2spmd<std::tuple<T...>> {
134  using type = std::tuple<typename simd2spmd<T>::type...>;
135 };
136 // * arithmetic type T converts to `uniform<T>`
137 template <class T>
138 struct simd2spmd<T, std::enable_if_t<std::is_arithmetic_v<T>>> {
139  using type = uniform<T>;
140 };
141 
142 // Determine number of elements in a simd type.
143 template <class T> struct simd_size {
144  static constexpr int value = 1; // 1 element in any type by default
145 };
146 
147 // * Specialization for the simd type.
148 template <class T, int N> struct simd_size<simd<T, N>> {
149  static constexpr int value = N;
150 };
151 
152 // Check if given type is uniform.
153 template <class T> struct is_uniform_type : std::false_type {};
154 template <class T> struct is_uniform_type<uniform<T>> : std::true_type {
155  using type = T;
156 };
157 
158 // Check if given type is simd or simd_mask.
159 template <class T> struct is_simd_or_mask_type : std::false_type {};
160 template <class T, int N>
161 struct is_simd_or_mask_type<simd<T, N>> : std::true_type {};
162 template <class T, int N>
163 struct is_simd_or_mask_type<simd_mask<T, N>> : std::true_type {};
164 
165 // Checks if all the types in the parameter pack are uniform<T>.
166 template <class... SpmdArgs> struct all_uniform_types {
167  constexpr operator bool() {
168  using TypeList = __MP11_NS::mp_list<SpmdArgs...>;
169  return __MP11_NS::mp_all_of<TypeList, is_uniform_type>::value;
170  }
171 };
172 
173 // "Unwraps" a value of the `uniform` type (used before passing to SPMD
174 // arguments to the __builtin_invoke_simd):
175 // - the case when there is nothing to unwrap
176 template <typename T> struct unwrap_uniform {
177  static auto impl(T val) { return val; }
178 };
179 
180 // - the real unwrapping case
181 template <typename T> struct unwrap_uniform<uniform<T>> {
182  static T impl(uniform<T> val) { return val; }
183 };
184 
185 // Deduces subgroup size of the caller based on given SIMD callable and
186 // corresponding SPMD arguments it is being invoke with via invoke_simd.
187 // Basically, for each supported subgroup size, this meta-function finds out if
188 // the callable can be invoked by C++ rules given the SPMD arguments transformed
189 // as prescribed by the spec assuming this subgroup size. One and only one
190 // subgroup size should conform.
191 template <class SimdCallable, class... SpmdArgs> struct sg_size {
192  template <class N>
193  using IsInvocableSgSize = __MP11_NS::mp_bool<std::is_invocable_v<
194  SimdCallable, typename spmd2simd<SpmdArgs, N::value>::type...>>;
195 
196  SYCL_EXTERNAL constexpr operator int() {
197  using SupportedSgSizes = __MP11_NS::mp_list_c<int, 1, 2, 4, 8, 16, 32>;
198  using InvocableSgSizes =
199  __MP11_NS::mp_copy_if<SupportedSgSizes, IsInvocableSgSize>;
200  static_assert((__MP11_NS::mp_size<InvocableSgSizes>::value == 1) &&
201  "no or multiple invoke_simd targets found");
202  return __MP11_NS::mp_front<InvocableSgSizes>::value;
203  }
204 };
205 
206 // Determine the return type of a SIMD callable.
207 template <int N, class SimdCallable, class... SpmdArgs>
208 using SimdRetType =
209  std::invoke_result_t<SimdCallable,
211 // Determine the return type of an invoke_simd based on the return type of a
212 // SIMD callable.
213 template <int N, class SimdCallable, class... SpmdArgs>
214 using SpmdRetType =
215  typename simd2spmd<SimdRetType<N, SimdCallable, SpmdArgs...>>::type;
216 
217 template <class SimdCallable, class... SpmdArgs>
218 static constexpr int get_sg_size() {
219  if constexpr (all_uniform_types<SpmdArgs...>()) {
220  using SimdRet = std::invoke_result_t<SimdCallable, SpmdArgs...>;
221 
222  if constexpr (is_simd_or_mask_type<SimdRet>::value) {
224  } else {
225  // fully uniform function - subgroup size does not matter
226  return 0;
227  }
228  } else {
229  return sg_size<SimdCallable, SpmdArgs...>();
230  }
231 }
232 
233 // This function is a wrapper around a call to a functor with field or a lambda
234 // with captures. Note __regcall - this is needed for efficient argument
235 // forwarding.
236 template <int N, class Callable, class... T>
237 [[intel::device_indirectly_callable]] SYCL_EXTERNAL __regcall detail::
238  SimdRetType<N, Callable, T...>
239  simd_obj_call_helper(const void *obj_ptr,
240  typename detail::spmd2simd<T, N>::type... simd_args) {
241  auto f =
242  *reinterpret_cast<const std::remove_reference_t<Callable> *>(obj_ptr);
243  return f(simd_args...);
244 }
245 
246 // This function is a wrapper around a call to a function.
247 template <int N, class Callable, class... T>
248 [[intel::device_indirectly_callable]] SYCL_EXTERNAL __regcall detail::
249  SimdRetType<N, Callable, T...>
251  typename detail::spmd2simd<T, N>::type... simd_args) {
252  return f(simd_args...);
253 }
254 
255 #ifdef _GLIBCXX_RELEASE
256 #if _GLIBCXX_RELEASE < 10
257 #define __INVOKE_SIMD_USE_STD_IS_FUNCTION_WA
258 #endif // _GLIBCXX_RELEASE < 10
259 #endif // _GLIBCXX_RELEASE
260 
261 #ifdef __INVOKE_SIMD_USE_STD_IS_FUNCTION_WA
262 // TODO This is a workaround for libstdc++ version 9 buggy behavior which
263 // returns false in the code below. Version 10 works fine. Once required
264 // minimum libstdc++ version is bumped to 10, this w/a should be removed.
265 // template <class F> bool foo(F &&f) {
266 // return std::is_function_v<std::remove_reference_t<F>>;
267 // }
268 // where F is a function type with __regcall.
269 template <class F> struct is_regcall_function_ptr_or_ref : std::false_type {};
270 
271 template <class Ret, class... Args>
272 struct is_regcall_function_ptr_or_ref<Ret(__regcall &)(Args...)>
273  : std::true_type {};
274 
275 template <class Ret, class... Args>
276 struct is_regcall_function_ptr_or_ref<Ret(__regcall *)(Args...)>
277  : std::true_type {};
278 
279 template <class Ret, class... Args>
280 struct is_regcall_function_ptr_or_ref<Ret(__regcall *&)(Args...)>
281  : std::true_type {};
282 
283 template <class F>
284 static constexpr bool is_regcall_function_ptr_or_ref_v =
285  is_regcall_function_ptr_or_ref<F>::value;
286 #endif // __INVOKE_SIMD_USE_STD_IS_FUNCTION_WA
287 
288 template <class Callable>
289 static constexpr bool is_function_ptr_or_ref_v =
290  std::is_function_v<std::remove_pointer_t<std::remove_reference_t<Callable>>>
291 #ifdef __INVOKE_SIMD_USE_STD_IS_FUNCTION_WA
292  || is_regcall_function_ptr_or_ref_v<Callable>
293 #endif // __INVOKE_SIMD_USE_STD_IS_FUNCTION_WA
294  ;
295 
296 template <typename Callable> struct remove_ref_from_func_ptr_ref_type {
297  using type = Callable;
298 };
299 
300 template <typename Ret, typename... Args>
301 struct remove_ref_from_func_ptr_ref_type<Ret(__regcall *&)(Args...)> {
302  using type = Ret(__regcall *)(Args...);
303 };
304 
305 template <typename T>
308 
309 } // namespace detail
310 
311 // --- The main API
312 
329 // TODO works only for functions and pointers to functions now,
330 // enable for lambda functions and functors.
331 template <class Callable, class... T>
332 __attribute__((always_inline)) auto invoke_simd(sycl::sub_group sg,
333  Callable &&f, T... args) {
334  // If the invoke_simd call site is fully uniform, then it does not matter
335  // what the subgroup size is and arguments don't need widening and return
336  // value does not need shrinking by this library or SPMD compiler, so 0
337  // is fine in this case.
338  constexpr int N = detail::get_sg_size<Callable, T...>();
339  using RetSpmd = detail::SpmdRetType<N, Callable, T...>;
340  constexpr bool is_function = detail::is_function_ptr_or_ref_v<Callable>;
341 
342  if constexpr (is_function) {
343  // The variables typed as pointer to a function become lvalue-reference
344  // when passed to invoke_simd() as universal pointers. That creates an
345  // additional indirection, which is resolved automatically by the compiler
346  // for the caller side of __builtin_invoke_simd, but which must be resolved
347  // manually during the creation of simd_func_call_helper.
348  // The class remove_ref_from_func_ptr_ref_type is used removes that
349  // unwanted indirection.
350  return __builtin_invoke_simd<true /*function*/, RetSpmd>(
353  f, detail::unwrap_uniform<T>::impl(args)...);
354  } else {
355  // TODO support functors and lambdas which are handled in this branch.
356  // The limiting factor for now is that the LLVMIR data flow analysis
357  // implemented in LowerInvokeSimd.cpp which, finds actual invoke_simd
358  // target function, can't handle this case yet.
359  return __builtin_invoke_simd<false /*functor/lambda*/, RetSpmd>(
360  detail::simd_obj_call_helper<N, Callable, T...>, &f,
362  }
363 // TODO Temporary macro and assert to enable API compilation testing.
364 // LowerInvokeSimd.cpp does not support this case yet.
365 #ifndef __INVOKE_SIMD_ENABLE_ALL_CALLABLES
366  static_assert(is_function &&
367  "invoke_simd does not support functors or lambdas yet");
368 #endif // __INVOKE_SIMD_ENABLE_ALL_CALLABLES
369 }
370 
371 } // namespace experimental
372 } // namespace oneapi
373 } // namespace ext
374 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
375 } // namespace sycl
sycl::_V1::ext::oneapi::experimental::detail::unwrap_uniform< uniform< T > >::impl
static T impl(uniform< T > val)
Definition: invoke_simd.hpp:182
sycl::_V1::ext::oneapi::experimental::detail::is_function_ptr_or_ref_v
static constexpr bool is_function_ptr_or_ref_v
Definition: invoke_simd.hpp:289
simd_mask
Definition: simd.hpp:1029
sub_group.hpp
sycl::_V1::ext::oneapi::experimental::detail::is_simd_or_mask_type
Definition: invoke_simd.hpp:159
simd.hpp
sycl::_V1::ext::oneapi::experimental::detail::simd_obj_call_helper
SYCL_EXTERNAL __regcall detail::SimdRetType< N, Callable, T... > simd_obj_call_helper(const void *obj_ptr, typename detail::spmd2simd< T, N >::type... simd_args)
Definition: invoke_simd.hpp:239
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:13
sycl::_V1::ext::oneapi::experimental::detail::spmd2simd< T, N, std::enable_if_t< std::is_arithmetic_v< T > > >::type
simd< T, N > type
Definition: invoke_simd.hpp:118
sycl::_V1::ext::oneapi::experimental::detail::SpmdRetType
typename simd2spmd< SimdRetType< N, SimdCallable, SpmdArgs... > >::type SpmdRetType
Definition: invoke_simd.hpp:215
uniform.hpp
sycl::_V1::ext::oneapi::experimental::detail::simd_func_call_helper
SYCL_EXTERNAL __regcall detail::SimdRetType< N, Callable, T... > simd_func_call_helper(Callable f, typename detail::spmd2simd< T, N >::type... simd_args)
Definition: invoke_simd.hpp:250
sycl::_V1::ext::oneapi::experimental::detail::all_uniform_types
Definition: invoke_simd.hpp:166
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:13
sycl::_V1::ext::oneapi::experimental::detail::get_sg_size
static constexpr int get_sg_size()
Definition: invoke_simd.hpp:218
sycl::_V1::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
sycl::_V1::ext::oneapi::experimental::detail::sg_size
Definition: invoke_simd.hpp:191
sycl::_V1::ext::oneapi::experimental::uniform
Definition: uniform.hpp:73
sycl::_V1::ext::oneapi::experimental::detail::spmd2simd< std::tuple< T... >, N >::type
std::tuple< typename spmd2simd< T, N >::type... > type
Definition: invoke_simd.hpp:111
sycl::_V1::ext::oneapi::experimental::detail::unwrap_uniform::impl
static auto impl(T val)
Definition: invoke_simd.hpp:177
sycl::_V1::ext::oneapi::experimental::detail::remove_ref_from_func_ptr_ref_type::type
Callable type
Definition: invoke_simd.hpp:297
sycl::_V1::ext::oneapi::experimental::detail::is_uniform_type
Definition: invoke_simd.hpp:153
SYCL_EXTERNAL
#define SYCL_EXTERNAL
Definition: defines_elementary.hpp:32
sycl::_V1::ext::oneapi::experimental::detail::remove_ref_from_func_ptr_ref_type
Definition: invoke_simd.hpp:296
sycl::_V1::ext::oneapi::experimental::detail::simd2spmd< std::tuple< T... > >::type
std::tuple< typename simd2spmd< T >::type... > type
Definition: invoke_simd.hpp:134
sycl::_V1::ext::oneapi::experimental::detail::spmd2simd< uniform< T >, N >::type
T type
Definition: invoke_simd.hpp:107
sycl::_V1::ext::oneapi::experimental::simd_abi::native_fixed_size
typename std::experimental::__simd_abi< std::experimental::_StorageKind::_VecExt, N > native_fixed_size
Definition: invoke_simd.hpp:84
sycl::_V1::ext::oneapi::experimental::detail::remove_ref_from_func_ptr_ref_type< Ret(__regcall *&)(Args...)>::type
Ret(__regcall *)(Args...) type
Definition: invoke_simd.hpp:302
__builtin_invoke_simd
SYCL_EXTERNAL __regcall SpmdRet __builtin_invoke_simd(HelperFunc helper, const void *obj, UserSimdFuncAndSpmdArgs... args)
Middle End - to - Back End interface to invoke explicit SIMD functions from SPMD SYCL context.
Definition: invoke_simd.hpp:45
sycl::_V1::ext::oneapi::experimental::detail::sg_size::IsInvocableSgSize
__MP11_NS::mp_bool< std::is_invocable_v< SimdCallable, typename spmd2simd< SpmdArgs, N::value >::type... > > IsInvocableSgSize
Definition: invoke_simd.hpp:194
simd
Definition: simd.hpp:1027
sycl::_V1::ext::oneapi::experimental::detail::simd2spmd< simd< T, N > >::type
T type
Definition: invoke_simd.hpp:130
sycl::_V1::ext::oneapi::experimental::detail::SimdRetType
std::invoke_result_t< SimdCallable, typename spmd2simd< SpmdArgs, N >::type... > SimdRetType
Definition: invoke_simd.hpp:210
std
Definition: accessor.hpp:3071
sycl::_V1::ext::oneapi::experimental::detail::spmd2simd
Definition: invoke_simd.hpp:104
sycl::_V1::ext::oneapi::experimental::detail::remove_ref_from_func_ptr_ref_type_t
typename remove_ref_from_func_ptr_ref_type< T >::type remove_ref_from_func_ptr_ref_type_t
Definition: invoke_simd.hpp:307
sycl::_V1::ext::oneapi::experimental::detail::unwrap_uniform
Definition: invoke_simd.hpp:176
sycl::_V1::ext::oneapi::experimental::detail::is_uniform_type< uniform< T > >::type
T type
Definition: invoke_simd.hpp:155
sycl::_V1::ext::oneapi::experimental::detail::simd2spmd
Definition: invoke_simd.hpp:123
sycl::_V1::ext::oneapi::experimental::__attribute__
__attribute__((always_inline)) auto invoke_simd(sycl
The invoke_simd free function invokes a SIMD function using all work-items in a sub_group.
Definition: invoke_simd.hpp:332
sycl::_V1::ext::oneapi::experimental::detail::simd_size
Definition: invoke_simd.hpp:143