DPC++ Runtime
Runtime libraries for oneAPI DPC++
functional.hpp
Go to the documentation of this file.
1 //==----------- functional.hpp --- SYCL functional -------------------------==//
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 
9 #pragma once
10 #include <CL/sycl/functional.hpp>
11 
12 #include <functional>
13 
15 namespace sycl {
16 namespace ext {
17 namespace oneapi {
18 
19 template <typename T = void> using plus = std::plus<T>;
20 template <typename T = void> using multiplies = std::multiplies<T>;
21 template <typename T = void> using bit_or = std::bit_or<T>;
22 template <typename T = void> using bit_xor = std::bit_xor<T>;
23 template <typename T = void> using bit_and = std::bit_and<T>;
24 template <typename T = void> using maximum = sycl::maximum<T>;
25 template <typename T = void> using minimum = sycl::minimum<T>;
26 
27 } // namespace oneapi
28 } // namespace ext
29 
30 #ifdef __SYCL_DEVICE_ONLY__
31 namespace detail {
32 
33 struct GroupOpISigned {};
34 struct GroupOpIUnsigned {};
35 struct GroupOpFP {};
36 
37 template <typename T, typename = void> struct GroupOpTag;
38 
39 template <typename T>
40 struct GroupOpTag<T, detail::enable_if_t<detail::is_sigeninteger<T>::value>> {
41  using type = GroupOpISigned;
42 };
43 
44 template <typename T>
45 struct GroupOpTag<T, detail::enable_if_t<detail::is_sugeninteger<T>::value>> {
46  using type = GroupOpIUnsigned;
47 };
48 
49 template <typename T>
50 struct GroupOpTag<T, detail::enable_if_t<detail::is_sgenfloat<T>::value>> {
51  using type = GroupOpFP;
52 };
53 
54 #define __SYCL_CALC_OVERLOAD(GroupTag, SPIRVOperation, BinaryOperation) \
55  template <typename T, __spv::GroupOperation O, __spv::Scope::Flag S> \
56  static T calc(GroupTag, T x, BinaryOperation) { \
57  using ConvertedT = detail::ConvertToOpenCLType_t<T>; \
58  \
59  using OCLT = \
60  conditional_t<std::is_same<ConvertedT, cl_char>() || \
61  std::is_same<ConvertedT, cl_short>(), \
62  cl_int, \
63  conditional_t<std::is_same<ConvertedT, cl_uchar>() || \
64  std::is_same<ConvertedT, cl_ushort>(), \
65  cl_uint, ConvertedT>>; \
66  OCLT Arg = x; \
67  OCLT Ret = \
68  __spirv_Group##SPIRVOperation(S, static_cast<unsigned int>(O), Arg); \
69  return Ret; \
70  }
71 
72 // calc for sycl function objects
73 __SYCL_CALC_OVERLOAD(GroupOpISigned, SMin, sycl::minimum<T>)
74 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, UMin, sycl::minimum<T>)
75 __SYCL_CALC_OVERLOAD(GroupOpFP, FMin, sycl::minimum<T>)
76 
77 __SYCL_CALC_OVERLOAD(GroupOpISigned, SMax, sycl::maximum<T>)
78 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, UMax, sycl::maximum<T>)
79 __SYCL_CALC_OVERLOAD(GroupOpFP, FMax, sycl::maximum<T>)
80 
81 __SYCL_CALC_OVERLOAD(GroupOpISigned, IAdd, sycl::plus<T>)
82 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IAdd, sycl::plus<T>)
83 __SYCL_CALC_OVERLOAD(GroupOpFP, FAdd, sycl::plus<T>)
84 
85 __SYCL_CALC_OVERLOAD(GroupOpISigned, IMulKHR, sycl::multiplies<T>)
86 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IMulKHR, sycl::multiplies<T>)
87 __SYCL_CALC_OVERLOAD(GroupOpFP, FMulKHR, sycl::multiplies<T>)
88 
89 __SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseOrKHR, sycl::bit_or<T>)
90 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseOrKHR, sycl::bit_or<T>)
91 __SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseXorKHR, sycl::bit_xor<T>)
92 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseXorKHR, sycl::bit_xor<T>)
93 __SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseAndKHR, sycl::bit_and<T>)
94 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseAndKHR, sycl::bit_and<T>)
95 
96 #undef __SYCL_CALC_OVERLOAD
97 
98 template <typename T, __spv::GroupOperation O, __spv::Scope::Flag S,
99  template <typename> class BinaryOperation>
100 static T calc(typename GroupOpTag<T>::type, T x, BinaryOperation<void>) {
101  return calc<T, O, S>(typename GroupOpTag<T>::type(), x, BinaryOperation<T>());
102 }
103 
104 } // namespace detail
105 #endif // __SYCL_DEVICE_ONLY__
106 
107 } // namespace sycl
108 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::ext::oneapi::bit_xor
std::bit_xor< T > bit_xor
Definition: functional.hpp:22
cl::sycl::ext::oneapi::plus
std::plus< T > plus
Definition: functional.hpp:19
cl::sycl::ext::oneapi::bit_or
std::bit_or< T > bit_or
Definition: functional.hpp:21
T
sycl
Definition: invoke_simd.hpp:68
__spv::Scope::Flag
Flag
Definition: spirv_types.hpp:27
cl::sycl::maximum
Definition: functional.hpp:43
cl::sycl::minimum
Definition: functional.hpp:26
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
__spv::GroupOperation
GroupOperation
Definition: spirv_types.hpp:105
cl::sycl::ext::oneapi::bit_and
std::bit_and< T > bit_and
Definition: functional.hpp:23
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
cl::sycl::ext::oneapi::multiplies
std::multiplies< T > multiplies
Definition: functional.hpp:20
functional.hpp
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12