16 inline namespace _V1 {
17 namespace ext::oneapi::experimental {
19 template <
typename GroupHelper,
typename T,
typename BinaryOperation>
21 BinaryOperation binary_op) {
22 #ifdef __SYCL_DEVICE_ONLY__
23 T *Memory =
reinterpret_cast<T *
>(group_helper.get_memory().data());
24 auto g = group_helper.get_group();
25 Memory[g.get_local_linear_id()] =
x;
29 for (
int i = 1; i < num_elements; i++) {
30 result = binary_op(result, Memory[i]);
36 std::ignore = group_helper;
38 std::ignore = num_elements;
39 std::ignore = binary_op;
41 "Group algorithms are not supported on host.");
47 template <
typename GroupHelper,
typename T,
typename BinaryOperation>
48 std::enable_if_t<(is_group_helper_v<GroupHelper>), T>
50 if constexpr (sycl::detail::is_native_op<T, BinaryOperation>::value) {
53 #ifdef __SYCL_DEVICE_ONLY__
55 group_helper,
x, group_helper.get_group().get_local_linear_range(),
59 "Group algorithms are not supported on host.");
63 template <
typename GroupHelper,
typename V,
typename T,
64 typename BinaryOperation>
65 std::enable_if_t<(is_group_helper_v<GroupHelper>), T>
67 BinaryOperation binary_op) {
68 if constexpr (sycl::detail::is_native_op<V, BinaryOperation>::value &&
69 sycl::detail::is_native_op<T, BinaryOperation>::value) {
73 #ifdef __SYCL_DEVICE_ONLY__
76 std::ignore = group_helper;
78 "Group algorithms are not supported on host.");
83 template <
typename GroupHelper,
typename Ptr,
typename BinaryOperation>
84 std::enable_if_t<(is_group_helper_v<GroupHelper> &&
85 sycl::detail::is_pointer_v<Ptr>),
88 BinaryOperation binary_op) {
89 if constexpr (sycl::detail::is_native_op<
91 BinaryOperation>::value) {
94 #ifdef __SYCL_DEVICE_ONLY__
99 auto g = group_helper.get_group();
100 T partial = *(first + g.get_local_linear_id());
101 Ptr second = first + g.get_local_linear_range();
103 [&](
const T &
x) { partial = binary_op(partial,
x); });
105 size_t num_elements = last - first;
106 num_elements = std::min(num_elements, g.get_local_linear_range());
110 std::ignore = group_helper;
113 std::ignore = binary_op;
115 "Group algorithms are not supported on host.");
119 template <
typename GroupHelper,
typename Ptr,
typename T,
120 typename BinaryOperation>
122 (is_group_helper_v<GroupHelper> && sycl::detail::is_pointer_v<Ptr>), T>
124 BinaryOperation binary_op) {
125 if constexpr (sycl::detail::is_native_op<T, BinaryOperation>::value) {
129 #ifdef __SYCL_DEVICE_ONLY__
130 return binary_op(init,
joint_reduce(group_helper, first, last, binary_op));
132 std::ignore = group_helper;
135 "Group algorithms are not supported on host.");
Function for_each(Group g, Ptr first, Ptr last, Function f)
T reduce_over_group_impl(GroupHelper group_helper, T x, size_t num_elements, BinaryOperation binary_op)
std::enable_if_t<(is_group_helper_v< GroupHelper >), T > reduce_over_group(GroupHelper group_helper, T x, BinaryOperation binary_op)
std::enable_if_t<(is_group_helper_v< GroupHelper > &&sycl::detail::is_pointer_v< Ptr >), typename std::iterator_traits< Ptr >::value_type > joint_reduce(GroupHelper group_helper, Ptr first, Ptr last, BinaryOperation binary_op)
std::enable_if_t<(is_group_v< std::decay_t< Group >> &&(std::is_trivially_copyable_v< T >||detail::is_vec< T >::value)), T > group_broadcast(Group g, T x, typename Group::id_type local_id)
std::enable_if_t<(is_group_v< std::decay_t< Group >> &&detail::is_pointer_v< Ptr > &&detail::is_arithmetic_or_complex< typename detail::remove_pointer< Ptr >::type >::value &&detail::is_arithmetic_or_complex< T >::value &&detail::is_plus_or_multiplies_if_complex< T, BinaryOperation >::value &&detail::is_native_op< T, BinaryOperation >::value), T > joint_reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op)
void group_barrier(ext::oneapi::experimental::root_group< dimensions > G, memory_scope FenceScope=decltype(G)::fence_scope)
std::enable_if_t<(is_group_v< std::decay_t< Group >> &&(detail::is_scalar_arithmetic< T >::value||(detail::is_complex< T >::value &&detail::is_multiplies< T, BinaryOperation >::value)) &&detail::is_native_op< T, BinaryOperation >::value), T > reduce_over_group(Group g, T x, BinaryOperation binary_op)
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()