Closed
Description
Summary
Compilation failure on sycl::joint_reduce
with sycl::sub_group
Although in SYCL2020 specification 4.17.4.5. reduce there is no restriction to use sycl::sub_group
as Group
:
template <typename Group, typename Ptr, typename BinaryOperation>
std::iterator_traits<Ptr>::value_type
joint_reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op);
1. Constraints: Available only if sycl::is_group_v<std::decay_t<Group>> is true, Ptr is a
pointer to a fundamental type, and BinaryOperation is a SYCL function object type.
Mandates: binary_op(*first, *first) must return a value of type std::iterator_traits<Ptr>::value_type.
Preconditions: first, last and the type of binary_op must be the same for all work-items
in the group. binary_op must be an instance of a SYCL function object.
Returns: The result of combining the values resulting from dereferencing all iterators
in the range [first, last) using the operator binary_op, where the values are combined
according to the generalized sum defined in standard C++.
To reproduce
#include <stdint.h>
#include <sycl.hpp>
#include <variant>
#include <iostream>
#include <algorithm>
int main() {
using AccT = sycl::accessor<int, 1>;
int v[256] = {};
std::fill_n(v, 255, 1);
constexpr size_t global_range_size = 256;
constexpr size_t local_range_size = 2;
sycl::range<1> global_range(global_range_size);
sycl::range<1> local_range(local_range_size);
{
sycl::buffer<int,1> v_buf(v, global_range);
sycl::nd_range<1> nd_range(global_range, local_range);
sycl::queue queue;
queue
.submit([&](sycl::handler &cgh) {
AccT v_acc( v_buf, cgh);
cgh.parallel_for(nd_range, [=](sycl::nd_item<1> item) {
auto group = item.get_group();
auto sub_group = item.get_sub_group();
int* v_begin = v_acc.get_pointer();
int* v_end = v_begin + v_acc.size();
auto op_plus = sycl::plus<int>();
sycl::joint_reduce( sub_group, v_begin, v_end, op_plus );
});
})
.wait_and_throw();
}
return EXIT_SUCCESS;
}
To build
clang++ -fsycl ./sample_joint_reduce.cpp -o sample
Output
In file included from ./sample_joint_reduce.cpp:3:
In file included from /opt/sycl/bin/../include/sycl/sycl.hpp:29:
In file included from /opt/sycl/bin/../include/sycl/ext/oneapi/experimental/group_sort.hpp:13:
In file included from /opt/sycl/bin/../include/sycl/detail/group_sort_impl.hpp:15:
/opt/sycl/bin/../include/sycl/group_algorithm.hpp:306:10: error: call to function 'joint_reduce' that is neither visible in the template definition nor found by argument-dependent lookup
return joint_reduce(g, first, last, init, binary_op);
^
./sample_joint_reduce.cpp:36:10: note: in instantiation of function template specialization 'sycl::joint_reduce<sycl::ext::oneapi::sub_group, int *, std::plus<int>>' requested here
sycl::joint_reduce( sub_group, v_begin, v_end, op_plus );
^
/opt/sycl/bin/../include/sycl/group_algorithm.hpp:330:1: note: 'joint_reduce' should be declared prior to the call site or in namespace 'sycl::ext::oneapi'
joint_reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op) {
^
1 error generated.