Are range algorithms on buffers functional enough to use?

[denommenator]

Discussed in #1786

^{Originally posted by denommenator August 20, 2024}
I've been trying to run some dpl algorithms on sub-ranges of sycl::buffers but have not been successful using the ranges library to do so. Here is a minimal example of what I'm talking about. The copy using all_view works, but the copies using a drop view in either the read or the write cause a compilation failure. (I'm using the intel/LLVM repo built on linux with a cuda backend).

#include <oneapi/dpl/execution>
#include <oneapi/dpl/iterator>
#include <oneapi/dpl/algorithm>
#include <oneapi/dpl/ranges>

#include <sycl/sycl.hpp>

#include <vector>

int main()
{
  std::vector<int> data = {1, 1, 2, 2, 3, 4, 5, 5, 6, 7, 7};
  std::vector<int> result(5, 0);
  {
    namespace rng = dpl::experimental::ranges;
    sycl::buffer dataB(data);
    sycl::buffer resultB(result);
    
    auto data_view = rng::views::all_read(dataB);
    auto result_view = rng::views::all_write(resultB);

    auto data_drop_view = data_view | rng::views::drop(1);
    auto result_drop_view = result_view | rng::views::drop(1);

    rng::copy(dpl::execution::dpcpp_default, data_view, result_view);
    //rng::copy(dpl::execution::dpcpp_default, data_drop_view, result_view); //compilation failure
    //rng::copy(dpl::execution::dpcpp_default, data_view, result_drop_view); //compilation failure
  }
}

Here is the compilation error output when uncommenting the result_drop_view copy.

/home/robert-denomme/sycl_workspace/sycl_sandbox/IceSYCL/extern/oneDPL/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h:241:17: error: no matching function for call to object of type 'const oneapi::dpl::unseq_backend::walk_n<const oneapi::dpl::execution::device_policy<> &, oneapi::dpl::__internal::__brick_copy<oneapi::dpl::__internal::__hetero_tagoneapi::dpl::__internal::__device_backend_tag, const oneapi::dpl::execution::device_policy<> &>>'
241 | __brick(__idx, __rngs...);
| ^~~~~~~
/home/robert-denomme/sycl_workspace/sycl_sandbox/IceSYCL/extern/oneDPL/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h:239:112: note: while substituting into a lambda expression here
239 | __cgh.parallel_for<_Name...>(sycl::range</dim=/1>(__count), [=](sycl::item</dim=/1> __item_id) {
| ^
/home/robert-denomme/sycl_workspace/sycl_sandbox/IceSYCL/extern/oneDPL/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h:235:100: note: while substituting into a lambda expression here
235 | auto __event = __exec.queue().submit([&__rngs..., &__brick, __count](sycl::handler& __cgh) {
| ^
/home/robert-denomme/sycl_workspace/sycl_sandbox/IceSYCL/extern/oneDPL/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl.h:258:12: note: in instantiation of function template specialization 'oneapi::dpl::__par_backend_hetero::__parallel_for_submitter<oneapi::dpl::__par_backend_hetero::__internal::__optional_kernel_name<>>::operator()<const oneapi::dpl::execution::device_policy<> &, oneapi::dpl::unseq_backend::walk_n<const oneapi::dpl::execution::device_policy<> &, oneapi::dpl::__internal::__brick_copy<oneapi::dpl::__internal::__hetero_tagoneapi::dpl::__internal::__device_backend_tag, const oneapi::dpl::execution::device_policy<> &>>, long, oneapi::dpl::__ranges::all_view<int, sycl::access::mode::read, sycl::access::target::global_buffer, sycl::access::placeholder::true_t> &, __nanorange::nano::drop_view<oneapi::dpl::__ranges::all_view<int, sycl::access::mode::write, sycl::access::target::global_buffer, sycl::access::placeholder::true_t>> &>' requested here
258 | return __parallel_for_submitter<_ForKernel>()(::std::forward<_ExecutionPolicy>(__exec), __brick, __count,
| ^
/home/robert-denomme/sycl_workspace/sycl_sandbox/IceSYCL/extern/oneDPL/include/oneapi/dpl/internal/../pstl/hetero/algorithm_ranges_impl_hetero.h:49:44: note: in instantiation of function template specialization 'oneapi::dpl::__par_backend_hetero::__parallel_for<const oneapi::dpl::execution::device_policy<> &, oneapi::dpl::unseq_backend::walk_n<const oneapi::dpl::execution::device_policy<> &, oneapi::dpl::__internal::__brick_copy<oneapi::dpl::__internal::__hetero_tagoneapi::dpl::__internal::__device_backend_tag, const oneapi::dpl::execution::device_policy<> &>>, long, oneapi::dpl::__ranges::all_view<int, sycl::access::mode::read, sycl::access::target::global_buffer, sycl::access::placeholder::true_t> &, __nanorange::nano::drop_view<oneapi::dpl::__ranges::all_view<int, sycl::access::mode::write, sycl::access::target::global_buffer, sycl::access::placeholder::true_t>> &>' requested here
49 | oneapi::dpl::__par_backend_hetero::__parallel_for(_BackendTag{}, ::std::forward<_ExecutionPolicy>(__exec),
| ^
/home/robert-denomme/sycl_workspace/sycl_sandbox/IceSYCL/extern/oneDPL/include/oneapi/dpl/internal/../pstl/glue_algorithm_ranges_impl.h:254:40: note: in instantiation of function template specialization 'oneapi::dpl::__internal::__ranges::__pattern_walk_n<oneapi::dpl::__internal::__device_backend_tag, const oneapi::dpl::execution::device_policy<> &, oneapi::dpl::__internal::__brick_copy<oneapi::dpl::__internal::__hetero_tagoneapi::dpl::__internal::__device_backend_tag, const oneapi::dpl::execution::device_policy<> &>, oneapi::dpl::__ranges::all_view<int, sycl::access::mode::read, sycl::access::target::global_buffer, sycl::access::placeholder::true_t> &, __nanorange::nano::drop_view<oneapi::dpl::__ranges::all_view<int, sycl::access::mode::write, sycl::access::target::global_buffer, sycl::access::placeholder::true_t>> &>' requested here
254 | oneapi::dpl::__internal::__ranges::__pattern_walk_n(
| ^
/home/robert-denomme/sycl_workspace/sycl_sandbox/IceSYCL/tests/particle_node_interactions.cpp:143:14: note: in instantiation of function template specialization 'oneapi::dpl::experimental::ranges::copy<const oneapi::dpl::execution::device_policy<> &, oneapi::dpl::__ranges::all_view<int, sycl::access::mode::read, sycl::access::target::global_buffer, sycl::access::placeholder::true_t> &, __nanorange::nano::drop_view<oneapi::dpl::__ranges::all_view<int, sycl::access::mode::write, sycl::access::target::global_buffer, sycl::access::placeholder::true_t>> &>' requested here
143 | rng::copy(dpl::execution::dpcpp_default, data_view, result_drop_view);
| ^
/home/robert-denomme/sycl_workspace/sycl_sandbox/IceSYCL/extern/oneDPL/include/oneapi/dpl/pstl/hetero/dpcpp/unseq_backend_sycl.h:106:5: note: candidate template ignored: substitution failure [with _ItemId = size_t, _Ranges = <const oneapi::dpl::__ranges::all_view<int, sycl::access::mode::read, sycl::access::target::global_buffer, sycl::access::placeholder::true_t> &, const __nanorange::nano::drop_view<oneapi::dpl::__ranges::all_view<int, sycl::access::mode::write, sycl::access::target::global_buffer, sycl::access::placeholder::true_t>> &>]: no viable overloaded operator[] for type 'const __nanorange::nano::drop_view<oneapi::dpl::__ranges::all_view<int, sycl::access::mode::write, sycl::access::target::global_buffer, sycl::access::placeholder::true_t>>'
106 | operator()(const _ItemId __idx, _Ranges&&... __rngs) const -> decltype(__f(__rngs[__idx]...))
| ^

[MikeDvorskiy]

@denommenator , please find a fix here
#1839