KhronosGroup / SYCL-Docs

SYCL Open Source Specification
Other
111 stars 67 forks source link

Load on vec does not allow pointer to non-const data. #540

Open steffenlarsen opened 6 months ago

steffenlarsen commented 6 months ago

The load member function on sycl::vec takes a multi_ptr<const DataT, ...> which seems to written that way to allow pointers with and without const data. However, because the function is dependent, the conversion from multi_ptr<DataT, ...> to multi_ptr<const DataT, ...> cannot be resolved. See https://godbolt.org/z/dTdc6K7a9.

Possible solution: Make two overloads of load - one with multi_ptr<std::remove_const_t<DataT>, ...> and one with multi_ptr<std::add_const_t<DataT>, ...>. See https://godbolt.org/z/xs175dc7v.

Inspired by https://github.com/KhronosGroup/SYCL-CTS/pull/870.

gmlueck commented 6 months ago

We had a similar issue with async_work_group_copy. In that case, we added another template parameter and a constraint. If we wanted to follow the same strategy with vec::load, I think we would add a template parameter SrcDataT and then add a constraint like:

Available only when: (std::is_same_v<DataT, std::remove_const_t<SrcDataT>> == true)

It seems like we would want to make a similar change to vec::store, adding a template parameter DestDataT and a constraint like:

Available only when: (std::is_same_v<std::remove_const_t<DataT>, DestDataT>> == true)

@AerialMantis, I think you were looking at the "const" issues with mult_ptr a little while ago. Was this an oversight with vec::load / vec::store, or was there a reason to keep it this way?

victor-eds commented 6 months ago

I'm wondering whether this should indeed be solved in sycl::multi_ptr instead. Issue to me is that we diverge from C++ pointers in the sense that we miss implicit cv-qualification of sycl::multi_ptr, i.e., we should add:

template <typename OtherElementType,
          typename = std::enable_if_t<std::is_same_v<OtherElementType, std::add_const_t<ElementType>> ||
                                      std::is_same_v<OtherElementType, std::add_volatile_t<ElementType>>>>
operator multi_ptr<OtherElementType, Space, DecorateAddress>() const;

Users would expect to be able to call a function receiving const T * with T *, not sure why that's not also the case with sycl::multi_ptr. WDYT?

gmlueck commented 6 months ago

@victor-eds: We already have these implicit conversions in multi_ptr, but it doesn't solve the problem. This is because template selection happens before implicit conversions are considered. Here is a slightly modified version of @steffenlarsen's godbolt example that demonstrates the problem (https://godbolt.org/z/5Pvo13M31).

victor-eds commented 6 months ago

True, then the two overloads look like a better solution. Omit my comment :+1: