From 3b2b9ecf8c647e03590923d208f2a5758314b429 Mon Sep 17 00:00:00 2001 From: Andrei Fedorov Date: Mon, 29 Jan 2024 19:35:35 +0100 Subject: [PATCH] Remove (3) and similar to it form default sorters --- .../sycl_ext_oneapi_group_sort.asciidoc | 115 ++++++------------ 1 file changed, 36 insertions(+), 79 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_group_sort.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_group_sort.asciidoc index 196c91ad6a682..b1e305ec649cd 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_group_sort.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_group_sort.asciidoc @@ -515,11 +515,7 @@ namespace sycl::ext::oneapi::experimental { template static size_t - memory_required(sycl::memory_scope scope, std::size_t range_size); // (3) - - template - static size_t - memory_required(sycl::device d, sycl::memory_scope scope, std::size_t range_size); // (4) + memory_required(sycl::device d, sycl::memory_scope scope, std::size_t range_size); // (3) }; template - group_sorter(sycl::span scratch, Compare comp = {}); // (5) + group_sorter(sycl::span scratch, Compare comp = {}); // (4) template - T operator()(Group g, T value); // (6) + T operator()(Group g, T value); // (5) template void operator()(Group g, sycl::span values, - Properties properties); // (7) - - static size_t - memory_required(sycl::memory_scope scope, std::size_t range_size); // (8) + Properties properties); // (6) static size_t - memory_required(sycl::device d, sycl::memory_scope scope, std::size_t range_size); // (9) + memory_required(sycl::device d, sycl::memory_scope scope, std::size_t range_size); // (7) }; template group_key_value_sorter(sycl::span scratch, - Compare comp = {}); // (10) + Compare comp = {}); // (8) template - std::tuple operator()(Group g, T key, U value); // (11) + std::tuple operator()(Group g, T key, U value); // (9) template void operator()(Group g, sycl::span keys, sycl::span values, - Properties property); // (12) + Properties property); // (10) static std::size_t - memory_required(sycl::memory_scope scope, std::size_t range_size); // (13) - - static std::size_t - memory_required(sycl::device d, sycl::memory_scope scope, std::size_t range_size); // (14) + memory_required(sycl::device d, sycl::memory_scope scope, std::size_t range_size); // (11) }; } @@ -583,13 +573,13 @@ namespace sycl::ext::oneapi::experimental { template joint_sorter(sycl::span scratch, const std::bitset mask = - std::bitset (std::numeric_limits::max())); // (15) + std::bitset (std::numeric_limits::max())); // (12) template - void operator()(Group g, Ptr first, Ptr last); // (16) + void operator()(Group g, Ptr first, Ptr last); // (13) static constexpr std::size_t - memory_required(sycl::memory_scope scope, std::size_t range_size); // (17) + memory_required(sycl::memory_scope scope, std::size_t range_size); // (14) }; template group_sorter(sycl::span scratch, const std::bitset mask = - std::bitset (std::numeric_limits::max())); // (18) + std::bitset (std::numeric_limits::max())); // (15) template - T operator()(Group g, T value); // (19) + T operator()(Group g, T value); // (16) template void operator()(Group g, sycl::span values, - Properties properties); // (20) + Properties properties); // (17) static constexpr std::size_t - memory_required(sycl::memory_scope scope, std::size_t range_size); // (21) + memory_required(sycl::memory_scope scope, std::size_t range_size); // (18) }; template group_key_value_sorter(sycl::span scratch, const std::bitset mask = - std::bitset (std::numeric_limits::max())); // (22) + std::bitset (std::numeric_limits::max())); // (19) template - std::tuple operator()(Group g, T key, U value); // (23) + std::tuple operator()(Group g, T key, U value); // (20) template void operator()(Group g, sycl::span keys, sycl::span values, - Properties properties); // (24) + Properties properties); // (21) static constexpr std::size_t - memory_required(sycl::memory_scope scope, std::size_t range_size); // (25) + memory_required(sycl::memory_scope scope, std::size_t range_size); // (22) }; } @@ -669,7 +659,7 @@ callers side. Size of required memory (bytes) is defined by calling `memory_required`. |=== -(1), (5), (10) create the object using `comp`. +(1), (4), (8) create the object using `comp`. Additional memory for the algorithm is provided using `scratch`. If `scratch.size()` is less than the value returned by `memory_required`, behavior of the corresponding sorting algorithm @@ -681,18 +671,6 @@ the `joint_sort` algorithm. _Complexity_: Let `N` be `last - first`. `O(N*log(N)*log(N))` comparisons. (3) Returns size of temporary memory (in bytes) that is required by -the default sorting algorithm defined by the sorter calling by `joint_sort`. -`range_size` represents a range size for sorting, -e.g. `last-first` from `operator()` arguments. -It must be called within a SYCL kernel. -Result depends on the `scope` parameter: -use `sycl::memory_scope::work_group` to get memory size required -for each work-group; -use `sycl::memory_scope::sub_group` to get memory size required -for each sub-group. -If other `scope` values are passed, behavior is unspecified. - -(4) Returns size of temporary memory (in bytes) that is required by the default sorting algorithm defined by the sorter calling by `joint_sort` depending on `d`. `range_size` represents a range size for sorting, @@ -705,30 +683,19 @@ use `sycl::memory_scope::sub_group` to get memory size required for each sub-group. If other `scope` values are passed, behavior is unspecified. -(6) Implements a default sorting algorithm to be called by +(5) Implements a default sorting algorithm to be called by the `sort_over_group` algorithm. _Complexity_: Let `N` be the `Group` size. `O(N*log(N)*log(N))` comparisons. -(7) Implements a default sorting algorithm that is called by +(6) Implements a default sorting algorithm that is called by `sort_over_group` and that accepts the `sycl::span` value as an input parameter. _Complexity_: Let `N` be the `Group` size multiplied by `ElementsPerWorkItem`. `O(N*log(N)*log(N))` comparisons. -(8) Returns the size of temporary memory (in bytes) that is required by the default -sorting algorithm defined by the sorter calling by `sort_over_group`. -`ElementsPerWorkItem` is the extent parameter for `sycl::span` -that is an input parameter for `sort_over_group`. -It must be called within a SYCL kernel. -If `scope == sycl::memory_scope::work_group`, -`range_size` is the size of the local range for `sycl::nd_range` -that was used to run the kernel; -if `scope = sycl::memory_scope::sub_group`, `range_size` is a sub-group size. -If other `scope` values are passed, behavior is unspecified. - -(9) Returns the size of temporary memory (in bytes) that is required by the default +(7) Returns the size of temporary memory (in bytes) that is required by the default sorting algorithm defined by the sorter calling by `sort_over_group` depending on `d`. `ElementsPerWorkItem` is the extent parameter for `sycl::span` @@ -740,31 +707,21 @@ that was used to run the kernel; if `scope = sycl::memory_scope::sub_group`, `range_size` is a sub-group size. If other `scope` values are passed, behavior is unspecified. -(11) Implements a default key-value sorting algorithm that is called +(9) Implements a default key-value sorting algorithm that is called by `sort_key_value_over_group` and that doesn't accept `sycl::span` values as input parameters. _Complexity_: Let `N` be the `Group` size multiplied by `ElementsPerWorkItem`. `O(N*log(N)*log(N))` comparisons. -(12) Implements a default key-value sorting algorithm that is called +(10) Implements a default key-value sorting algorithm that is called by `sort_key_value_over_group` and that accepts `sycl::span` values as input parameters. _Complexity_: Let `N` be the `Group` size multiplied by `ElementsPerWorkItem`. `O(N*log(N)*log(N))` comparisons. -(13) Returns size of temporary memory (in bytes) that is required by -the default key-value -sorting algorithm defined by the sorter calling by `sort_key_value_over_group`. -It must be called within a SYCL kernel. -If `scope == sycl::memory_scope::work_group`, -`range_size` is the size of the local range for `sycl::nd_range` -that was used to run the kernel; -if `scope = sycl::memory_scope::sub_group`, `range_size` is a sub-group size. -If other `scope` values are passed, behavior is unspecified. - -(14) Returns size of temporary memory (in bytes) that is required by +(11) Returns size of temporary memory (in bytes) that is required by the default key-value sorting algorithm defined by the sorter calling by `sort_key_value_over_group` depending on `d`. @@ -775,17 +732,17 @@ that was used to run the kernel; if `scope = sycl::memory_scope::sub_group`, `range_size` is a sub-group size. If other `scope` values are passed, behavior is unspecified. -(15), (18), (22) create +(12), (15), (19) create the class object to sort values considering only bits that corresponds to 1 in `mask`. Additional memory for the algorithm is provided using `scratch`. If `scratch.size()` is less than the value returned by `memory_required`, behavior of the corresponding sorting algorithm is undefined. -(16) Implements the radix sorting algorithm to be called by +(13) Implements the radix sorting algorithm to be called by the `joint_sort` algorithm. -(17) Returns size of temporary memory (in bytes) that is required by +(14) Returns size of temporary memory (in bytes) that is required by the radix sort algorithm calling by `joint_sort`. `range_size` represents a range size for sorting, @@ -797,14 +754,14 @@ use `sycl::memory_scope::sub_group` to get memory size required for each sub-group. If other `scope` values are passed, behavior is unspecified. -(19) Implements the radix sorting algorithm to be called by +(16) Implements the radix sorting algorithm to be called by the `sort_over_group` algorithm. -(20) Implements the radix sorting algorithm that is called by +(17) Implements the radix sorting algorithm that is called by `sort_over_group` and that accepts the `sycl::span` value as an input parameter. -(21) Returns size of temporary memory (in bytes) that is required by the radix +(18) Returns size of temporary memory (in bytes) that is required by the radix sorting algorithm defined by the sorter calling by `sort_over_group`. `ElementsPerWorkItem` is a parameter for `sycl::span` that is an input parameter for `sort_over_group`, where `T` is @@ -815,15 +772,15 @@ that was used to run the kernel; if `scope = sycl::memory_scope::sub_group`, `range_size` is a sub-group size. If other `scope` values are passed, behavior is unspecified. -(23) Implements the radix sorting algorithm that is called +(20) Implements the radix sorting algorithm that is called by `sort_key_value_over_group` and that doesn't accept `sycl::span` values as input parameters. -(24) Implements the radix key-value sorting algorithm that is called +(21) Implements the radix key-value sorting algorithm that is called by `sort_key_value_over_group` and that accepts `sycl::span` values as input parameters. -(25) Returns size of temporary memory (in bytes) that is required by the radix key-value +(22) Returns size of temporary memory (in bytes) that is required by the radix key-value sorting algorithm defined by the sorter calling by `sort_key_value_over_group` with `sycl::span` and `sycl::span` as input parameters.