[SYClomatic] Add cub::RadixSort.SortBlockedToStriped header by abhilash1910 · Pull Request #1562 · oneapi-src/SYCLomatic · GitHub

PR to add SortBlockedToStriped sorting ,
depends on #1483 (scatter_to_striped added in exchange)
cc @mmichel11 @danhoeflinger @yihanwg

In addition to the comments here, I think a test case would also be nice to ensure correctness in the SYCLomatic-test repo.

Do we need to synchronize here before we exit the loop and call twiddle_out over the keys?

I do not think a local barrier is required here since we call a barrier outside the exchange loop before twiddling out.

My understanding is that twiddle_in and twiddle_out throw a runtime exception when it is not a supported type. Maybe it is not worth addressing here, but would it be better to catch this at compile time instead based on the type of T?

Thanks @mmichel11 for the suggestions, I think the modifications can be addressed later as well since this PR is for basic functionality . I am in progress of adding test suite for this sorting algos. Since there is a feature freeze on 8th, if it looks ok I can put it in for merging . Any modifications on this can be addressed on another PR as well . Let me know your thoughts. Thanks.

Also tagging @zhimingwang36 for review.

Why should we convert to uint32_t instead of using T directly?

How was this resolved?
(In general, I'd recommend letting the reviewer resolve threads when they are satisfied, rather than resolving them yourself.)

@abhilash1910 Thank you working on this, can you please also provide a test patch in SYCLomatic-test repo?

At a high level, is the only difference between this and the radix_sort function that we are doing the final exchange from scatter_to_striped rather than scatter_to_blocked?

If so, it seems like we shouldn't copy / rewrite the entire function to accomplish this. We should be able to just adjust that final exchange and rely on the same code otherwise. This can be done a number of ways, with a simple helper and a branch, or via some functor / lambda.

At a high level, is the only difference between this and the radix_sort function that we are doing the final exchange from scatter_to_striped rather than scatter_to_blocked?

If so, it seems like we shouldn't copy / rewrite the entire function to accomplish this. We should be able to just adjust that final exchange and rely on the same code otherwise. This can be done a number of ways, with a simple helper and a branch, or via some functor / lambda.

Yeah thought so, reduced it to a helper + branch. Addressed other issues as well.
Can this be reviewed on priority ? @danhoeflinger @mmichel11 @yihanwg .Thanks

@abhilash1910 I am seeing clang-format failing in CI. Could you resolve this?

@mmichel11 @danhoeflinger fixed the issues. Lets wait for CI to run . Thanks.

LGTM

Add SortBlockedToStriped header

438c7ab

abhilash1910 requested a review from a team as a code owner December 27, 2023 06:44

abhilash1910 requested review from tangjj11 and wenhuiNi December 27, 2023 06:44

abhilash1910 marked this pull request as draft December 27, 2023 06:44

abhilash1910 added 2 commits January 2, 2024 16:13

clang-format

4dc9b75

Merge branch 'oneapi-src:SYCLomatic' into block_radixsort_variants

08087b4

abhilash1910 marked this pull request as ready for review January 2, 2024 10:44

clang-format

c50035d

mmichel11 reviewed Jan 4, 2024

View reviewed changes

yihwang-nv reviewed Jan 4, 2024

View reviewed changes

Comment thread clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h Outdated

mmichel11 reviewed Jan 4, 2024

View reviewed changes

Comment thread clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h Outdated

yihwang-nv reviewed Jan 4, 2024

View reviewed changes

review commits/bugs

7fb2c48

danhoeflinger reviewed Jan 4, 2024

View reviewed changes

abhilash1910 added 2 commits January 5, 2024 11:49

update loop and remove repeat code

5a96f80

clang-format

63e1f66

danhoeflinger reviewed Jan 5, 2024

View reviewed changes

Comment thread clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h Outdated

danhoeflinger reviewed Jan 5, 2024

View reviewed changes

Comment thread clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h

mmichel11 reviewed Jan 5, 2024

View reviewed changes

Comment thread clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h Outdated

mmichel11 reviewed Jan 5, 2024

View reviewed changes

Comment thread clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h Outdated

danhoeflinger reviewed Jan 5, 2024

View reviewed changes

Comment thread clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h Outdated

abhilash1910 added 2 commits January 5, 2024 23:51

fix bugs and review

d28e838

clang

287f4ed

mmichel11 reviewed Jan 5, 2024

View reviewed changes

Comment thread clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h

abhilash1910 added 2 commits January 7, 2024 22:57

fix bugs

f711383

clang-format

b89badb

abhilash1910 added 3 commits January 8, 2024 09:08

Merge branch 'oneapi-src:SYCLomatic' into block_radixsort_variants

7fe2046

clang-format

a3c1502

clang

5810fb8

danhoeflinger reviewed Jan 8, 2024

View reviewed changes

Comment thread clang/runtime/dpct-rt/include/dpct/dpl_extras/dpcpp_extensions.h

review commits

1b60be4

abhilash1910 added 2 commits January 8, 2024 22:08

format & clean

d7d22df

format

38b58db

danhoeflinger approved these changes Jan 8, 2024

View reviewed changes

mmichel11 approved these changes Jan 8, 2024

View reviewed changes

zhimingwang36 merged commit aeee36e into oneapi-src:SYCLomatic Jan 9, 2024

abhilash1910 mentioned this pull request Jan 17, 2024

[SYCLomatic-test] Add Help Function for Sort Headers oneapi-src/SYCLomatic-test#601

Merged

abhilash1910 mentioned this pull request Apr 14, 2024

[SYCLomatic] Create separate util file for groups #1784

Merged

+                    if (begin_bit >= end_bit) {
+                      exchange<T, VALUES_PER_THREAD>(_local_memory)
+                          .scatter_to_striped(item, keys);

+              #pragma unroll
+                  for (int i = 0; i < VALUES_PER_THREAD; ++i) {
+                    unsigned_keys[i] = detail::traits<T>::twiddle_in(unsigned_keys[i]);

+                                     int begin_bit = 0, int end_bit = 8 * sizeof(T)) {
+                  uint32_t(&unsigned_keys)[VALUES_PER_THREAD] =
+                      reinterpret_cast<uint32_t(&)[VALUES_PER_THREAD]>(keys);

[SYClomatic] Add cub::RadixSort.SortBlockedToStriped header by abhilash1910 · Pull Request #1562 · oneapi-src/SYCLomatic · GitHub

Uh oh!

[SYClomatic] Add cub::RadixSort.SortBlockedToStriped header - #1562

[SYClomatic] Add cub::RadixSort.SortBlockedToStriped header#1562zhimingwang36 merged 17 commits intooneapi-src:SYCLomaticoneapi-src/SYCLomatic:SYCLomaticfrom abhilash1910:block_radixsort_variantsabhilash1910/SYCLomatic:block_radixsort_variantsCopy head branch name to clipboard

Conversation

abhilash1910 commented Dec 27, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

mmichel11 left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

mmichel11 Jan 4, 2024

Choose a reason for hiding this comment

Uh oh!

abhilash1910 Jan 4, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

mmichel11 Jan 4, 2024

Choose a reason for hiding this comment

Uh oh!

abhilash1910 commented Jan 4, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

Uh oh!

yihwang-nv Jan 4, 2024

Choose a reason for hiding this comment

Uh oh!

danhoeflinger Jan 8, 2024

Choose a reason for hiding this comment

Uh oh!

yihwang-nv commented Jan 4, 2024

Uh oh!

danhoeflinger left a comment

Choose a reason for hiding this comment

Uh oh!

abhilash1910 commented Jan 5, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

mmichel11 commented Jan 8, 2024

Uh oh!

abhilash1910 commented Jan 8, 2024

Uh oh!

danhoeflinger left a comment

Choose a reason for hiding this comment

Uh oh!

mmichel11 left a comment

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

[SYClomatic] Add cub::RadixSort.SortBlockedToStriped header#1562
zhimingwang36 merged 17 commits into
oneapi-src:SYCLomaticfrom
abhilash1910:block_radixsort_variants

abhilash1910 commented Dec 27, 2023 •
edited
Loading

abhilash1910 Jan 4, 2024 •
edited
Loading

abhilash1910 commented Jan 4, 2024 •
edited
Loading

abhilash1910 commented Jan 5, 2024 •
edited
Loading