[SYCL] Add support for SYCL 2020 in class group #5447

aobolensk · 2022-02-01T16:17:00Z

Implement missing methods for class group according to SYCL 2020 4.9.1.7.

Pennycook · 2022-02-01T16:53:10Z

This change is good, but I think if we're going to add some SYCL 2020 functionality to the group class we should go all the way. If we only partially implement the SYCL 2020 interface, developers might be confused.

The SYCL 2020 group class is documented here: https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#group-class, and I don't think many extra changes are required:

Deprecate get_id()
Add get_local_id()
Add get_max_local_range()
Add get_group_linear_id()
Add get_local_linear_id()
Add get_group_linear_range()
Add get_local_linear_range()
Add leader()

aobolensk · 2022-02-10T09:44:13Z

@Pennycook added functions you asked for in your comment. Also extended tests for group and added test for local_id on device (intel/llvm-test-suite#821).
Limitation: get_local_id() host implementation requires ABI breaking change, so it was stubbed with an exception. On device it works well.

Pennycook · 2022-02-10T15:32:06Z

sycl/include/CL/sycl/group.hpp

+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 1), size_t>
+  get_local_linear_id() const {
+    id<Dimensions> localId = get_local_id();
+    return localId[0];
+  }


I'd like @gmlueck to double-check this. Greg, can this be declared in our headers as a template function if it isn't a template function in the specification? My concern here is that this might be user-observable: a user could call get_local_linear_id<Dimensions + 1>, for example, or might have to insert extra template keywords in some situations.

If I'm right to be concerned, I think we can work around this by moving the specialization to a helper function, something like the following:

namespace sycl { namespace detail { size_t get_local_linear_id(sycl::group<0>& g) const { return g.get_local_id(); } // + Specializations for group<1> and group<2> } } size_t get_local_linear_id() const { return sycl::detail::get_local_linear_id(*this); }

Yes, it does seem weird that we would implement this function as a template when it is not that way in the spec. It seems like it would be easy to fix by calling some separate templated function as John proposes. I assume the separate function could also be a private member function of group?

class group { public: size_t get_local_linear_id() const { return get_local_linear_id_impl<Dimensions>(); } private: template <int dims> typename detail::enable_if_t<(dims == 1), size_t> get_local_linear_id_impl() const { id<Dimensions> localId = get_local_id(); return localId[0]; } }

Or, you could use constexpr-if, but that would require the compiler to be C++17.

Pennycook · 2022-02-10T15:36:35Z

sycl/include/CL/sycl/group.hpp

+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 1), range<Dimensions>>
+  get_max_local_range() const {
+    return range<Dimensions>{
+        static_cast<size_t>(info::device::max_work_group_size)};
+  }
+
+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 2), range<Dimensions>>
+  get_max_local_range() const {
+    return range<Dimensions>{
+        static_cast<size_t>(info::device::max_work_group_size),
+        static_cast<size_t>(info::device::max_work_group_size)};
+  }
+
+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 3), range<Dimensions>>
+  get_max_local_range() const {
+    return range<Dimensions>{
+        static_cast<size_t>(info::device::max_work_group_size),
+        static_cast<size_t>(info::device::max_work_group_size),
+        static_cast<size_t>(info::device::max_work_group_size)};
+  }


I appreciate the specification could be clearer here, but the intent of this function is to return the maximum number of work-items in any work-group for the kernel that is currently executing, rather than the maximum number of work-items supported by the device.

Because we don't support non-uniform work-group sizes, I think you can replace all of this with:

size_t get_max_local_range() const { return get_local_range(); }

The reason both functions exist is that they return different values for the sub_group class. Supporting both in group allows developers to write generic code accepting either a sub_group or group.

Pennycook · 2022-02-10T15:38:16Z

sycl/include/CL/sycl/group.hpp

    return (index[0] * groupRange[1] * groupRange[2]) +
           (index[1] * groupRange[2]) + index[2];
  }

