Skip to content

Compilation failure on sycl::joint_reduce with sub_group #8348

Closed
@aleksmesh

Description

@aleksmesh

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.

Metadata

Metadata

Assignees

Labels

CTSImpacts Khronos SYCL CTSbugSomething isn't workingconfirmed

Type

No type

Projects

No projects

Milestone

No milestone

Relationships

None yet

Development

No branches or pull requests

Issue actions