Out of Place Esimd Radix Sort and Tests #1439

danhoeflinger · 2024-03-11T13:10:03Z

Adding out of place support to ESIMD radix sort [by key] (high level refactoring).
Adding tests to test out of place sort [by key].
Removing some unnecessary extra kernel names (Odd vs Even should be the same kernel, just with different runtime params)

APIs are disambiguated between iterator and range interfaces using enable_if to check if a type is an iterator. Documentation updates will be made in a separate PR.

Signed-off-by: Dan Hoeflinger <[email protected]>

SergeyKopienko · 2024-03-26T08:34:42Z

include/oneapi/dpl/experimental/kt/esimd_radix_sort.h

+radix_sort_out_of_place(sycl::queue __q, _KeysRng1&& __keys_rng, _KeysRng2&& __keys_rng_out, _KernelParam __param = {})
+{
+    __impl::__check_esimd_sort_params<__radix_bits, _KernelParam::data_per_workitem, _KernelParam::workgroup_size>();
+    if (__keys_rng.size() <= 0)


Why you wrote here <= ?
Usually we use the code like

if (__n == 0) return;

fixed, thanks.

SergeyKopienko · 2024-03-26T08:41:38Z

include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_dispatchers.h


    __event_chain = __radix_sort_onesweep_scan_submitter<__stage_count, __bin_count, _EsimdRadixSortScan>()(
        __q, __mem_holder.__global_hist_ptr(), __n, __event_chain);

-    for (::std::uint32_t __stage = 0; __stage < __stage_count; __stage++)
+    auto __submit_iteration = [&](auto __input_, auto __output_, auto __p_global_hist_, auto __p_group_hists_,


Probably the constructions like this may decrease performance, I think.
Sometimes they are required, of course. But in this place it's sooks like for me as the simple code style.
Of course, it's only my opinion.

My intention is to reduce repetitive code to make more apparent what is actually different between these invocations of __radix_sort_onesweep_submitter, but I don't think it is required technically.

Also, my changes unroll the first iteration of the submission before the loop, as its inputs / outputs are special cased so there are three redundant calls, rather than just two.

I can go back to something more similar to the previous way of doing it without the lambda. I don't have a strong preference here.

I've removed the lambda __submit_iteration.

SergeyKopienko · 2024-03-26T08:47:29Z

include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_dispatchers.h

+template <::std::uint8_t __radix_bits, ::std::uint16_t __data_per_work_item, ::std::uint16_t __work_group_size,
+          typename _RngPack, typename _MemHolder>
+void
+__allocate_temp(_MemHolder& __mem_holder, ::std::size_t __n)


May be better to return _MemHolder instance from this method?
In this case we are able to not declare some their methods as const...

I originally abstracted these two functions out to be separate functions to be able to reuse from multiple places, but then later unified out_of_place and in_place versions together anyway pulling that decision up the stack, so they don't even need to be separate functions. I'll return them both to be within __onesweep

I've restored this to be within the __onesweep function.

SergeyKopienko · 2024-03-26T09:01:20Z

test/kt/esimd_radix_sort_by_key_out_of_place.cpp

+    auto expected_first = oneapi::dpl::make_zip_iterator(std::begin(expected_keys), std::begin(expected_values));
+    std::stable_sort(expected_first, expected_first + size, CompareKey<isAscending>{});
+
+    std::string parameters_msg = ", n: " + std::to_string(size) + ", sizeof(key): " + std::to_string(sizeof(KeyT)) +


If you use std::strstream you may to build this message without a lot of std::to_string calls.

fixed, thanks.

SergeyKopienko · 2024-03-26T09:07:46Z

include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_dispatchers.h

    using _EsimRadixSortKernel =
        oneapi::dpl::__par_backend_hetero::__internal::__kernel_name_provider<__esimd_radix_sort_one_wg<_KernelName>>;

    return __radix_sort_one_wg_submitter<__is_ascending, __radix_bits, __data_per_work_item, __work_group_size, _KeyT,
-                                         _EsimRadixSortKernel>()(__q, ::std::forward<_RngPack>(__pack), __n);
+                                         _EsimRadixSortKernel>()(__q, __pack_in, __pack_out, __n);


Why you not use std::forward for pass __pack_in and __pack_out ?

fixed, thanks.

include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_dispatchers.h

Signed-off-by: Dan Hoeflinger <[email protected]>

include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_dispatchers.h

dmitriy-sobolev

The PR looks good to me besides _out_of_place function name part, which is mentioned in the description.

akukanov · 2024-04-17T15:23:12Z

include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_dispatchers.h

+        if (__stage % 2 != 0)
+        {
+            __event_chain = __radix_sort_onesweep_submitter<__is_ascending, __radix_bits, __data_per_work_item,
+                                                            __work_group_size, _EsimdRadixSortSweep>()(
+                __q, __virt_pack1, __virt_pack2, __p_global_hist, __p_group_hists, __sweep_work_group_count, __n,
+                __stage, __event_chain);
+        }
+        else
+        {
+            __event_chain = __radix_sort_onesweep_submitter<__is_ascending, __radix_bits, __data_per_work_item,
+                                                            __work_group_size, _EsimdRadixSortSweep>()(
+                __q, __virt_pack2, __virt_pack1, __p_global_hist, __p_group_hists, __sweep_work_group_count, __n,
+                __stage, __event_chain);
+        }


Unless I miss something obvious, these two kernel submissions will not generate identical kernels, and so should really have distinct names (like with the Even/Odd suffixes before the patch).

Now in hindsight, I think you are correct, because the types of __input_pack, __virt_pack1 and __virt_pack2 may not match.

However, if they do match, do you agree that these would generate an identical kernel?

My proposal would be to merely include the types of the input & output in the kernel name, that way they aren't required to generate two identical kernels with different names when the types match.

I'm doing some testing with no unnamed lambda, and seeing the issue. I'll provide an update when it is fixed.

I've fixed the issue as proposed by including the decayed RngPack typenames in the kernel name.
I've also fixed kernel name clashing in the tests for this and the in-place version of esimd radix sort, and added a test utility to generate a new kernel_params struct with a new kernelname.

Thanks for catching this.

Signed-off-by: Dan Hoeflinger <[email protected]>

danhoeflinger · 2024-04-19T01:37:08Z

The PR looks good to me besides _out_of_place function name part, which is mentioned in the description.

@dmitriy-sobolev
I've updated the APIs to use enable_if to disambiguate the range and iterator versions. If we decide to remove the range interfaces, it will be an easy change from the current state of the PR. Please take a look at the __is_iterator_type_v check.

dmitriy-sobolev

The approach to get rid of out_of_place suffix looks good to me. I've left some minor comments, though.

include/oneapi/dpl/experimental/kt/esimd_radix_sort.h

include/oneapi/dpl/pstl/utils.h

dmitriy-sobolev · 2024-04-22T16:09:06Z

include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_dispatchers.h

+    using _EsimdRadixSortSweepInitial = oneapi::dpl::__par_backend_hetero::__internal::__kernel_name_provider<
+        ::std::conditional_t<__has_values, __esimd_radix_sort_onesweep_by_key<std::decay_t<_RngPack1>, std::decay_t<_RngPack2>, _KernelName>,
+                             __esimd_radix_sort_onesweep<std::decay_t<_RngPack1>, std::decay_t<_RngPack2>, _KernelName>>>;


I'm not strictly opposed of that approach, but there is a drawback: kernel names will become longer (I guess, much longer).

The motivation of the change is to separate kernel names for the potential arrangements of _RngPack1, _RngPack2, _RngPack3, isn't it? If so, what do you think about using some constant instead, e.g.:

__esimd_radix_sort_onesweep<0, ...> instead of __esimd_radix_sort_onesweep<std::decay_t<_RngPack1>, std::decay_t<_RngPack2>, ...>,

__esimd_radix_sort_onesweep<1, ...> instead of __esimd_radix_sort_onesweep<std::decay_t<_RngPack3>, std::decay_t<_RngPack2>, ...>?

The goal is to have a unique kernel name if and only if we have a unique kernel. We want to avoid same mangled name error, but also we don't want to add extra kernel compilation only to support a differently named kernel when the types match exactly.

Lets say _RngPack1 is the exact same as _RngPack3, we then want __esimd_radix_sort_onesweep<std::decay_t<_RngPack1>, std::decay_t<_RngPack2>, ...>
and __esimd_radix_sort_onesweep<std::decay_t<_RngPack3>, std::decay_t<_RngPack2>, ...> to be the exact same as well. If not, I think we require an extra kernel compilation only to satisfy having a different name.

If all three exactly match, we only want one kernel and one kernel name.

That's reasonable, thanks for describing.

Probably, it is still possible to use shorter kernel names. I will think about it separately as it is rather a "nice to have" thing.

Currently, I see the following options:

Rng1 type Rng2 type Rng3 type Kernels Constants

A B C A->B, B->C, C->B 0,1, 2

A A A A->A, A->A, A->A 3, 3, 3

A A B A->A, A->B, B->A 3, 0, 4

A B A A->B, B->A, A->B 0, 4, 0

B A A B->A, A->A, A->A 4, 3, 3

I'm interested in the idea, but to create such constants, you would need to basically check is_same<decay_t<_RngPack1>, decay_t<_RngPack2>> for all three combos and assign constants appropriately.
The code is ugly either way. Perhaps the shorter kernel name is worth it.

include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_dispatchers.h

include/oneapi/dpl/experimental/kt/esimd_radix_sort.h

dmitriy-sobolev

I've checked the code once again. I'm ready to approve once minor issues noted above have been solved.

Signed-off-by: Dan Hoeflinger <[email protected]>

dmitriy-sobolev

LGTM

danhoeflinger requested a review from dmitriy-sobolev March 11, 2024 20:59

danhoeflinger marked this pull request as ready for review March 12, 2024 13:01

danhoeflinger added 4 commits March 20, 2024 13:47

out of place esimd radix sort and tests

d854024

Signed-off-by: Dan Hoeflinger <[email protected]>

refactor to consolidate redundant code

8c7cb76

Signed-off-by: Dan Hoeflinger <[email protected]>

bugfix for in place mode

3c7e048

Signed-off-by: Dan Hoeflinger <[email protected]>

remove formatting only changes

12cd664

Signed-off-by: Dan Hoeflinger <[email protected]>

danhoeflinger force-pushed the dev/dhoeflin/out_of_place_esimd_radixsort branch from 49465ec to 12cd664 Compare March 20, 2024 17:47

danhoeflinger added this to the 2022.6.0 milestone Mar 25, 2024

danhoeflinger requested review from julianmi, timmiesmith, MikeDvorskiy, SergeyKopienko, mmichel11 and adamfidel March 25, 2024 14:35

SergeyKopienko reviewed Mar 26, 2024

View reviewed changes

adamfidel reviewed Mar 26, 2024

View reviewed changes

include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_dispatchers.h Outdated Show resolved Hide resolved

mmichel11 reviewed Mar 26, 2024

View reviewed changes

include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_dispatchers.h Outdated Show resolved Hide resolved

danhoeflinger added 9 commits March 27, 2024 13:06

remove copies and simplify constexpr pack selection

5b36d4b

Signed-off-by: Dan Hoeflinger <[email protected]>

switching to n==0

66386db

Signed-off-by: Dan Hoeflinger <[email protected]>

returning allocation and temp pack creation to onesweep func

2e54fab

Signed-off-by: Dan Hoeflinger <[email protected]>

forwarding packs

ed55ee3

Signed-off-by: Dan Hoeflinger <[email protected]>

removing submit_iteration lambda

5e0ce6a

Signed-off-by: Dan Hoeflinger <[email protected]>

using string streams instead of multiple to_string calls

92984ec

Signed-off-by: Dan Hoeflinger <[email protected]>

adding missing .str()

b58622c

Signed-off-by: Dan Hoeflinger <[email protected]>

formatting

75de1f7

Signed-off-by: Dan Hoeflinger <[email protected]>

clang format

cd22360

Signed-off-by: Dan Hoeflinger <[email protected]>

dmitriy-sobolev reviewed Apr 16, 2024

View reviewed changes

include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_dispatchers.h Show resolved Hide resolved

dmitriy-sobolev reviewed Apr 16, 2024

View reviewed changes

akukanov reviewed Apr 17, 2024

View reviewed changes

enable_if iterator check for clashing

fef6e96

Signed-off-by: Dan Hoeflinger <[email protected]>

danhoeflinger added 5 commits April 18, 2024 16:43

fix for kernel name clashing

6444081

Signed-off-by: Dan Hoeflinger <[email protected]>

fix for unnamed lambda

4213e17

Signed-off-by: Dan Hoeflinger <[email protected]>

formatting

54cbcfd

Signed-off-by: Dan Hoeflinger <[email protected]>

formatting

88b5821

Signed-off-by: Dan Hoeflinger <[email protected]>

wrapping KT kernelname in dpcpp backend

a9fb79c

Signed-off-by: Dan Hoeflinger <[email protected]>

danhoeflinger mentioned this pull request Apr 19, 2024

Single-pass scan kernel template #1320

Merged

dmitriy-sobolev reviewed Apr 22, 2024

View reviewed changes

include/oneapi/dpl/experimental/kt/esimd_radix_sort.h Show resolved Hide resolved

include/oneapi/dpl/pstl/utils.h Outdated Show resolved Hide resolved

dmitriy-sobolev reviewed Apr 22, 2024

View reviewed changes

include/oneapi/dpl/experimental/kt/internal/esimd_radix_sort_dispatchers.h Outdated Show resolved Hide resolved

dmitriy-sobolev reviewed Apr 22, 2024

View reviewed changes

include/oneapi/dpl/experimental/kt/esimd_radix_sort.h Show resolved Hide resolved

dmitriy-sobolev reviewed Apr 22, 2024

View reviewed changes

address review feedback

96ccbda

Signed-off-by: Dan Hoeflinger <[email protected]>

dmitriy-sobolev approved these changes Apr 23, 2024

View reviewed changes

danhoeflinger merged commit 6e7b290 into main Apr 23, 2024
20 checks passed

danhoeflinger deleted the dev/dhoeflin/out_of_place_esimd_radixsort branch April 23, 2024 12:10

danhoeflinger mentioned this pull request May 14, 2024

[Docs] Out of place esimd radix sort documentation update #1583

Merged

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Out of Place Esimd Radix Sort and Tests #1439

Out of Place Esimd Radix Sort and Tests #1439

danhoeflinger commented Mar 11, 2024 •

edited

Loading

SergeyKopienko Mar 26, 2024

danhoeflinger Mar 27, 2024

SergeyKopienko Mar 26, 2024

danhoeflinger Mar 27, 2024

danhoeflinger Mar 27, 2024

SergeyKopienko Mar 26, 2024

danhoeflinger Mar 27, 2024

danhoeflinger Mar 27, 2024

SergeyKopienko Mar 26, 2024

danhoeflinger Mar 27, 2024

SergeyKopienko Mar 26, 2024

danhoeflinger Mar 27, 2024

dmitriy-sobolev left a comment

akukanov Apr 17, 2024

danhoeflinger Apr 17, 2024 •

edited

Loading

danhoeflinger Apr 17, 2024

danhoeflinger Apr 18, 2024

danhoeflinger commented Apr 19, 2024 •

edited

Loading

dmitriy-sobolev left a comment

dmitriy-sobolev Apr 22, 2024

danhoeflinger Apr 22, 2024

dmitriy-sobolev Apr 22, 2024

danhoeflinger Apr 22, 2024

dmitriy-sobolev left a comment

dmitriy-sobolev left a comment

Rng1 type	Rng2 type	Rng3 type	Kernels	Constants
A	B	C	A->B, B->C, C->B	0,1, 2
A	A	A	A->A, A->A, A->A	3, 3, 3
A	A	B	A->A, A->B, B->A	3, 0, 4
A	B	A	A->B, B->A, A->B	0, 4, 0
B	A	A	B->A, A->A, A->A	4, 3, 3

Out of Place Esimd Radix Sort and Tests #1439

Out of Place Esimd Radix Sort and Tests #1439

Conversation

danhoeflinger commented Mar 11, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

dmitriy-sobolev left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

danhoeflinger Apr 17, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

danhoeflinger commented Apr 19, 2024 • edited Loading

dmitriy-sobolev left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

dmitriy-sobolev left a comment

Choose a reason for hiding this comment

dmitriy-sobolev left a comment

Choose a reason for hiding this comment

danhoeflinger commented Mar 11, 2024 •

edited

Loading

danhoeflinger Apr 17, 2024 •

edited

Loading

danhoeflinger commented Apr 19, 2024 •

edited

Loading