15 inline namespace _V1 {
16 namespace ext::oneapi::experimental {
18 template <
typename GroupHelper,
typename T,
typename BinaryOperation>
20 BinaryOperation binary_op) {
21 #ifdef __SYCL_DEVICE_ONLY__
22 T *Memory =
reinterpret_cast<T *
>(group_helper.get_memory().data());
23 auto g = group_helper.get_group();
24 Memory[g.get_local_linear_id()] =
x;
28 for (
int i = 1; i < num_elements; i++) {
29 result = binary_op(result, Memory[i]);
35 std::ignore = group_helper;
37 std::ignore = num_elements;
38 std::ignore = binary_op;
39 throw runtime_error(
"Group algorithms are not supported on host.",
40 PI_ERROR_INVALID_DEVICE);
46 template <
typename GroupHelper,
typename T,
typename BinaryOperation>
47 std::enable_if_t<(is_group_helper_v<GroupHelper>), T>
49 if constexpr (sycl::detail::is_native_op<T, BinaryOperation>::value) {
52 #ifdef __SYCL_DEVICE_ONLY__
54 group_helper,
x, group_helper.get_group().get_local_linear_range(),
57 throw runtime_error(
"Group algorithms are not supported on host.",
58 PI_ERROR_INVALID_DEVICE);
62 template <
typename GroupHelper,
typename V,
typename T,
63 typename BinaryOperation>
64 std::enable_if_t<(is_group_helper_v<GroupHelper>), T>
66 BinaryOperation binary_op) {
67 if constexpr (sycl::detail::is_native_op<V, BinaryOperation>::value &&
68 sycl::detail::is_native_op<T, BinaryOperation>::value) {
72 #ifdef __SYCL_DEVICE_ONLY__
75 std::ignore = group_helper;
76 throw runtime_error(
"Group algorithms are not supported on host.",
77 PI_ERROR_INVALID_DEVICE);
82 template <
typename GroupHelper,
typename Ptr,
typename BinaryOperation>
83 std::enable_if_t<(is_group_helper_v<GroupHelper> &&
84 sycl::detail::is_pointer_v<Ptr>),
87 BinaryOperation binary_op) {
88 if constexpr (sycl::detail::is_native_op<
90 BinaryOperation>::value) {
93 #ifdef __SYCL_DEVICE_ONLY__
98 auto g = group_helper.get_group();
99 T partial = *(first + g.get_local_linear_id());
100 Ptr second = first + g.get_local_linear_range();
102 [&](
const T &
x) { partial = binary_op(partial,
x); });
104 size_t num_elements = last - first;
105 num_elements = std::min(num_elements, g.get_local_linear_range());
109 std::ignore = group_helper;
112 std::ignore = binary_op;
113 throw runtime_error(
"Group algorithms are not supported on host.",
114 PI_ERROR_INVALID_DEVICE);
118 template <
typename GroupHelper,
typename Ptr,
typename T,
119 typename BinaryOperation>
121 (is_group_helper_v<GroupHelper> && sycl::detail::is_pointer_v<Ptr>), T>
123 BinaryOperation binary_op) {
124 if constexpr (sycl::detail::is_native_op<T, BinaryOperation>::value) {
128 #ifdef __SYCL_DEVICE_ONLY__
129 return binary_op(init,
joint_reduce(group_helper, first, last, binary_op));
131 std::ignore = group_helper;
133 throw runtime_error(
"Group algorithms are not supported on host.",
134 PI_ERROR_INVALID_DEVICE);
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)