Unpacking a sycl::item using structured bindings

[ivan-pi]

Is there any plan to support structured bindings (C++17) when using sycl items?

For example to iterate across a 2D range, I’d like to do something like this:

    // Apply to all elements
    q.parallel_for({N,N},[=](auto item) {
        const auto [i, j] = get_linear(item);
        a(i,j) = 8.0;
    });

    // Apply to inner elements 
    q.parallel_for({N-2,N-2},[=](auto item) {
        const auto [i, j] = get_linear_with_offset(item,{1,1});
        a(i,j) = 4.0;
    });

and not have to unpack them individually using the access operator:

std::array<size_t,2> offset{1,1};

const size_t i = item[0] + offset[0];
const size_t j = item[1] + offset[1];

A full example of this in action can be found here: https://godbolt.org/z/enYbzbobz

Edit: The template helper function used to unpack the indexes is,

#include <tuple>
#include <utility>
#include <sycl/sycl.hpp>

template<typename Item, size_t... I>
auto linear_impl_(const Item &it, std::index_sequence<I...>) {
    return std::make_tuple(it[I]...);
}

// Helper function to unpack indices using a structured binding
template<int Dimensions>
inline auto get_linear(sycl::item<Dimensions> it) {
    return linear_impl_(it,std::make_index_sequence<Dimensions>{});
}

[ivan-pi]

This would be desirable also for compatibility with the C++ standard parallelism:

auto v = stdv::cartesian_product(
   stdv::iota(0,N), stdv::iota(0,N));

auto a = mdspan(a_ptr, N, N);

std::for_each(ex::par_unseq, 
    begin(v), end(v), [=](auto idx) {
        auto [i, j] = idx;
        a[i,j] = 8.0;
});

[keryell1]

I am convinced this is useful as this is implemented as an extension in https://github.com/triSYCL/triSYCL/blob/master/include/triSYCL/detail/small_array.hpp#L392 and you might enjoy the test https://github.com/triSYCL/triSYCL/blob/master/tests/id/tuple_like_protocol.cpp

[ivan-pi]

Thanks for the pointer. I wasn’t aware an interface exists for binding tuple-like types.

I suppose one could bind both from id’s and item’s, as both support the operator[]. Since 'id`'s support arithmetic operators, it would make it simple to apply an offset:

auto [i, j] = item.get_id() + id(1,1); // offset

The tuple_size and tuple_element are straightforward:

    template<int Dimensions>
    struct tuple_size<sycl::id<Dimensions>> : 
        std::integral_constant<std::size_t,Dimensions> {};

    template <std::size_t N, int Dimensions>
    struct tuple_element<N,sycl::id<Dimensions>> {
        using type = size_t;
    };

From the compiler error message I’ve inferred that what is missing is a get() function.

    sycl::id<2> idx{1,2};
    auto [i, j] = idx;

<source>:35:11: error: use of undeclared identifier 'get'
    auto [i, j] = idx;

This would need to be a member function of the sycl::id:

    template< std::size_t I, int Dimensions >
    size_t& get(sycl::id<Dimensions>& idx ) noexcept {
        return idx[I];
    }

[keryell1]

In triSYCL, which is a very sparse implementation of SYCL, we experiment a lot of modern things, like mdspan interface, tuple-like interface, etc. Since the tuple-like protocol is implemented by the common small array class, all the SYCL types like id, range, vec, marray, etc. support structure binding.
I have not provided structured binding on item itself since there are actually 2 or 3 small arrays behind the scene and it might not be obvious which one is used. But probably the one you suggest is the more useful.
As the SYCL committee is working on the next version of SYCL, keep suggesting.
If you have suggestions, just open an issue on GitHub - KhronosGroup/SYCL-Docs: SYCL Open Source Specification so it is easier to track.
I guess that some low-hanging fruits extensions like the one you propose will be adopted.

[ivan-pi]

An mdspan interface would be nice; AFAIK only the NVIDIA compilers currently ship a version of mdspan derived from the Kokkos view. The mdspan is probably easier than having users write their own container views like this one (one could go with the buffer accessor model though):

// A custom 2D view of an array
template<typename T>
struct view_2d {
public:
    // Constructor
    view_2d(T *p_, int N_, int M_) : N(N_), M(M_), p(p_) {} 

    // Accesor, i.e. A(i,j)
    inline T &operator()(size_t i, size_t j) const { // <- const is important here
        return p[i*N + j]; 
    }

private:
    int N, M;
    double *p;
};

It took me a while before I realized I need const on the element access operator, to be able to pass the view to a lambda by value.

Would you suggest I just create an issue for the structured bindings for the start, or I also try to make draft a PR with the required changes?