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 
11 #include <sycl/detail/spirv.hpp>
12 #include <sycl/functional.hpp> // for maximum, minimum
13 
14 #include <functional> // for bit_and, bit_or, bit_xor, multiplies
15 
16 namespace sycl {
17 inline namespace _V1 {
18 namespace ext::oneapi {
19 
20 template <typename T = void> using plus = std::plus<T>;
21 template <typename T = void> using multiplies = std::multiplies<T>;
22 template <typename T = void> using bit_or = std::bit_or<T>;
23 template <typename T = void> using bit_xor = std::bit_xor<T>;
24 template <typename T = void> using bit_and = std::bit_and<T>;
25 template <typename T = void> using maximum = sycl::maximum<T>;
26 template <typename T = void> using minimum = sycl::minimum<T>;
27 
28 } // namespace ext::oneapi
29 
30 #ifdef __SYCL_DEVICE_ONLY__
31 namespace detail {
32 
33 struct GroupOpISigned {};
34 struct GroupOpIUnsigned {};
35 struct GroupOpFP {};
36 struct GroupOpC {};
37 struct GroupOpBool {};
38 
39 template <typename T, typename = void> struct GroupOpTag;
40 
41 template <typename T>
42 struct GroupOpTag<T, std::enable_if_t<detail::is_sigeninteger_v<T>>> {
43  using type = GroupOpISigned;
44 };
45 
46 template <typename T>
47 struct GroupOpTag<T, std::enable_if_t<detail::is_sugeninteger_v<T>>> {
48  using type = GroupOpIUnsigned;
49 };
50 
51 template <typename T>
52 struct GroupOpTag<T, std::enable_if_t<detail::is_sgenfloat_v<T>>> {
53  using type = GroupOpFP;
54 };
55 
56 template <typename T>
57 struct GroupOpTag<T, std::enable_if_t<detail::is_genbool_v<T>>> {
58  using type = GroupOpBool;
59 };
60 
61 // GroupOpC (std::complex) is handled in sycl/stl_wrappers/complex.
62 
63 #define __SYCL_CALC_OVERLOAD(GroupTag, SPIRVOperation, BinaryOperation) \
64  template <__spv::GroupOperation O, typename Group, typename T> \
65  static T calc(Group g, GroupTag, T x, BinaryOperation) { \
66  return sycl::detail::spirv::Group##SPIRVOperation<O>(g, x); \
67  }
68 
69 // calc for sycl function objects
70 __SYCL_CALC_OVERLOAD(GroupOpISigned, SMin, sycl::minimum<T>)
71 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, UMin, sycl::minimum<T>)
72 __SYCL_CALC_OVERLOAD(GroupOpFP, FMin, sycl::minimum<T>)
73 
74 __SYCL_CALC_OVERLOAD(GroupOpISigned, SMax, sycl::maximum<T>)
75 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, UMax, sycl::maximum<T>)
76 __SYCL_CALC_OVERLOAD(GroupOpFP, FMax, sycl::maximum<T>)
77 
78 __SYCL_CALC_OVERLOAD(GroupOpISigned, IAdd, sycl::plus<T>)
79 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IAdd, sycl::plus<T>)
80 __SYCL_CALC_OVERLOAD(GroupOpFP, FAdd, sycl::plus<T>)
81 
82 __SYCL_CALC_OVERLOAD(GroupOpISigned, IMul, sycl::multiplies<T>)
83 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, IMul, sycl::multiplies<T>)
84 __SYCL_CALC_OVERLOAD(GroupOpFP, FMul, sycl::multiplies<T>)
85 __SYCL_CALC_OVERLOAD(GroupOpC, CMulINTEL, sycl::multiplies<T>)
86 
87 __SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseOr, sycl::bit_or<T>)
88 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseOr, sycl::bit_or<T>)
89 __SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseXor, sycl::bit_xor<T>)
90 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseXor, sycl::bit_xor<T>)
91 __SYCL_CALC_OVERLOAD(GroupOpISigned, BitwiseAnd, sycl::bit_and<T>)
92 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, BitwiseAnd, sycl::bit_and<T>)
93 
94 __SYCL_CALC_OVERLOAD(GroupOpBool, LogicalAnd, sycl::logical_and<T>)
95 __SYCL_CALC_OVERLOAD(GroupOpISigned, LogicalAnd, sycl::logical_and<T>)
96 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, LogicalAnd, sycl::logical_and<T>)
97 __SYCL_CALC_OVERLOAD(GroupOpFP, LogicalAnd, sycl::logical_and<T>)
98 
99 __SYCL_CALC_OVERLOAD(GroupOpBool, LogicalOr, sycl::logical_or<T>)
100 __SYCL_CALC_OVERLOAD(GroupOpISigned, LogicalOr, sycl::logical_or<T>)
101 __SYCL_CALC_OVERLOAD(GroupOpIUnsigned, LogicalOr, sycl::logical_or<T>)
102 __SYCL_CALC_OVERLOAD(GroupOpFP, LogicalOr, sycl::logical_or<T>)
103 
104 #undef __SYCL_CALC_OVERLOAD
105 
106 template <__spv::GroupOperation O, typename Group, typename T,
107  template <typename> class BinaryOperation>
108 static T calc(Group g, typename GroupOpTag<T>::type, T x,
109  BinaryOperation<void>) {
110  return calc<O>(g, typename GroupOpTag<T>::type(), x, BinaryOperation<T>());
111 }
112 
113 } // namespace detail
114 #endif // __SYCL_DEVICE_ONLY__
115 
116 } // namespace _V1
117 } // namespace sycl
std::plus< T > plus
Definition: functional.hpp:20
std::multiplies< T > multiplies
Definition: functional.hpp:21
std::bit_or< T > bit_or
Definition: functional.hpp:22
std::bit_xor< T > bit_xor
Definition: functional.hpp:23
std::bit_and< T > bit_and
Definition: functional.hpp:24
std::bit_and< T > bit_and
Definition: functional.hpp:20
std::multiplies< T > multiplies
Definition: functional.hpp:19
std::bit_xor< T > bit_xor
Definition: functional.hpp:22
std::plus< T > plus
Definition: functional.hpp:18
std::bit_or< T > bit_or
Definition: functional.hpp:21
Definition: access.hpp:18