Skip to content

Commit

Permalink
[SYCL] Properties must be a properties list in group_load/store (#…
Browse files Browse the repository at this point in the history 
…16422)

Before this patch we were failing with more obscure errors like

```
no member named 'has_property' in ...property_value<...data_placement_key...>
```
  • Loading branch information
@aelovikov-intel
aelovikov-intel authored Dec 23, 2024
1 parent 51008bf commit 15929c6
Show file tree
Hide file tree
Showing 5 changed files with 150 additions and 117 deletions.
30 changes: 21 additions & 9 deletions sycl/doc/extensions/experimental/sycl_ext_oneapi_group_load_store.asciidoc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -99,6 +99,7 @@ in the group.
* Value type of `InputIteratorT` must be convertible to `OutputT`.
* Value type of `InputIteratorT` and `OutputT` must be trivially copyable
and default constructible.
* `Properties` is an instance of `sycl::ext::oneapi::experimental::properties`

_Effects_: Loads single element from `in_iter` to `out` by using the `g` group
object to identify memory location as `in_iter` + `g.get_local_linear_id()`.
Expand Down Expand Up @@ -129,6 +130,7 @@ in the group.
* Value type of `InputIteratorT` must be convertible to `OutputT`.
* Value type of `InputIteratorT` and `OutputT` must be trivially copyable
and default constructible.
* `Properties` is an instance of `sycl::ext::oneapi::experimental::properties`

_Effects_: Loads `N` elements from `in_iter` to `out`
using the `g` group object.
Expand Down Expand Up @@ -165,6 +167,7 @@ work-group or sub-group.
* Value type of `InputIteratorT` must be convertible to `OutputT`.
* Value type of `InputIteratorT` and `OutputT` must be trivially copyable
and default constructible.
* `Properties` is an instance of `sycl::ext::oneapi::experimental::properties`

_Effects_: Loads `ElementsPerWorkItem` elements from `in_iter` to `out`
using the `g` group object.
Expand Down Expand Up @@ -204,6 +207,7 @@ in the group.
* `InputT` must be convertible to value type of `OutputIteratorT`.
* `InputT` and value type of `OutputIteratorT` must be trivially copyable
and default constructible.
* `Properties` is an instance of `sycl::ext::oneapi::experimental::properties`

_Effects_: Stores single element `in` to `out_iter` by using the `g` group
object to identify memory location as `out_iter` + `g.get_local_linear_id()`
Expand Down Expand Up @@ -235,6 +239,7 @@ in the group.
* `InputT` must be convertible to value type of `OutputIteratorT`.
* `InputT` and value type of `OutputIteratorT` must be trivially copyable
and default constructible.
* `Properties` is an instance of `sycl::ext::oneapi::experimental::properties`

_Effects_: Stores `N` elements from `in` vec to `out_iter`
using the `g` group object.
Expand Down Expand Up @@ -273,6 +278,7 @@ work-group or sub-group.
* `InputT` must be convertible to value type of `OutputIteratorT`.
* `InputT` and value type of `OutputIteratorT` must be trivially copyable
and default constructible.
* `Properties` is an instance of `sycl::ext::oneapi::experimental::properties`

_Effects_: Stores `ElementsPerWorkItem` elements from `in` span to `out_iter`
using the `g` group object.
Expand Down Expand Up @@ -370,7 +376,7 @@ Specifies data layout used in group_load/store for `sycl::vec` or fixed-size
arrays functions.

Example:
`group_load(g, input, output_span, data_placement_blocked);`
`group_load(g, input, output_span, properties{data_placement_blocked});`

=== Optimization Properties

Expand Down Expand Up @@ -398,7 +404,7 @@ inline constexpr contiguous_memory_key::value_t contiguous_memory;
----

For example, we can assert that `input` is a contiguous iterator:
`group_load(g, input, output_span, contiguous_memory);`
`group_load(g, input, output_span, properties{contiguous_memory});`

If `input` isn't a contiguous iterator, the behavior is undefined.

Expand Down Expand Up @@ -432,7 +438,7 @@ inline constexpr full_group_key::value_t full_group;

For example, we can assert that there is no uneven group partition,
so the implementation can rely on `get_max_local_range()` range size:
`group_load(sg, input, output_span, full_group);`
`group_load(sg, input, output_span, properties{full_group});`

If partition is uneven the behavior is undefined.