+  bool leader() const { return false; }


The specification for this function says:

Return true for exactly one work-item in the work-group, if the calling work-item is the leader of the work-group, and false for all other work-items in the work-group.

The leader of the work-group is determined during construction of the work-group, and is invariant for the lifetime of the work-group. The leader of the work-group is guaranteed to be the work-item with a local id of 0.

Because of the last sentence, I think the correct implementation is:

bool leader() const { return (get_local_linear_id() == 0); }

sycl/include/CL/sycl/group.hpp

bader · 2022-02-11T13:37:39Z

@alexbatashev, take a look at this failure - https://github.com/intel/llvm/runs/5155471215?check_suite_focus=true. I think we should improve the CI scripts to avoid this.

Pennycook

Sorry for not being clear in my initial review -- my comment about get_local_linear_id() applies to all of the new template functions you've introduced here. I think there are three such functions left, and I've called out the first occurrence of each.

Pennycook · 2022-02-11T15:07:47Z

sycl/include/CL/sycl/group.hpp

  range<Dimensions> get_local_range() const { return localRange; }

  size_t get_local_range(int dimension) const { return localRange[dimension]; }

+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 1), size_t>
+  get_local_linear_range() const {


This should call a get_local_linear_range_impl().

Pennycook · 2022-02-11T15:08:03Z

sycl/include/CL/sycl/group.hpp

  range<Dimensions> get_group_range() const { return groupRange; }

  size_t get_group_range(int dimension) const {
    return get_group_range()[dimension];
  }

+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 1), size_t>
+  get_group_linear_range() const {


This should call a get_group_linear_range_impl().

Pennycook · 2022-02-11T15:08:27Z

sycl/include/CL/sycl/group.hpp

+
+  template <int dims = Dimensions>
+  typename detail::enable_if_t<(dims == 1), size_t>
+  get_group_linear_id() const {


This should call a get_group_linear_id_impl().

aobolensk · 2022-02-11T15:53:07Z

Sorry for not being clear in my initial review -- my comment about get_local_linear_id() applies to all of the new template functions you've introduced here. I think there are three such functions left, and I've called out the first occurrence of each.

Done for rest of the functions

againull · 2022-02-11T18:57:05Z

@aobolensk Could you please check if this failure is related to your PR:

Failed Tests (1):
SYCL :: Basic/handler/handler_mem_op.cpp

againull · 2022-02-11T19:02:28Z

@aobolensk Could you please check if this failure is related to your PR:

Failed Tests (1): SYCL :: Basic/handler/handler_mem_op.cpp

Never mind, it is unrelated. Found a PR isolated to L0 plugin which has the same failure: #5541

sycl/include/CL/sycl/group.hpp

…/llvm into refactor_existing_workflows * 'refactor_existing_workflows' of github.com:alexbatashev/llvm: (2051 commits) [SYCL][L0] Honor property::queue::enable_profiling (intel#5543) [SYCL][CI] Enable sccache on Windows (intel#5589) [SYCL][Doc] Move internal design docs (intel#5556) [sycl-post-link] Initialize the integer Value variable (intel#5585) [CI] Fix nightly builds (intel#5584) [SYCL][L0] Fix use of copy-engines in L0 interop queue (intel#5579) Update OpenCL headers tag to dcd5bed (intel#5575) [SYCL] Fix warning for InOrderQueueSyncCheck unit test build (intel#5577) [SYCL][HIP] Remove arch requirement for running lit tests (intel#5253) [SYCL][L0] Fix timestamp calculation (in ns) (intel#5555) [SYCL] Fix sync of host task vs kernel for in-order queue (intel#5551) [sycl-post-link] Add a check for device globals with device_image_scope (intel#5517) [SYCL] Fix SYCL Kernel Body Check (intel#5546) [SYCL] Add support for SYCL 2020 in class group (intel#5447) Fix tests after 1c729d7 Use align attribute for kernel pointer arg alignment Fix tests after 18834dc Mark pointer-typed kernel arguments as ABI aligned [CI] Add experimental Windows build to GitHub Actions nightly (intel#5560) [sycl-post-link][NFC] Address clang-tidy concerns in the sycl-post-link (intel#5552) Fix lit test after changes in Clang Improve backward compatibility for DISubRange ...

aobolensk requested a review from a team as a code owner February 1, 2022 16:17

aobolensk requested a review from againull February 1, 2022 16:17

aobolensk force-pushed the add_get_group_id branch from c384248 to f4921ad Compare February 1, 2022 16:21

aobolensk changed the title ~~Add missing get_group_id function in sycl::group~~ Add support for SYCL 2020 in class group Feb 10, 2022

aobolensk mentioned this pull request Feb 10, 2022

Add tests for local_id in group class on device intel/llvm-test-suite#821

Merged

aobolensk requested a review from Pennycook February 10, 2022 09:46

Pennycook requested changes Feb 10, 2022

View reviewed changes

vladimirlaz suggested changes Feb 11, 2022

View reviewed changes

sycl/include/CL/sycl/group.hpp Outdated Show resolved Hide resolved

sycl/include/CL/sycl/group.hpp Outdated Show resolved Hide resolved

aobolensk added 4 commits February 11, 2022 14:15

Add missing get_group_id function in sycl::group

5668fff

Add rest of SYCL 2020 functions for class group

51d4d77

Extend test for class group

5dd827b

Address review comments

5af7ed6

aobolensk force-pushed the add_get_group_id branch from f72fbd2 to 1fc564a Compare February 11, 2022 11:16

vladimirlaz previously approved these changes Feb 11, 2022

View reviewed changes

aobolensk dismissed vladimirlaz’s stale review via 2258997 February 11, 2022 11:19

aobolensk force-pushed the add_get_group_id branch from 1fc564a to 2258997 Compare February 11, 2022 11:19

bader changed the title ~~Add support for SYCL 2020 in class group~~ [SYCL] Add support for SYCL 2020 in class group Feb 11, 2022

bader requested review from vladimirlaz and Pennycook February 11, 2022 12:56

vladimirlaz previously approved these changes Feb 11, 2022

View reviewed changes

JackAKirk mentioned this pull request Feb 11, 2022

[SYCL][CUDA] Add async_group_copy free function implementation #4907

Closed

Pennycook requested changes Feb 11, 2022

View reviewed changes

aobolensk dismissed vladimirlaz’s stale review via 7fbb9a7 February 11, 2022 15:44

Fix deprecation macro and add deprecation tests

3cba04f

aobolensk force-pushed the add_get_group_id branch from 7fbb9a7 to 3cba04f Compare February 11, 2022 15:49

aobolensk requested a review from Pennycook February 11, 2022 15:53

Pennycook approved these changes Feb 11, 2022

View reviewed changes

againull approved these changes Feb 11, 2022

View reviewed changes

bader reviewed Feb 12, 2022

View reviewed changes

sycl/include/CL/sycl/group.hpp Show resolved Hide resolved

bader merged commit 73d59ce into intel:sycl Feb 14, 2022

[SYCL] Add support for SYCL 2020 in class group #5447

[SYCL] Add support for SYCL 2020 in class group #5447

Uh oh!

Conversation

aobolensk commented Feb 1, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Pennycook commented Feb 1, 2022

Uh oh!

aobolensk commented Feb 10, 2022

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

bader commented Feb 11, 2022

Uh oh!

Pennycook left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

aobolensk commented Feb 11, 2022

Uh oh!

againull commented Feb 11, 2022

Uh oh!

againull commented Feb 11, 2022

Uh oh!

Uh oh!

Uh oh!

aobolensk commented Feb 1, 2022 •

edited

Loading