From f5439406dc7e3f31720673e73957093fd9392fd9 Mon Sep 17 00:00:00 2001 From: "Fedorov, Andrey" Date: Thu, 8 Apr 2021 19:41:55 +0300 Subject: [PATCH 1/7] added initial version for group sort Signed-off-by: Fedorov, Andrey --- .../SYCL_INTEL_group_sort.asciidoc | 328 ++++++++++++++++++ sycl/doc/extensions/README.md | 1 + 2 files changed, 329 insertions(+) create mode 100755 sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc diff --git a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc new file mode 100755 index 0000000000000..9d7c249523e85 --- /dev/null +++ b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc @@ -0,0 +1,328 @@ += SYCL_EXT_ONEAPI_GROUP_SORT +:source-highlighter: coderay +:coderay-linenums-mode: table + +// This section needs to be after the document title. +:doctype: book +:toc2: +:toc: left +:encoding: utf-8 +:lang: en + +:blank: pass:[ +] + +// Set the default source code type in this document to C++, +// for syntax highlighting purposes. This is needed because +// docbook uses c++ and html5 uses cpp. +:language: {basebackend@docbook:c++:cpp} + +== Notice + +Copyright (c) 2021-2021 Intel Corporation. All rights reserved. + +IMPORTANT: This specification is a draft. + +NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are +trademarks of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. +used by permission by Khronos. + +NOTE: This document is better viewed when rendered as html with asciidoctor. +GitHub does not render image icons. + +This extension is written against the SYCL 2020 revision 3 specification. All +references below to the "core SYCL specification" or to section numbers in the +SYCL specification refer to that revision. + +== Introduction + +This extension introduces sorting functions to the group algorithms library, +Sorter objects and Sorter type traits. + +== Feature test macro + +This extension provides a feature-test macro as described in the core SYCL +specification section 6.3.3 "Feature test macros". Therefore, an +implementation supporting this extension must predefine the macro +`SYCL_EXT_ONEAPI_GROUP_SORT` to one of the values defined in the table below. +Applications can test for the existence of this macro to determine if the +implementation supports this feature, or applications can test the macro's +value to determine which of the extension's APIs the implementation supports. + +[%header,cols="1,5"] +|=== +|Value |Description +|1 |Initial extension version. Base features are supported. +|=== + +==== Sorter + +Sorter is a special type that encapsulates a sorting algorithm. +Sorter may contain parameters that help to get better performance. +Those sort data provided to the `operator()` +that should contain an implementation of a sorting algorithm. +Sorters must have a template +argument representing binary comparison function object and must define +`compare_type` as an alias for such comparison object. +Semantics of `operator()` is following: + +[source,c++] +---- +// to call by joint_sort +template +void operator()(Group g, Ptr begin, Ptr end); + +// and/or to call by sort_over_group +template +T operator()(Group g, T val); + +---- +_Constraints_: Available only if `sycl::is_group_v>` is true. + +_Preconditions_: `first`, `last` must be the same for all work-items in the group. + +Example of custom Sorter: +[source,c++] +---- +template +class bubble_sort{ +public: + using compare_type = Compare; + compare_type comp; + + template + void operator()(Group g, Ptr begin, Ptr end){ + size_t n = end - begin; + size_t idx = g.get_local_id().get(0); + if(idx == 0) + for(size_t i = 0; i < n; ++i) + for(size_t j = i + 1; j < n; ++j) + if(comp(begin[j], begin[i])) + std::swap(begin[i], begin[j]); + } +}; +---- + +==== Predefined Sorters +SYCL provides the following predefined classes: + +[source,c++] +---- +namespace sycl::ext::oneapi { + + template> + class default_sorter { + public: + using compare_type = Compare; + + default_sorter(Compare comp = Compare()); + + template + void operator()(Group g, Ptr begin, Ptr end); + + template + T operator()(Group g, T val); + }; + + template, unsigned int BitsPerPass = 4> + class radix_sorter { + public: + using compare_type = Compare; + + radix_sorter(const std::bitset mask = + std::bitset (std::numeric_limits::max())); + + template + void operator()(Group g, Ptr begin, Ptr end); + + template + T operator()(Group g, T val); + }; + +} +---- + +Table. Description of predefined Sorters. +|=== +|Sorter|Description + +|`template> +default_sorter` +|Use a default sorting method based on an implementation-defined heuristic +using `Compare` as the binary comparison function object. + +|`template, unsigned int BitsPerPass = 4> +radix_sorter` +|Use radix sort as a sorting method. `Compare` is used as the binary comparison +function object. The method can provide better runtime performance. +Only arithmetic types as `T` and `std::less`, `std::greater` as `Compare` +can be passed to radix_sorter. +`BitsPerPass` is a number of bits that values are split by. +For example, if a sequence of `int32_t` is sorted using `BitsPerPass == 4` then one +pass of the radix sort algorithm considers only 4 bits. The number of passes is `32/4=8`. +|=== + +Table. Constructors of the `default_sorter` class. +|=== +|Constructor|Description + +|`default_sorter(Compare comp = Compare())` +|Creates the `default_sorter` object using `comp`. +|=== + +Table. Member functions of the `default_sorter` class. +|=== +|Member function|Description + +|`template +void operator()(Group g, Ptr begin, Ptr end)` +|Implements a default sorting algorithm. It's callable by the `joint_sort` algorithm. + +_Complexity_: Let `N` be `last - first`. `O(Nlog^2(N))` comparisons. + +|`template +T operator()(Group g, T val)` +|Implements a default sorting algorithm. It's callable by the `sort_over_group` algorithm. + +_Complexity_: Let `N` be the work group size. `O(Nlog^2(N))` comparisons. +|=== + +Table. Constructors of the `radix_sorter` class. +|=== +|Constructor|Description + +|`radix_sorter(const std::bitset mask = std::bitset +(std::numeric_limits::max()));` +|Creates the `radix_sorter` object to sort values considering only bits +that corresponds to 1 in `mask`. +|=== + +Table. Member functions of the `radix_sorter` class. +|=== +|Member function|Description + +|`template +void operator()(Group g, Ptr begin, Ptr end)` +|Implements the radix sort algorithm. It's callable by the `joint_sort` algorithm. + +|`template +T operator()(Group g, T val)` +|Implements the radix sort algorithm. It's callable by the `sort_over_group` algorithm. +|=== + +==== Sorter type trait + +[source,c++] +---- +namespace sycl::ext::oneapi { + template + class is_sorter; + + template + inline constexpr bool is_sorter_v = is_sorter::value; +} +---- + +The `is_sorter` type trait is used to determine which Sorters are supported +by `joint_sort` and `sort_over_group` functions, and to control when sorting +functions participate in overload resolution. + +`is_sorter` is `std::true_type` if `T` contains a member type named `compare_type` and +`std::­false_type` otherwise. A SYCL implementation may introduce additional +specializations of `is_sorter` for implementation-defined Sorters. If an application +defines a customer Sorter `S`, it should ensure that `S` contains a member type named +`compare_type` and the `operator()` member function or provide a specialization `is_sorter`. + +==== Sort +The sort function from the {cpp} standard sorts elements with respect to +the binary comparison function object. + +SYCL provides two similar algorithms: + +`joint_sort` uses the work-items in a group to execute the corresponding +algorithm in parallel. + +`sort_over_group` performs a sort over values held directly by the work-items +in a group, and results returned to work-item `i` represent values that are in +position `i` in the ordered range. + +[source,c++] +---- +namespace sycl::ext::oneapi { + template + void joint_sort(Group g, Ptr first, Ptr last); // (1) + + template + void joint_sort(Group g, Ptr first, Ptr last, Compare comp); // (2) + + template + void joint_sort(Group g, Ptr first, Ptr last, Sorter sorter); // (3) + + template + T sort_over_group(Group g, T val); // (4) + + template + T sort_over_group(Group g, T val, Compare comp); // (5) + + template + T sort_over_group(Group g, T val, Sorter sorter); // (6) +} +---- + +_Constraints_: All functions are available only if `sycl::is_group_v>` +is true and `sycl::ext::oneapi::is_sorter_v>` is true. + +_Preconditions_: `first`, `last` must be the same for all work-items in the group. + +1._Effects_: Sort the elements in the range `[first, last_)`. +Elements are compared by the `operator<`. + +_Complexity_: Let `N` be `last - first`. `O(Nlog^2(N))` comparisons. + +2._Mandates_: `comp` must satisfy the requirements of `Compare` from +the {cpp} standard. + +_Effects_: Sort the elements in the range `[first, last)` with respect to the +binary comparison function object `comp`. + +_Complexity_: Let `N` be `last - first`. `O(Nlog^2(N))` comparisons. + +3._Effects_: Equivalent to: `sorter(g, first, last)`. + +4._Returns_: The value returned on work-item `i` is the value in position `i` +of the ordered range resulting from sorting `val` from all work-items in the +`g` group. Elements are compared by the `operator<`. +For multi-dimensional groups, the order of work-items in the group is +determined by their linear id. + +_Complexity_: Let `N` be the work group size. `O(Nlog^2(N))`` comparisons. + +2._Mandates_: `comp` must satisfy the requirements of `Compare` from +the {cpp} standard. + +_Returns_: The value returned on work-item `i` is the value in position `i` +of the ordered range resulting from sorting `val` from all work-items in the +`g` group with respect to the binary comparison function object `comp`. +For multi-dimensional groups, the order of work-items in the group is +determined by their linear id. + +_Complexity_: Let `N` be the work group size. `O(Nlog^2(N))`` comparisons. + +6._Effects_: Equivalent to: `return sorter(g, val)`. + +== Issues + +. Sort function can have interfaces with static arrays in private memory +as well. The concern is that it needs to check the performance gain of such +interfaces more closely. +. It can be a separate proposal for key-value sorting basing on Projections. +It needs to be investigated what is the response for that. + +== Revision History + +[cols="5,15,15,70"] +[grid="rows"] +[options="header"] +|======================================== +|Rev|Date|Author|Changes +|1|{docdate}|Andrey Fedorov|Initial public working draft +|======================================== diff --git a/sycl/doc/extensions/README.md b/sycl/doc/extensions/README.md index 0faddf8955ee2..de45b10137aa3 100755 --- a/sycl/doc/extensions/README.md +++ b/sycl/doc/extensions/README.md @@ -38,6 +38,7 @@ DPC++ extensions status: | [Level-Zero backend specification](LevelZeroBackend/LevelZeroBackend.md) | Supported | | | [ITT annotations support](ITTAnnotations/ITTAnnotations.rst) | Supported | | | [SYCL_EXT_ONEAPI_DEVICE_IF](DeviceIf/device_if.asciidoc) | Proposal | | +| [SYCL_INTEL_group_sort](GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc) | Proposal | | Legend: From d0d9eadac2ead6258d775621dd7d0a8812afbcb4 Mon Sep 17 00:00:00 2001 From: "Fedorov, Andrey" Date: Tue, 20 Apr 2021 17:58:01 +0300 Subject: [PATCH 2/7] update radix_sorter and remove is_sorter Signed-off-by: Fedorov, Andrey --- .../SYCL_INTEL_group_sort.asciidoc | 88 +++++++++---------- 1 file changed, 41 insertions(+), 47 deletions(-) diff --git a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc index 9d7c249523e85..802161276ca62 100755 --- a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc +++ b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc @@ -56,29 +56,37 @@ value to determine which of the extension's APIs the implementation supports. ==== Sorter -Sorter is a special type that encapsulates a sorting algorithm. -Sorter may contain parameters that help to get better performance. -Those sort data provided to the `operator()` +Sorter is a special type that encapsulates a sorting algorithm. Sorter may contain parameters +that help to get better performance. Data for sorting are provided to the `operator()` that should contain an implementation of a sorting algorithm. -Sorters must have a template -argument representing binary comparison function object and must define -`compare_type` as an alias for such comparison object. Semantics of `operator()` is following: [source,c++] ---- -// to call by joint_sort template void operator()(Group g, Ptr begin, Ptr end); -// and/or to call by sort_over_group template T operator()(Group g, T val); - ---- -_Constraints_: Available only if `sycl::is_group_v>` is true. -_Preconditions_: `first`, `last` must be the same for all work-items in the group. +At least one overload for `operator()` is required. + +Table. `operator()` for Sorters. +|=== +|`operator()`|Description + +|`template +void operator()(Group g, Ptr begin, Ptr end);` +|Implements a sorting algorithm that calls by `joint_sort`. +Available only if `sycl::is_group_v>` is true. +`begin`, `end` must be the same for all work-items in the group. + +|`template +T operator()(Group g, T val);` +|Implements a sorting algorithm that calls by `sort_over_group`. +Available only if `sycl::is_group_v>` is true. +|=== Example of custom Sorter: [source,c++] @@ -86,8 +94,7 @@ Example of custom Sorter: template class bubble_sort{ public: - using compare_type = Compare; - compare_type comp; + Compare comp; template void operator()(Group g, Ptr begin, Ptr end){ @@ -103,6 +110,22 @@ public: ---- ==== Predefined Sorters + +`radix_order` is a `enum` that defines the sorting order when `radix_sorter` is used. +Only ascending and descending orders are applicable. + +[source,c++] +---- +namespace sycl::ext::oneapi { + + enum radix_order { + ascending, + descending + }; + +} +---- + SYCL provides the following predefined classes: [source,c++] @@ -112,8 +135,6 @@ namespace sycl::ext::oneapi { template> class default_sorter { public: - using compare_type = Compare; - default_sorter(Compare comp = Compare()); template @@ -123,11 +144,9 @@ namespace sycl::ext::oneapi { T operator()(Group g, T val); }; - template, unsigned int BitsPerPass = 4> + template class radix_sorter { public: - using compare_type = Compare; - radix_sorter(const std::bitset mask = std::bitset (std::numeric_limits::max())); @@ -152,10 +171,8 @@ using `Compare` as the binary comparison function object. |`template, unsigned int BitsPerPass = 4> radix_sorter` -|Use radix sort as a sorting method. `Compare` is used as the binary comparison -function object. The method can provide better runtime performance. -Only arithmetic types as `T` and `std::less`, `std::greater` as `Compare` -can be passed to radix_sorter. +|Use radix sort as a sorting method. `Order` specify the sorting order. +Only arithmetic types as `T` can be passed to radix_sorter. `BitsPerPass` is a number of bits that values are split by. For example, if a sequence of `int32_t` is sorted using `BitsPerPass == 4` then one pass of the radix sort algorithm considers only 4 bits. The number of passes is `32/4=8`. @@ -177,7 +194,7 @@ Table. Member functions of the `default_sorter` class. void operator()(Group g, Ptr begin, Ptr end)` |Implements a default sorting algorithm. It's callable by the `joint_sort` algorithm. -_Complexity_: Let `N` be `last - first`. `O(Nlog^2(N))` comparisons. +_Complexity_: Let `N` be `end - begin`. `O(Nlog^2(N))` comparisons. |`template T operator()(Group g, T val)` @@ -209,29 +226,6 @@ T operator()(Group g, T val)` |Implements the radix sort algorithm. It's callable by the `sort_over_group` algorithm. |=== -==== Sorter type trait - -[source,c++] ----- -namespace sycl::ext::oneapi { - template - class is_sorter; - - template - inline constexpr bool is_sorter_v = is_sorter::value; -} ----- - -The `is_sorter` type trait is used to determine which Sorters are supported -by `joint_sort` and `sort_over_group` functions, and to control when sorting -functions participate in overload resolution. - -`is_sorter` is `std::true_type` if `T` contains a member type named `compare_type` and -`std::­false_type` otherwise. A SYCL implementation may introduce additional -specializations of `is_sorter` for implementation-defined Sorters. If an application -defines a customer Sorter `S`, it should ensure that `S` contains a member type named -`compare_type` and the `operator()` member function or provide a specialization `is_sorter`. - ==== Sort The sort function from the {cpp} standard sorts elements with respect to the binary comparison function object. @@ -296,7 +290,7 @@ determined by their linear id. _Complexity_: Let `N` be the work group size. `O(Nlog^2(N))`` comparisons. -2._Mandates_: `comp` must satisfy the requirements of `Compare` from +5._Mandates_: `comp` must satisfy the requirements of `Compare` from the {cpp} standard. _Returns_: The value returned on work-item `i` is the value in position `i` From 85cc411ca9ce193583d92ec46c3703f7de59f896 Mon Sep 17 00:00:00 2001 From: "Fedorov, Andrey" Date: Tue, 20 Apr 2021 18:12:36 +0300 Subject: [PATCH 3/7] reworked issues Signed-off-by: Fedorov, Andrey --- .../GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc index 802161276ca62..145478f25fcd3 100755 --- a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc +++ b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc @@ -305,11 +305,14 @@ _Complexity_: Let `N` be the work group size. `O(Nlog^2(N))`` comparisons. == Issues -. Sort function can have interfaces with static arrays in private memory -as well. The concern is that it needs to check the performance gain of such -interfaces more closely. +. Sort function can have interfaces with static arrays in private memory as well. +The concern is that it can require changes for other group algortihms as well since sort +basing on private memory is not very useful if other algorithms in the chain use local +memory only. . It can be a separate proposal for key-value sorting basing on Projections. It needs to be investigated what is the response for that. +. Sorter trait can be useful if there are Finder, Reducer or other objects +will be added to the Spec to be used with other Group algorithms, e.g. find, reduce. == Revision History From 17efe737cfb249cd87f60e2b9c0f65de6fbabebb Mon Sep 17 00:00:00 2001 From: "Fedorov, Andrey" Date: Wed, 21 Apr 2021 18:06:55 +0300 Subject: [PATCH 4/7] apply comments Signed-off-by: Fedorov, Andrey --- .../SYCL_INTEL_group_sort.asciidoc | 26 +++++++++---------- 1 file changed, 13 insertions(+), 13 deletions(-) diff --git a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc index 145478f25fcd3..bf5af2c6e839a 100755 --- a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc +++ b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc @@ -8,6 +8,7 @@ :toc: left :encoding: utf-8 :lang: en +:stem: :blank: pass:[ +] @@ -35,8 +36,7 @@ SYCL specification refer to that revision. == Introduction -This extension introduces sorting functions to the group algorithms library, -Sorter objects and Sorter type traits. +This extension introduces sorting functions to the group algorithms library and Sorter objects. == Feature test macro @@ -118,7 +118,7 @@ Only ascending and descending orders are applicable. ---- namespace sycl::ext::oneapi { - enum radix_order { + enum class radix_order { ascending, descending }; @@ -153,7 +153,7 @@ namespace sycl::ext::oneapi { template void operator()(Group g, Ptr begin, Ptr end); - template + template T operator()(Group g, T val); }; @@ -169,7 +169,7 @@ default_sorter` |Use a default sorting method based on an implementation-defined heuristic using `Compare` as the binary comparison function object. -|`template, unsigned int BitsPerPass = 4> +|`template radix_sorter` |Use radix sort as a sorting method. `Order` specify the sorting order. Only arithmetic types as `T` can be passed to radix_sorter. @@ -194,13 +194,13 @@ Table. Member functions of the `default_sorter` class. void operator()(Group g, Ptr begin, Ptr end)` |Implements a default sorting algorithm. It's callable by the `joint_sort` algorithm. -_Complexity_: Let `N` be `end - begin`. `O(Nlog^2(N))` comparisons. +_Complexity_: Let `N` be `end - begin`:stem:[O(N \times log_2(N))] comparisons. |`template T operator()(Group g, T val)` |Implements a default sorting algorithm. It's callable by the `sort_over_group` algorithm. -_Complexity_: Let `N` be the work group size. `O(Nlog^2(N))` comparisons. +_Complexity_: Let `N` be the work group size:stem:[O(N \times log_2(N))] comparisons. |=== Table. Constructors of the `radix_sorter` class. @@ -263,14 +263,14 @@ namespace sycl::ext::oneapi { ---- _Constraints_: All functions are available only if `sycl::is_group_v>` -is true and `sycl::ext::oneapi::is_sorter_v>` is true. +is true and `Sorter` is a SYCL Sorter. _Preconditions_: `first`, `last` must be the same for all work-items in the group. -1._Effects_: Sort the elements in the range `[first, last_)`. +1._Effects_: Sort the elements in the range `[first, last)`. Elements are compared by the `operator<`. -_Complexity_: Let `N` be `last - first`. `O(Nlog^2(N))` comparisons. +_Complexity_: Let `N` be `last - first`:stem:[O(N \times log_2(N))] comparisons. 2._Mandates_: `comp` must satisfy the requirements of `Compare` from the {cpp} standard. @@ -278,7 +278,7 @@ the {cpp} standard. _Effects_: Sort the elements in the range `[first, last)` with respect to the binary comparison function object `comp`. -_Complexity_: Let `N` be `last - first`. `O(Nlog^2(N))` comparisons. +_Complexity_: Let `N` be `last - first`:stem:[O(N \times log_2(N))] comparisons. 3._Effects_: Equivalent to: `sorter(g, first, last)`. @@ -288,7 +288,7 @@ of the ordered range resulting from sorting `val` from all work-items in the For multi-dimensional groups, the order of work-items in the group is determined by their linear id. -_Complexity_: Let `N` be the work group size. `O(Nlog^2(N))`` comparisons. +_Complexity_: Let `N` be the work group size:stem:[O(N \times log_2(N))] comparisons. 5._Mandates_: `comp` must satisfy the requirements of `Compare` from the {cpp} standard. @@ -299,7 +299,7 @@ of the ordered range resulting from sorting `val` from all work-items in the For multi-dimensional groups, the order of work-items in the group is determined by their linear id. -_Complexity_: Let `N` be the work group size. `O(Nlog^2(N))`` comparisons. +_Complexity_: Let `N` be the work group size:stem:[O(N \times log_2(N))] comparisons. 6._Effects_: Equivalent to: `return sorter(g, val)`. From 7c75ab3bbd9bc9c959e182e50197d100b23283b9 Mon Sep 17 00:00:00 2001 From: "Fedorov, Andrey" Date: Fri, 23 Apr 2021 10:36:54 +0300 Subject: [PATCH 5/7] add O(N*log_2(N)) Signed-off-by: Fedorov, Andrey --- .../SYCL_INTEL_group_sort.asciidoc | 15 +++++++-------- 1 file changed, 7 insertions(+), 8 deletions(-) diff --git a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc index bf5af2c6e839a..4a4be38c6bae9 100755 --- a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc +++ b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc @@ -8,7 +8,6 @@ :toc: left :encoding: utf-8 :lang: en -:stem: :blank: pass:[ +] @@ -194,13 +193,13 @@ Table. Member functions of the `default_sorter` class. void operator()(Group g, Ptr begin, Ptr end)` |Implements a default sorting algorithm. It's callable by the `joint_sort` algorithm. -_Complexity_: Let `N` be `end - begin`:stem:[O(N \times log_2(N))] comparisons. +_Complexity_: Let `N` be `end - begin`. `O(N*log_2(N))` comparisons. |`template T operator()(Group g, T val)` |Implements a default sorting algorithm. It's callable by the `sort_over_group` algorithm. -_Complexity_: Let `N` be the work group size:stem:[O(N \times log_2(N))] comparisons. +_Complexity_: Let `N` be the work group size. `O(N*log_2(N))` comparisons. |=== Table. Constructors of the `radix_sorter` class. @@ -270,7 +269,7 @@ _Preconditions_: `first`, `last` must be the same for all work-items in the grou 1._Effects_: Sort the elements in the range `[first, last)`. Elements are compared by the `operator<`. -_Complexity_: Let `N` be `last - first`:stem:[O(N \times log_2(N))] comparisons. +_Complexity_: Let `N` be `last - first`. `O(N*log_2(N))` comparisons. 2._Mandates_: `comp` must satisfy the requirements of `Compare` from the {cpp} standard. @@ -278,7 +277,7 @@ the {cpp} standard. _Effects_: Sort the elements in the range `[first, last)` with respect to the binary comparison function object `comp`. -_Complexity_: Let `N` be `last - first`:stem:[O(N \times log_2(N))] comparisons. +_Complexity_: Let `N` be `last - first`. `O(N*log_2(N))` comparisons. 3._Effects_: Equivalent to: `sorter(g, first, last)`. @@ -288,7 +287,7 @@ of the ordered range resulting from sorting `val` from all work-items in the For multi-dimensional groups, the order of work-items in the group is determined by their linear id. -_Complexity_: Let `N` be the work group size:stem:[O(N \times log_2(N))] comparisons. +_Complexity_: Let `N` be the work group size. `O(N*log_2(N))` comparisons. 5._Mandates_: `comp` must satisfy the requirements of `Compare` from the {cpp} standard. @@ -299,7 +298,7 @@ of the ordered range resulting from sorting `val` from all work-items in the For multi-dimensional groups, the order of work-items in the group is determined by their linear id. -_Complexity_: Let `N` be the work group size:stem:[O(N \times log_2(N))] comparisons. +_Complexity_: Let `N` be the work group size. `O(N*log_2(N))` comparisons. 6._Effects_: Equivalent to: `return sorter(g, val)`. @@ -311,7 +310,7 @@ basing on private memory is not very useful if other algorithms in the chain use memory only. . It can be a separate proposal for key-value sorting basing on Projections. It needs to be investigated what is the response for that. -. Sorter trait can be useful if there are Finder, Reducer or other objects +. Sorter traits can be useful if there are Finder, Reducer or other objects will be added to the Spec to be used with other Group algorithms, e.g. find, reduce. == Revision History From 0869b0519b6e111ab34391fd4037a6cc8d4e5ebf Mon Sep 17 00:00:00 2001 From: "Fedorov, Andrey" Date: Fri, 23 Apr 2021 16:45:15 +0300 Subject: [PATCH 6/7] fixed a wording --- .../SYCL_INTEL_group_sort.asciidoc | 40 +++++++++---------- 1 file changed, 20 insertions(+), 20 deletions(-) diff --git a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc index 4a4be38c6bae9..8cbcc8b17a3d4 100755 --- a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc +++ b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc @@ -63,7 +63,7 @@ Semantics of `operator()` is following: [source,c++] ---- template -void operator()(Group g, Ptr begin, Ptr end); +void operator()(Group g, Ptr first, Ptr last); template T operator()(Group g, T val); @@ -76,10 +76,10 @@ Table. `operator()` for Sorters. |`operator()`|Description |`template -void operator()(Group g, Ptr begin, Ptr end);` +void operator()(Group g, Ptr first, Ptr last);` |Implements a sorting algorithm that calls by `joint_sort`. Available only if `sycl::is_group_v>` is true. -`begin`, `end` must be the same for all work-items in the group. +`first`, `last` must be the same for all work-items in the group. |`template T operator()(Group g, T val);` @@ -96,14 +96,14 @@ public: Compare comp; template - void operator()(Group g, Ptr begin, Ptr end){ - size_t n = end - begin; + void operator()(Group g, Ptr first, Ptr last){ + size_t n = last - first; size_t idx = g.get_local_id().get(0); if(idx == 0) for(size_t i = 0; i < n; ++i) for(size_t j = i + 1; j < n; ++j) - if(comp(begin[j], begin[i])) - std::swap(begin[i], begin[j]); + if(comp(first[j], first[i])) + std::swap(first[i], first[j]); } }; ---- @@ -137,7 +137,7 @@ namespace sycl::ext::oneapi { default_sorter(Compare comp = Compare()); template - void operator()(Group g, Ptr begin, Ptr end); + void operator()(Group g, Ptr first, Ptr last); template T operator()(Group g, T val); @@ -150,7 +150,7 @@ namespace sycl::ext::oneapi { std::bitset (std::numeric_limits::max())); template - void operator()(Group g, Ptr begin, Ptr end); + void operator()(Group g, Ptr first, Ptr last); template T operator()(Group g, T val); @@ -171,7 +171,7 @@ using `Compare` as the binary comparison function object. |`template radix_sorter` |Use radix sort as a sorting method. `Order` specify the sorting order. -Only arithmetic types as `T` can be passed to radix_sorter. +Only arithmetic types as `T` can be passed to `radix_sorter`. `BitsPerPass` is a number of bits that values are split by. For example, if a sequence of `int32_t` is sorted using `BitsPerPass == 4` then one pass of the radix sort algorithm considers only 4 bits. The number of passes is `32/4=8`. @@ -190,14 +190,14 @@ Table. Member functions of the `default_sorter` class. |Member function|Description |`template -void operator()(Group g, Ptr begin, Ptr end)` -|Implements a default sorting algorithm. It's callable by the `joint_sort` algorithm. +void operator()(Group g, Ptr first, Ptr last)` +|Implements a default sorting algorithm to be called by the `joint_sort` algorithm. -_Complexity_: Let `N` be `end - begin`. `O(N*log_2(N))` comparisons. +_Complexity_: Let `N` be `last - first`. `O(N*log_2(N))` comparisons. |`template T operator()(Group g, T val)` -|Implements a default sorting algorithm. It's callable by the `sort_over_group` algorithm. +|Implements a default sorting algorithm to be called by the `sort_over_group` algorithm. _Complexity_: Let `N` be the work group size. `O(N*log_2(N))` comparisons. |=== @@ -217,12 +217,12 @@ Table. Member functions of the `radix_sorter` class. |Member function|Description |`template -void operator()(Group g, Ptr begin, Ptr end)` -|Implements the radix sort algorithm. It's callable by the `joint_sort` algorithm. +void operator()(Group g, Ptr first, Ptr last)` +|Implements the radix sort algorithm to be called by the `joint_sort` algorithm. -|`template +|`template T operator()(Group g, T val)` -|Implements the radix sort algorithm. It's callable by the `sort_over_group` algorithm. +|Implements the radix sort algorithm to be called by the `sort_over_group` algorithm. |=== ==== Sort @@ -267,7 +267,7 @@ is true and `Sorter` is a SYCL Sorter. _Preconditions_: `first`, `last` must be the same for all work-items in the group. 1._Effects_: Sort the elements in the range `[first, last)`. -Elements are compared by the `operator<`. +Elements are compared by `operator<`. _Complexity_: Let `N` be `last - first`. `O(N*log_2(N))` comparisons. @@ -283,7 +283,7 @@ _Complexity_: Let `N` be `last - first`. `O(N*log_2(N))` comparisons. 4._Returns_: The value returned on work-item `i` is the value in position `i` of the ordered range resulting from sorting `val` from all work-items in the -`g` group. Elements are compared by the `operator<`. +`g` group. Elements are compared by `operator<`. For multi-dimensional groups, the order of work-items in the group is determined by their linear id. From 1420c39a3aeda593750bbe7b461cd5decafacfd8 Mon Sep 17 00:00:00 2001 From: "Fedorov, Andrey" Date: Mon, 26 Apr 2021 18:42:38 +0300 Subject: [PATCH 7/7] update copyright Signed-off-by: Fedorov, Andrey --- .../extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc index 8cbcc8b17a3d4..8f342878e155d 100755 --- a/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc +++ b/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc @@ -18,7 +18,7 @@ == Notice -Copyright (c) 2021-2021 Intel Corporation. All rights reserved. +Copyright (c) 2021 Intel Corporation. All rights reserved. IMPORTANT: This specification is a draft.