-
Notifications
You must be signed in to change notification settings - Fork 770
[SYCL] Add support for key/value sorting APIs #13942
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.
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
Conversation
range_size, comp, | ||
aligned_scratch_ptr + range_size * sizeof(T)); | ||
val = *local_copy; | ||
} |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Should we add something like assert(false && "Scratch size less than value results in UB")
in the else part? It will be helpful in debugging.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Unfortunately native assert is not supported by all devices and I believe fallback assert can add overhead (as it creates assert buffer and copies it back) and I'm not sure if it is reliable enough to use it in default execution path. That's why I didn't put assert here.
Tag @andreyfe1 |
: comp(comp_), scratch(scratch_) {} | ||
|
||
template <typename Group> | ||
T operator()([[maybe_unused]] Group g, [[maybe_unused]] T val) { |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Is it planned to add another operator()
: void operator()(Group g, sycl::span<T, ElementsPerWorkItem> values, Properties properties);
?
The same for group_key_value_sorter
and radix_sorters
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I didn't add them on purpose to split functionality into several PRs. Do you think it is better to create a single PR with all functionality?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
No, I needed to make sure that it's not forgotten. I'm fine if it's covered in a separate PR
a6187e6
to
6aa7995
Compare
@uditagarwal97 @aelovikov-intel This PR was rebased on top of the latest changes and is ready for review. |
auto this_comp = this->comp; | ||
auto comp_key_value = [this_comp](const KeyValue &lhs, |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
auto this_comp = this->comp; | |
auto comp_key_value = [this_comp](const KeyValue &lhs, | |
auto comp_key_value = [this_comp = this->comp](const KeyValue &lhs, |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Fixed.
@@ -343,6 +343,41 @@ class group_sorter { | |||
} | |||
}; | |||
|
|||
template <typename T, typename U, typename CompareT = std::less<>, |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Please rename T
/U
to KeyTy
/ValueTy
.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Renamed.
typename std::is_same<std::tuple<Key, Value>, | ||
decltype(std::declval<Sorter>()( | ||
std::declval<G>(), std::declval<Key>(), | ||
std::declval<Value>()))>; |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I think std::invoke_result_t
would allow to simplify this a bit.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Thanks, simplified.
g, KeyValue(key, value)); | ||
} | ||
|
||
static constexpr std::size_t memory_required(sycl::memory_scope scope, |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
According to the spec memory_required
is not constexpr
for default_sorters
. Please, make corresponding changes for all default sorters
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Thanks, fixed.
@@ -343,6 +343,41 @@ class group_sorter { | |||
} |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
This comment is regarding https://github.com/intel/llvm/blob/sycl/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp#L252
It's expected for joint_sorter
(default and radix sorters both) to support oneDPL zip_iterator
(https://github.com/intel/llvm/blob/sycl/sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp#L252 : Example 3).
Do you plan to add this support (no matter in this PR or a separate one)?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I will look into this but not as part of this PR.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
ok. Thanks
Add group_key_value_sorter sorters and sort_key_value_over_group APIs based on https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_group_sort.asciidoc extension. This PR was split out from larger PR: intel#13713 Co-authored-by: "Andrei Fedorov [[email protected]](mailto:[email protected])" Co-authored-by: "Romanov Vlad [[email protected]](mailto:[email protected])"
Add group_key_value_sorter sorters and sort_key_value_over_group APIs based on https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_group_sort.asciidoc extension.
This PR was split out from larger PR: #13713
Co-authored-by: "Andrei Fedorov [email protected]"
Co-authored-by: "Romanov Vlad [email protected]"