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 
20 #include <CL/sycl/sub_group.hpp>
22 #include <sycl/detail/boost/mp11.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 
40 template <bool IsFunc, class SpmdRet, class SimdCallee, class... SpmdArgs,
41  class = std::enable_if_t<!IsFunc>>
42 SYCL_EXTERNAL __regcall SpmdRet
43 __builtin_invoke_simd(SimdCallee target, const void *obj, SpmdArgs... args)
44 #ifdef __SYCL_DEVICE_ONLY__
45  ;
46 #else
47 {
48  // __builtin_invoke_simd is not supported on the host device yet
49  throw sycl::exception(sycl::errc::feature_not_supported,
50  "__builtin_invoke_simd is not supported on host");
51 }
52 #endif // __SYCL_DEVICE_ONLY__
53 
54 template <bool IsFunc, class SpmdRet, class SimdCallee, class... SpmdArgs,
55  class = std::enable_if_t<IsFunc>>
56 SYCL_EXTERNAL __regcall SpmdRet __builtin_invoke_simd(SimdCallee target,
57  SpmdArgs... args)
58 #ifdef __SYCL_DEVICE_ONLY__
59  ;
60 #else
61 {
62  // __builtin_invoke_simd is not supported on the host device yet
63  throw sycl::exception(sycl::errc::feature_not_supported,
64  "__builtin_invoke_simd is not supported on host");
65 }
66 #endif // __SYCL_DEVICE_ONLY__
67 
68 namespace sycl {
69 namespace ext {
70 namespace oneapi {
71 namespace experimental {
72 
73 // --- Basic definitions prescribed by the spec.
74 namespace simd_abi {
75 // "Fixed-size simd width of N" ABI based on clang vectors - used as the ABI for
76 // SIMD objects this implementation of invoke_simd spec is based on.
77 template <class T, int N>
78 using native_fixed_size = typename std::experimental::__simd_abi<
79  std::experimental::_StorageKind::_VecExt, N>;
80 } // namespace simd_abi
81 
82 // The SIMD object type, which is the generic std::experimental::simd type with
83 // the native fixed size ABI.
84 template <class T, int N>
85 using simd = std::experimental::simd<T, simd_abi::native_fixed_size<T, N>>;
86 
87 // The SIMD mask object type.
88 template <class T, int N>
89 using simd_mask =
90  std::experimental::simd_mask<T, simd_abi::native_fixed_size<T, N>>;
91 
92 // --- Helpers
93 namespace detail {
94 
95 namespace __MP11_NS = sycl::detail::boost::mp11;
96 
97 // This structure performs the SPMD-to-SIMD parameter type conversion as defined
98 // by the spec.
99 template <class T, int N, class = void> struct spmd2simd;
100 // * `uniform<T>` converts to `T`
101 template <class T, int N> struct spmd2simd<uniform<T>, N> {
102  using type = T;
103 };
104 // * tuple of types converts to tuple of converted tuple element types.
105 template <class... T, int N> struct spmd2simd<std::tuple<T...>, N> {
106  using type = std::tuple<typename spmd2simd<T, N>::type...>;
107 };
108 // * arithmetic type converts to a simd vector with this element type and the
109 // width equal to caller's subgroup size and passed as the `N` template
110 // argument.
111 template <class T, int N>
112 struct spmd2simd<T, N, std::enable_if_t<std::is_arithmetic_v<T>>> {
113  using type = simd<T, N>;
114 };
115 
116 // This structure performs the SIMD-to-SPMD return type conversion as defined
117 // by the spec.
118 template <class, class = void> struct simd2spmd;
119 // * `uniform<T>` stays the same
120 template <class T> struct simd2spmd<uniform<T>> {
121  using type = uniform<T>;
122 };
123 // * `simd<T, N>` converts to T
124 template <class T, int N> struct simd2spmd<simd<T, N>> {
125  using type = T;
126 };
127 // * tuple of types converts to tuple of converted tuple element types.
128 template <class... T> struct simd2spmd<std::tuple<T...>> {
129  using type = std::tuple<typename simd2spmd<T>::type...>;
130 };
131 // * arithmetic type T converts to `uniform<T>`
132 template <class T>
133 struct simd2spmd<T, std::enable_if_t<std::is_arithmetic_v<T>>> {
134  using type = uniform<T>;
135 };
136 
137 // Check if given type is uniform.
138 template <class T> struct is_uniform_type : std::false_type {};
139 template <class T> struct is_uniform_type<uniform<T>> : std::true_type {
140  using type = T;
141 };
142 
143 // Check if given type is simd or simd_mask.
144 template <class T> struct is_simd_or_mask_type : std::false_type {};
145 template <class T, int N>
146 struct is_simd_or_mask_type<simd<T, N>> : std::true_type {};
147 template <class T, int N>
148 struct is_simd_or_mask_type<simd_mask<T, N>> : std::true_type {};
149 
150 // Checks if the return value type and the types of arguments of given
151 // SimdCallable are all uniform.
152 template <class SimdCallable, class... SpmdArgs> struct has_uniform_signature {
153  constexpr operator bool() {
154  using ArgTypeList = __MP11_NS::mp_list<SpmdArgs...>;
155 
156  if constexpr (__MP11_NS::mp_all_of<ArgTypeList, is_uniform_type>::value) {
157  using SimdRet = std::invoke_result_t<SimdCallable, SpmdArgs...>;
160  } else {
161  return false;
162  }
163  }
164 };
165 
166 // "Unwraps" a value of the `uniform` type (used before passing to SPMD
167 // arguments to the __builtin_invoke_simd):
168 // - the case when there is nothing to unwrap
169 template <typename T> struct unwrap_uniform {
170  static auto impl(T val) { return val; }
171 };
172 
173 // - the real unwrapping case
174 template <typename T> struct unwrap_uniform<uniform<T>> {
175  static T impl(uniform<T> val) { return val; }
176 };
177 
178 // Deduces subgroup size of the caller based on given SIMD callable and
179 // corresponding SPMD arguments it is being invoke with via invoke_simd.
180 // Basically, for each supported subgroup size, this meta-function finds out if
181 // the callable can be invoked by C++ rules given the SPMD arguments transformed
182 // as prescribed by the spec assuming this subgroup size. One and only one
183 // subgroup size should conform.
184 template <class SimdCallable, class... SpmdArgs> struct sg_size {
185  template <class N>
186  using IsInvocableSgSize = __MP11_NS::mp_bool<std::is_invocable_v<
187  SimdCallable, typename spmd2simd<SpmdArgs, N::value>::type...>>;
188 
189  constexpr operator int() {
190  using SupportedSgSizes = __MP11_NS::mp_list_c<int, 1, 2, 4, 8, 16, 32>;
191  using InvocableSgSizes =
192  __MP11_NS::mp_copy_if<SupportedSgSizes, IsInvocableSgSize>;
193  static_assert((__MP11_NS::mp_size<InvocableSgSizes>::value == 1) &&
194  "no or multiple invoke_simd targets found");
195  return __MP11_NS::mp_front<InvocableSgSizes>::value;
196  }
197 };
198 
199 // Determine the return type of a SIMD callable.
200 template <int N, class SimdCallable, class... SpmdArgs>
201 using SimdRetType =
202  std::invoke_result_t<SimdCallable,
204 // Determine the return type of an invoke_simd based on the return type of a
205 // SIMD callable.
206 template <int N, class SimdCallable, class... SpmdArgs>
207 using SpmdRetType =
208  typename simd2spmd<SimdRetType<N, SimdCallable, SpmdArgs...>>::type;
209 
210 template <class SimdCallable, class... SpmdArgs>
211 static constexpr int get_sg_size() {
213  return 0; // subgroup size does not matter then
214  } else {
215  return sg_size<SimdCallable, SpmdArgs...>();
216  }
217 }
218 
219 // This function is a wrapper around a call to a functor with field or a lambda
220 // with captures. Note __regcall - this is needed for efficient argument
221 // forwarding.
222 template <int N, class Callable, class... T>
223 __regcall detail::SimdRetType<N, Callable, T...>
224 simd_call_helper(const void *obj_ptr,
225  typename detail::spmd2simd<T, N>::type... simd_args) {
226  auto f =
227  *reinterpret_cast<const std::remove_reference_t<Callable> *>(obj_ptr);
228  return f(simd_args...);
229 };
230 
231 #ifdef _GLIBCXX_RELEASE
232 #if _GLIBCXX_RELEASE < 10
233 #define __INVOKE_SIMD_USE_STD_IS_FUNCTION_WA
234 #endif // _GLIBCXX_RELEASE < 10
235 #endif // _GLIBCXX_RELEASE
236 
237 #ifdef __INVOKE_SIMD_USE_STD_IS_FUNCTION_WA
238 // TODO This is a workaround for libstdc++ version 9 buggy behavior which
239 // returns false in the code below. Version 10 works fine. Once required
240 // minimum libstdc++ version is bumped to 10, this w/a should be removed.
241 // template <class F> bool foo(F &&f) {
242 // return std::is_function_v<std::remove_reference_t<F>>;
243 // }
244 // where F is a function type with __regcall.
245 template <class F> struct is_regcall_function_ptr_or_ref : std::false_type {};
246 
247 template <class Ret, class... Args>
248 struct is_regcall_function_ptr_or_ref<Ret(__regcall &)(Args...)>
249  : std::true_type {};
250 
251 template <class Ret, class... Args>
252 struct is_regcall_function_ptr_or_ref<Ret(__regcall *)(Args...)>
253  : std::true_type {};
254 
255 template <class Ret, class... Args>
256 struct is_regcall_function_ptr_or_ref<Ret(__regcall *&)(Args...)>
257  : std::true_type {};
258 
259 template <class F>
260 static constexpr bool is_regcall_function_ptr_or_ref_v =
261  is_regcall_function_ptr_or_ref<F>::value;
262 #endif // __INVOKE_SIMD_USE_STD_IS_FUNCTION_WA
263 
264 template <class Callable>
265 static constexpr bool is_function_ptr_or_ref_v =
266  std::is_function_v<std::remove_pointer_t<std::remove_reference_t<Callable>>>
267 #ifdef __INVOKE_SIMD_USE_STD_IS_FUNCTION_WA
268  || is_regcall_function_ptr_or_ref_v<Callable>
269 #endif // __INVOKE_SIMD_USE_STD_IS_FUNCTION_WA
270  ;
271 } // namespace detail
272 
273 // --- The main API
274 
291 // TODO works only for functions now, enable for other callables.
292 template <class Callable, class... T>
293 __attribute__((always_inline)) auto invoke_simd(sycl::sub_group sg,
294  Callable &&f, T... args) {
295  // If the invoke_simd call site is fully uniform, then it does not matter
296  // what the subgroup size is and arguments don't need widening and return
297  // value does not need shrinking by this library or SPMD compiler, so 0
298  // is fine in this case.
299  constexpr int N = detail::get_sg_size<Callable, T...>();
300  using RetSpmd = detail::SpmdRetType<N, Callable, T...>;
301  constexpr bool is_function = detail::is_function_ptr_or_ref_v<Callable>;
302 
303  if constexpr (is_function) {
304  return __builtin_invoke_simd<is_function, RetSpmd>(
305  f, detail::unwrap_uniform<T>::impl(args)...);
306  } else {
307  // TODO support functors and lambdas which are handled in this branch.
308  // The limiting factor for now is that the LLVMIR data flow analysis
309  // implemented in LowerInvokeSimd.cpp which, finds actual invoke_simd
310  // target function, can't handle this case yet.
311  return __builtin_invoke_simd<is_function, RetSpmd>(
312  detail::simd_call_helper<N, Callable, T...>, &f,
314  }
315 // TODO Temporary macro and assert to enable API compilation testing.
316 // LowerInvokeSimd.cpp does not support this case yet.
317 #ifndef __INVOKE_SIMD_ENABLE_ALL_CALLABLES
318  static_assert(is_function &&
319  "invoke_simd does not support functors or lambdas yet");
320 #endif // __INVOKE_SIMD_ENABLE_ALL_CALLABLES
321 }
322 
323 } // namespace experimental
324 } // namespace oneapi
325 } // namespace ext
326 } // namespace sycl
simd_mask
Definition: simd.hpp:1029
sub_group.hpp
simd.hpp
SYCL_EXTERNAL
#define SYCL_EXTERNAL
Definition: defines_elementary.hpp:34
sycl::ext::oneapi::experimental::detail::SimdRetType
std::invoke_result_t< SimdCallable, typename spmd2simd< SpmdArgs, N >::type... > SimdRetType
Definition: invoke_simd.hpp:203
sycl::ext::oneapi::experimental::detail::unwrap_uniform
Definition: invoke_simd.hpp:169
uniform.hpp
sycl
Definition: invoke_simd.hpp:68
sycl::ext::oneapi::experimental::detail::has_uniform_signature
Definition: invoke_simd.hpp:152
sycl::ext::oneapi::experimental::detail::is_uniform_type< uniform< T > >::type
T type
Definition: invoke_simd.hpp:140
sycl::ext::oneapi::experimental::detail::sg_size
Definition: invoke_simd.hpp:184
sycl::ext::oneapi::experimental::detail::unwrap_uniform< uniform< T > >::impl
static T impl(uniform< T > val)
Definition: invoke_simd.hpp:175
sycl::ext::oneapi::experimental::detail::SpmdRetType
typename simd2spmd< SimdRetType< N, SimdCallable, SpmdArgs... > >::type SpmdRetType
Definition: invoke_simd.hpp:208
sycl::ext::oneapi::experimental::detail::is_simd_or_mask_type
Definition: invoke_simd.hpp:144
__builtin_invoke_simd
SYCL_EXTERNAL __regcall SpmdRet __builtin_invoke_simd(SimdCallee target, const void *obj, SpmdArgs... args)
Middle End - to - Back End interface to invoke explicit SIMD functions from SPMD SYCL context.
Definition: invoke_simd.hpp:43
sycl::ext::oneapi::experimental::detail::spmd2simd
Definition: invoke_simd.hpp:99
sycl::ext::oneapi::experimental::detail::spmd2simd< uniform< T >, N >::type
T type
Definition: invoke_simd.hpp:102
cl::sycl::access::target
target
Definition: access.hpp:17
sycl::ext::oneapi::experimental::uniform
Definition: uniform.hpp:69
sycl::ext::oneapi::experimental::detail::simd2spmd< simd< T, N > >::type
T type
Definition: invoke_simd.hpp:125
sycl::ext::oneapi::experimental::detail::is_uniform_type
Definition: invoke_simd.hpp:138
simd
Definition: simd.hpp:1027
sycl::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:79
sycl::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:293
std
Definition: accessor.hpp:2617
sycl::ext::oneapi::experimental::detail::simd_call_helper
__regcall detail::SimdRetType< N, Callable, T... > simd_call_helper(const void *obj_ptr, typename detail::spmd2simd< T, N >::type... simd_args)
Definition: invoke_simd.hpp:224
sycl::ext::oneapi::experimental::detail::simd2spmd< std::tuple< T... > >::type
std::tuple< typename simd2spmd< T >::type... > type
Definition: invoke_simd.hpp:129
sycl::ext::oneapi::experimental::detail::get_sg_size
static constexpr int get_sg_size()
Definition: invoke_simd.hpp:211
sycl::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:187
sycl::ext::oneapi::experimental::detail::spmd2simd< std::tuple< T... >, N >::type
std::tuple< typename spmd2simd< T, N >::type... > type
Definition: invoke_simd.hpp:106
sycl::ext::oneapi::experimental::detail::simd2spmd
Definition: invoke_simd.hpp:118
sycl::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:113
sycl::ext::oneapi::experimental::detail::is_function_ptr_or_ref_v
static constexpr bool is_function_ptr_or_ref_v
Definition: invoke_simd.hpp:265
sycl::ext::oneapi::experimental::detail::unwrap_uniform::impl
static auto impl(T val)
Definition: invoke_simd.hpp:170
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24