Skip to content

[SYCL] Properties must be a properties list in group_load/store #16422

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
aelovikov-intel merged 1 commit into from
Dec 23, 2024
Merged

[SYCL] Properties must be a properties list in group_load/store #16422

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
30 changes: 21 additions & 9 deletions sycl/doc/extensions/experimental/sycl_ext_oneapi_group_load_store.asciidoc
View file
Open in desktop
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
Open in desktop
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
Loading