Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[SYCL][Group algorithms] Add group sorting algorithms implementation #4439

Merged
merged 17 commits into from
Sep 15, 2021
Merged

[SYCL][Group algorithms] Add group sorting algorithms implementation #4439

merged 17 commits into from
Sep 15, 2021

Conversation

andreyfe1
Copy link
Contributor

@andreyfe1 andreyfe1 commented Aug 31, 2021 •
edited
Loading

This PR introduce the implementation for the Group Sort extension.

The PR includes:

  • Feature macro
  • joint_sort and sort_over_group functions
  • default_sorter
  • group_with_scratchpad

Algorithms are quite general. It should work with custom data types, custom comparators, custom sorters.

The PR doesn't include:

  • radix_sorter
  • optimized specialization for arithmetic types

Tests are here: intel/llvm-test-suite#438

Signed-off-by: Fedorov, Andrey andrey.fedorov@intel.com

@andreyfe1
initial commit for group sort implementation
a8a5b33 
Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
@andreyfe1 andreyfe1 requested a review from a team as a code owner August 31, 2021 12:21
@andreyfe1 andreyfe1 requested a review from s-kanaev August 31, 2021 12:21
@bader bader marked this pull request as draft August 31, 2021 12:24
@bader
Copy link
Contributor

bader commented Aug 31, 2021

@capatober, please, create draft pull request for "work in progress".

@andreyfe1
fixed compilation issues
c27f0ac 
Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
@andreyfe1 andreyfe1 changed the title WIP: [Group algorithms]: Initial implementation for Group Sorting Algorithm WIP: [SYCL][Group algorithms]: Initial implementation for Group Sorting Algorithm Aug 31, 2021
@andreyfe1
Merge remote-tracking branch 'upstream/sycl' into group_sort_impl
a1d9c09
@andreyfe1
applied clang-format
0159ff8 
Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
@andreyfe1
applied clang-format
d8aab75 
Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
@andreyfe1
applied clang-format
f40e9cc 
Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
@andreyfe1 andreyfe1 marked this pull request as ready for review September 1, 2021 09:10
@andreyfe1 andreyfe1 changed the title WIP: [SYCL][Group algorithms]: Initial implementation for Group Sorting Algorithm [SYCL][Group algorithms]: Initial implementation for Group Sorting Algorithm Sep 1, 2021
@andreyfe1
Copy link
Contributor Author

The PR is ready for review.
clang-format check differs from clang-format from the repo. Seems it should be improved on infrastructure side.

@andreyfe1 andreyfe1 mentioned this pull request Sep 2, 2021
Merged
@andreyfe1
fixed when tests are passed
6b55761 
Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
@andreyfe1
Copy link
Contributor Author

@intel/llvm-reviewers-runtime, just kind reminder. Patch is ready for review

@andreyfe1
found that fence_scope is implemented
cb54021 
Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
alexbatashev
alexbatashev reviewed Sep 3, 2021
View reviewed changes
sycl/include/CL/sycl/detail/group_sort_impl.hpp Outdated Show resolved Hide resolved
sycl/include/CL/sycl/detail/group_sort_impl.hpp Outdated Show resolved Hide resolved
sycl/include/CL/sycl/detail/group_sort_impl.hpp Show resolved Hide resolved
sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp Outdated
Comment on lines 34 to 35
std::uint8_t *scratch;
std::size_t scratch_size;
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is very similar to my previous comment. A pair of a pointer to an array and its size is either range or span. Since we have span backported to SYCL 2020, can we use it here?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We consider this option previously during the API discussion. The thing is that if we replaced this with sycl::span, we would need an additional template parameter for sorter (extent). However, extent is not relevant to the default_sorter, but to the parameter of its ctor.

sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp Outdated Show resolved Hide resolved
sycl/include/sycl/ext/oneapi/group_sort.hpp Outdated Show resolved Hide resolved
sycl/include/sycl/ext/oneapi/group_sort.hpp Outdated Show resolved Hide resolved
sycl/include/sycl/ext/oneapi/group_sort.hpp Outdated Show resolved Hide resolved
@andreyfe1
applied comments from reviewers
68f2d8e 
Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
@andreyfe1
fix for get_pointer usage
339d043 
Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
@andreyfe1
align with clang-format from infrastructure
eddd7a3 
Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
cperkinsintel
cperkinsintel reviewed Sep 3, 2021
View reviewed changes
Copy link
Contributor

@cperkinsintel cperkinsintel left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Does this extension have a doc (under llvm/sycl/doc/extensions) ?

sycl/include/sycl/ext/oneapi/experimental/group_helpers_sorters.hpp Outdated Show resolved Hide resolved
sycl/include/CL/sycl/detail/group_sort_impl.hpp Show resolved Hide resolved
@cperkinsintel
Copy link
Contributor

Also, we need some tests.

@andreyfe1
Copy link
Contributor Author

