16 #define SYCL_EXT_ONEAPI_INVOKE_SIMD 1
21 #include <sycl/detail/boost/mp11.hpp>
42 template <
bool IsFunc,
class SpmdRet,
class HelperFunc,
43 class... UserSimdFuncAndSpmdArgs,
class = std::enable_if_t<!IsFunc>>
46 UserSimdFuncAndSpmdArgs... args)
47 #ifdef __SYCL_DEVICE_ONLY__
52 throw sycl::exception(sycl::errc::feature_not_supported,
53 "__builtin_invoke_simd is not supported on host");
57 template <
bool IsFunc,
class SpmdRet,
class HelperFunc,
58 class... UserSimdFuncAndSpmdArgs,
class = std::enable_if_t<IsFunc>>
61 #ifdef __SYCL_DEVICE_ONLY__
66 throw sycl::exception(sycl::errc::feature_not_supported,
67 "__builtin_invoke_simd is not supported on host");
74 namespace ext::oneapi::experimental {
80 template <
class T,
int N>
82 std::experimental::_StorageKind::_VecExt, N>;
87 template <
class T,
int N>
88 using simd = std::experimental::simd<T, simd_abi::native_fixed_size<T, N>>;
91 template <
class T,
int N>
93 std::experimental::simd_mask<T, simd_abi::native_fixed_size<T, N>>;
98 namespace __MP11_NS = sycl::detail::boost::mp11;
102 template <
class T,
int N,
class =
void>
struct spmd2simd;
109 using type = std::tuple<typename spmd2simd<T, N>::type...>;
114 template <
class T,
int N>
132 using type = std::tuple<typename simd2spmd<T>::type...>;
144 static constexpr
int value = 1;
149 static constexpr
int value = N;
160 template <
class T,
int N>
162 template <
class T,
int N>
167 constexpr
operator bool() {
168 using TypeList = __MP11_NS::mp_list<SpmdArgs...>;
169 return __MP11_NS::mp_all_of<TypeList, is_uniform_type>::value;
177 static auto impl(T val) {
return val; }
191 template <
class SimdCallable,
class... SpmdArgs>
struct sg_size {
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;
207 template <
int N,
class SimdCallable,
class... SpmdArgs>
209 std::invoke_result_t<SimdCallable,
213 template <
int N,
class SimdCallable,
class... SpmdArgs>
217 template <
class SimdCallable,
class... SpmdArgs>
220 using SimdRet = std::invoke_result_t<SimdCallable, SpmdArgs...>;
229 return sg_size<SimdCallable, SpmdArgs...>();
236 template <
int N,
class Callable,
class... T>
237 [[intel::device_indirectly_callable]]
SYCL_EXTERNAL __regcall detail::
238 SimdRetType<N, Callable, T...>
242 *
reinterpret_cast<const std::remove_reference_t<Callable> *
>(obj_ptr);
243 return f(simd_args...);
247 template <
int N,
class Callable,
class... T>
248 [[intel::device_indirectly_callable]]
SYCL_EXTERNAL __regcall detail::
249 SimdRetType<N, Callable, T...>
252 return f(simd_args...);
255 #ifdef _GLIBCXX_RELEASE
256 #if _GLIBCXX_RELEASE < 10
257 #define __INVOKE_SIMD_USE_STD_IS_FUNCTION_WA
261 #ifdef __INVOKE_SIMD_USE_STD_IS_FUNCTION_WA
269 template <
class F>
struct is_regcall_function_ptr_or_ref : std::false_type {};
271 template <
class Ret,
class... Args>
272 struct is_regcall_function_ptr_or_ref<Ret(__regcall &)(Args...)>
275 template <
class Ret,
class... Args>
276 struct is_regcall_function_ptr_or_ref<Ret(__regcall *)(Args...)>
279 template <
class Ret,
class... Args>
280 struct is_regcall_function_ptr_or_ref<Ret(__regcall *&)(Args...)>
284 static constexpr
bool is_regcall_function_ptr_or_ref_v =
285 is_regcall_function_ptr_or_ref<F>::value;
288 template <
class Callable>
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>
300 template <
typename Ret,
typename... Args>
302 using type = Ret(__regcall *)(Args...);
305 template <
typename T>
331 template <
class Callable,
class... T>
333 Callable &&f, T... args) {
340 constexpr
bool is_function = detail::is_function_ptr_or_ref_v<Callable>;
342 if constexpr (is_function) {
365 #ifndef __INVOKE_SIMD_ENABLE_ALL_CALLABLES
366 static_assert(is_function &&
367 "invoke_simd does not support functors or lambdas yet");
#define __SYCL_INLINE_VER_NAMESPACE(X)
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.
typename std::enable_if< B, T >::type enable_if_t
std::invoke_result_t< SimdCallable, typename spmd2simd< SpmdArgs, N >::type... > SimdRetType
typename simd2spmd< SimdRetType< N, SimdCallable, SpmdArgs... > >::type SpmdRetType
static constexpr bool is_function_ptr_or_ref_v
SYCL_EXTERNAL __regcall detail::SimdRetType< N, Callable, T... > simd_obj_call_helper(const void *obj_ptr, typename detail::spmd2simd< T, N >::type... simd_args)
typename remove_ref_from_func_ptr_ref_type< T >::type remove_ref_from_func_ptr_ref_type_t
SYCL_EXTERNAL __regcall detail::SimdRetType< N, Callable, T... > simd_func_call_helper(Callable f, typename detail::spmd2simd< T, N >::type... simd_args)
static constexpr int get_sg_size()
typename std::experimental::__simd_abi< std::experimental::_StorageKind::_VecExt, N > native_fixed_size
__attribute__((always_inline)) auto invoke_simd(sycl
The invoke_simd free function invokes a SIMD function using all work-items in a sub_group.
---— Error handling, matching OpenCL plugin semantics.
Ret(__regcall *)(Args...) type
__MP11_NS::mp_bool< std::is_invocable_v< SimdCallable, typename spmd2simd< SpmdArgs, N::value >::type... > > IsInvocableSgSize
std::tuple< typename simd2spmd< T >::type... > type
std::tuple< typename spmd2simd< T, N >::type... > type