Skip to content

Commit

Permalink
Remove (3) and similar to it form default sorters
Browse files Browse the repository at this point in the history
  • Loading branch information
@andreyfe1
andreyfe1 committed Jan 30, 2024
1 parent 8a8ba3b commit 3b2b9ec
Showing 1 changed file with 36 additions and 79 deletions.
115 changes: 36 additions & 79 deletions sycl/doc/extensions/proposed/sycl_ext_oneapi_group_sort.asciidoc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -515,11 +515,7 @@ namespace sycl::ext::oneapi::experimental {
template<typename T>
static size_t
memory_required(sycl::memory_scope scope, std::size_t range_size); // (3)
template<typename T>
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<typename T,
Expand All @@ -528,21 +524,18 @@ namespace sycl::ext::oneapi::experimental {
class group_sorter{
public:
template<std::size_t Extent>
group_sorter(sycl::span<std::byte, Extent> scratch, Compare comp = {}); // (5)
group_sorter(sycl::span<std::byte, Extent> scratch, Compare comp = {}); // (4)
template<typename Group>
T operator()(Group g, T value); // (6)
T operator()(Group g, T value); // (5)
template<typename Group, typename Properties>
void operator()(Group g,
sycl::span<T, ElementsPerWorkItem> 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<typename T,
Expand All @@ -553,22 +546,19 @@ namespace sycl::ext::oneapi::experimental {
public:
template<std::size_t Extent>
group_key_value_sorter(sycl::span<std::byte, Extent> scratch,
Compare comp = {}); // (10)
Compare comp = {}); // (8)
template<typename Group>
std::tuple<T, U> operator()(Group g, T key, U value); // (11)
std::tuple<T, U> operator()(Group g, T key, U value); // (9)
template<typename Group, typename Properties>
void operator()(Group g,
sycl::span<T, ElementsPerWorkItem> keys,
sycl::span<U, ElementsPerWorkItem> 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)
};
}
Expand All @@ -583,13 +573,13 @@ namespace sycl::ext::oneapi::experimental {
template<std::size_t Extent>
joint_sorter(sycl::span<std::byte, Extent> scratch,
const std::bitset<sizeof(T) * CHAR_BIT> mask =
std::bitset<sizeof(T) * CHAR_BIT> (std::numeric_limits<unsigned long long>::max())); // (15)
std::bitset<sizeof(T) * CHAR_BIT> (std::numeric_limits<unsigned long long>::max())); // (12)
template<typename Group, typename Ptr>
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<typename T,
Expand All @@ -602,18 +592,18 @@ namespace sycl::ext::oneapi::experimental {
template<std::size_t Extent>
group_sorter(sycl::span<std::byte, Extent> scratch,
const std::bitset<sizeof(T) * CHAR_BIT> mask =
std::bitset<sizeof(T) * CHAR_BIT> (std::numeric_limits<unsigned long long>::max())); // (18)
std::bitset<sizeof(T) * CHAR_BIT> (std::numeric_limits<unsigned long long>::max())); // (15)
template<typename Group>
T operator()(Group g, T value); // (19)
T operator()(Group g, T value); // (16)
template<typename Group, typename Properties>
void operator()(Group g,
sycl::span<T, ElementsPerWorkItem> 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<typename T,
Expand All @@ -627,19 +617,19 @@ namespace sycl::ext::oneapi::experimental {
template<std::size_t Extent>
group_key_value_sorter(sycl::span<std::byte, Extent> scratch,
const std::bitset<sizeof(T) * CHAR_BIT> mask =
std::bitset<sizeof(T) * CHAR_BIT> (std::numeric_limits<unsigned long long>::max())); // (22)
std::bitset<sizeof(T) * CHAR_BIT> (std::numeric_limits<unsigned long long>::max())); // (19)
template<typename Group>
std::tuple<T, U> operator()(Group g, T key, U value); // (23)
std::tuple<T, U> operator()(Group g, T key, U value); // (20)
template<typename Group, typename Properties>
void operator()(Group g,
sycl::span<T, ElementsPerWorkItem> keys,
sycl::span<U, ElementsPerWorkItem> 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)
};
}
Expand Down Expand Up @@ -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
Expand All @@ -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,
Expand All @@ -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`
Expand All @@ -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`.
Expand All @@ -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,
Expand All @@ -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<T, ElementsPerWorkItem>`
that is an input parameter for `sort_over_group`, where `T` is
Expand All @@ -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<T, ElementsPerWorkItem>` and
`sycl::span<U, ElementsPerWorkItem>` as input parameters.
Expand Down

0 comments on commit 3b2b9ec

Please sign in to comment.