Does this extension have a doc (under llvm/sycl/doc/extensions) ?

Sure. It's mentioned in the description of the PR. (https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_sort.asciidoc)

Also, we need some tests.

It's also mentioned in the description :) See here: intel/llvm-test-suite#438

@andreyfe1 andreyfe1 mentioned this pull request Sep 6, 2021
Closed
@andreyfe1
apply some comments from review
b4aa9c3 
Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
@andreyfe1
replace more runtime_errors with sycl::exception
768d3a3 
Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
@andreyfe1
add placement new instead of assignment
972f6ef 
Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
@andreyfe1 andreyfe1 requested a review from a team as a code owner September 6, 2021 10:26
@andreyfe1
remove comments
8a07d59 
Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
alexbatashev
alexbatashev previously approved these changes Sep 6, 2021
View reviewed changes
Copy link
Contributor

@alexbatashev alexbatashev left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM

@andreyfe1
Copy link
Contributor Author

@intel/llvm-reviewers-runtime, is anything else needed to merge the code to the repo? I don't have enough rights for this action

@alexbatashev
Copy link
Contributor

@intel/llvm-reviewers-runtime, is anything else needed to merge the code to the repo? I don't have enough rights for this action

At least you need to make sure CI is green. Also approve from @intel/dpcpp-specification-reviewers is required.

@andreyfe1
add more macros
b6ae2e5 
Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
@andreyfe1
Copy link
Contributor Author

@intel/llvm-reviewers-runtime, CI is green. Feel free to approve and merge. Changes for spec are obvious and doesn't require to be approved additionally

alexbatashev
alexbatashev previously approved these changes Sep 14, 2021
View reviewed changes
@bader bader requested a review from Pennycook September 14, 2021 08:59
@andreyfe1
made changes for multi-dimensional groups
e162168 
Signed-off-by: Fedorov, Andrey <andrey.fedorov@intel.com>
@andreyfe1 andreyfe1 requested a review from Pennycook September 15, 2021 08:06
alexbatashev
alexbatashev approved these changes Sep 15, 2021
View reviewed changes
Copy link
Contributor

@alexbatashev alexbatashev left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Still LGTM

@bader bader changed the title [SYCL][Group algorithms]: Initial implementation for Group Sorting Algorithm [SYCL][Group algorithms] Add group sorting algorithms implementation Sep 15, 2021
@romanovvlad romanovvlad merged commit 932ae56 into intel:sycl Sep 15, 2021
alexbatashev added a commit to alexbatashev/llvm that referenced this pull request Sep 18, 2021
@alexbatashev
Merge remote-tracking branch 'upstream/sycl' into xpti_tests_enable
bc2716c 
* upstream/sycl: (36 commits)
  [SYCL] Add SYCL2020 target::device enumeration value (intel#4587)
  [SYCL][Doc] Update ITT instrumentation docs (intel#4503)
  [SYCL][L0] Make all L0 events have device-visibility (intel#4534)
  [SYCL] Updated Level-Zero backend spec according to SYCL 2020 standard (intel#4560)
  [SYCL] Add error_code support for SYCL 1.2.1 exception classes (intel#4574)
  [SYCL][CI] Provide --ci-defaults option for config script (intel#4583)
  [CI] Switch GitHub Actions to Ubuntu 20.04 (intel#4582)
  [SYCL][CUDA] Fix context clearing in PiCuda tests (intel#4483)
  [SYCL] Hide SYCL service kernels (intel#4519)
  [SYCL][L0] Fix mismatched ZE call count (intel#4559)
  [SYCL] Remove function pointers extension (intel#4459)
  [GitHub Actions] Uplift clang version in post-commit validation (intel#4581)
  [SYCL] Ignore usm prefetch dummy flag (intel#4568)
  [SYCL][Group algorithms] Add group sorting algorithms implementation (intel#4439)
  [SYCL] Resolve name clash with a user defined symbol (intel#4570)
  [clang-offload-wrapper] Do not create .tgtimg section with -emit-reg-funcs=0 (intel#4577)
  [SYCL][FPGA] Remove deprecated attribute functionality (intel#4532)
  [SYCL] Remove _class aliases (intel#4465)
  [SYCL][CUDA][HIP] Report every device in its own platform (intel#4571)
  [SYCL][L0] make_device shouldn't need platform as an input (intel#4561)
  ...
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@cperkinsintel cperkinsintel cperkinsintel left review comments

@v-klochkov v-klochkov v-klochkov left review comments

@alexbatashev alexbatashev alexbatashev approved these changes

@Pennycook Pennycook Pennycook approved these changes

@s-kanaev s-kanaev Awaiting requested review from s-kanaev

Assignees
No one assigned
Labels
None yet
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
7 participants
@andreyfe1 @bader @cperkinsintel @alexbatashev @Pennycook @v-klochkov @romanovvlad