Expand Down Expand Up @@ -466,11 +472,13 @@ q.submit([&](sycl::handler& cgh) {
auto offset = g.get_group_id(0) * g.get_local_range(0) *
items_per_thread;
sycl_exp::group_load(g, input + offset, sycl::span{ data }, sycl_exp::contiguous_memory);
auto props = sycl_exp::properties{sycl_exp::contiguous_memory};
sycl_exp::group_load(g, input + offset, sycl::span{ data }, props);
// Work with data...
sycl_exp::group_store(g, output + offset, sycl::span{ data }, sycl_exp::contiguous_memory);
sycl_exp::group_store(g, output + offset, sycl::span{ data }, props);
});
});
----
Expand Down Expand Up @@ -546,11 +554,13 @@ q.submit([&](sycl::handler& cgh) {
sycl_exp::group_with_scratchpad gh{ g,
sycl::span{ buf_ptr, temp_memory_size } };
sycl_exp::group_load(gh, input + offset, sycl::span{ data }, sycl_exp::contiguous_memory);
auto props = sycl_exp::properties{sycl_exp::contiguous_memory};
sycl_exp::group_load(gh, input + offset, sycl::span{ data }, props);
// Work with data...
sycl_exp::group_store(gh, output + offset, sycl::span{ data }, sycl_exp::contiguous_memory);
sycl_exp::group_store(gh, output + offset, sycl::span{ data }, props);
});
});
----
Expand Down Expand Up @@ -583,11 +593,13 @@ q.submit([&](sycl::handler& cgh) {
sycl_exp::group_with_scratchpad gh{ g,
sycl::span{ buf_ptr, temp_memory_size } };
sycl_exp::group_load(gh, input + offset, sycl::span{ data }, sycl_exp::data_placement_striped);
auto striped = sycl_exp::properties{sycl_exp::data_placement_striped};
sycl_exp::group_load(gh, input + offset, sycl::span{ data }, striped);
// Work with data...
sycl_exp::group_store(gh, output + offset, sycl::span{ data }, sycl_exp::data_placement_striped);
sycl_exp::group_store(gh, output + offset, sycl::span{ data }, striped);
});
});
----
Expand Down
18 changes: 12 additions & 6 deletions sycl/include/sycl/ext/oneapi/experimental/group_load_store.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -233,7 +233,8 @@ template <typename Group, typename InputIteratorT, typename OutputT,
std::size_t ElementsPerWorkItem,
typename Properties = decltype(properties())>
std::enable_if_t<detail::verify_load_types<InputIteratorT, OutputT> &&
detail::is_generic_group_v<Group>>
detail::is_generic_group_v<Group> &&
is_property_list_v<Properties>>
group_load(Group g, InputIteratorT in_ptr,
span<OutputT, ElementsPerWorkItem> out, Properties props = {}) {
constexpr bool blocked = detail::isBlocked(props);
Expand Down Expand Up @@ -305,7 +306,8 @@ template <typename Group, typename InputT, std::size_t ElementsPerWorkItem,
typename OutputIteratorT,
typename Properties = decltype(properties())>
std::enable_if_t<detail::verify_store_types<InputT, OutputIteratorT> &&
detail::is_generic_group_v<Group>>
detail::is_generic_group_v<Group> &&
is_property_list_v<Properties>>
group_store(Group g, const span<InputT, ElementsPerWorkItem> in,
OutputIteratorT out_ptr, Properties props = {}) {
constexpr bool blocked = detail::isBlocked(props);
Expand Down Expand Up @@ -352,7 +354,8 @@ group_store(Group g, const span<InputT, ElementsPerWorkItem> in,
template <typename Group, typename InputIteratorT, typename OutputT,
typename Properties = decltype(properties())>
std::enable_if_t<detail::verify_load_types<InputIteratorT, OutputT> &&
detail::is_generic_group_v<Group>>
detail::is_generic_group_v<Group> &&
is_property_list_v<Properties>>
group_load(Group g, InputIteratorT in_ptr, OutputT &out,
Properties properties = {}) {
group_load(g, in_ptr, span<OutputT, 1>(&out, 1), properties);
Expand All @@ -362,7 +365,8 @@ group_load(Group g, InputIteratorT in_ptr, OutputT &out,
template <typename Group, typename InputT, typename OutputIteratorT,
typename Properties = decltype(properties())>
std::enable_if_t<detail::verify_store_types<InputT, OutputIteratorT> &&
detail::is_generic_group_v<Group>>
detail::is_generic_group_v<Group> &&
is_property_list_v<Properties>>
group_store(Group g, const InputT &in, OutputIteratorT out_ptr,
Properties properties = {}) {
group_store(g, span<const InputT, 1>(&in, 1), out_ptr, properties);
Expand All @@ -372,7 +376,8 @@ group_store(Group g, const InputT &in, OutputIteratorT out_ptr,
template <typename Group, typename InputIteratorT, typename OutputT, int N,
typename Properties = decltype(properties())>
std::enable_if_t<detail::verify_load_types<InputIteratorT, OutputT> &&
detail::is_generic_group_v<Group>>
detail::is_generic_group_v<Group> &&
is_property_list_v<Properties>>
group_load(Group g, InputIteratorT in_ptr, sycl::vec<OutputT, N> &out,
Properties properties = {}) {
group_load(g, in_ptr, span<OutputT, N>(&out[0], N), properties);
Expand All @@ -382,7 +387,8 @@ group_load(Group g, InputIteratorT in_ptr, sycl::vec<OutputT, N> &out,
template <typename Group, typename InputT, int N, typename OutputIteratorT,
typename Properties = decltype(properties())>
std::enable_if_t<detail::verify_store_types<InputT, OutputIteratorT> &&
detail::is_generic_group_v<Group>>
detail::is_generic_group_v<Group> &&
is_property_list_v<Properties>>
group_store(Group g, const sycl::vec<InputT, N> &in, OutputIteratorT out_ptr,
Properties properties = {}) {
group_store(g, span<const InputT, N>(&in[0], N), out_ptr, properties);
Expand Down
Loading

0 comments on commit 15929c6

Please sign in to comment.