Skip to content

[SYCL] Implement SYCL2020 sub_group class #3765

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

Merged
merged 7 commits into from
May 21, 2021
Merged

[SYCL] Implement SYCL2020 sub_group class #3765

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
b8ea144
[SYCL] Implement SYCL2020 sub_group class
May 17, 2021
2d386e9
Revert some changes to avoid ABI breakage
May 17, 2021
85c3c55
fix build
May 18, 2021
b1a8a11
Address some feedback, implement barrier
May 19, 2021
54e68a5
clang-format
May 19, 2021
3507d53
address feedback
May 21, 2021
9d3fbdf
Update comment
May 21, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions sycl/include/CL/sycl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -52,6 +52,7 @@
#include <CL/sycl/sampler.hpp>
#include <CL/sycl/specialization_id.hpp>
#include <CL/sycl/stream.hpp>
#include <CL/sycl/sub_group.hpp>
#include <CL/sycl/types.hpp>
#include <CL/sycl/usm.hpp>
#include <CL/sycl/version.hpp>
30 changes: 30 additions & 0 deletions sycl/include/CL/sycl/ONEAPI/sub_group.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,6 +18,7 @@
#include <CL/sycl/detail/spirv.hpp>
#include <CL/sycl/detail/type_traits.hpp>
#include <CL/sycl/id.hpp>
#include <CL/sycl/memory_enums.hpp>
#include <CL/sycl/range.hpp>
#include <CL/sycl/types.hpp>

Expand Down Expand Up @@ -104,6 +105,8 @@ struct sub_group {
using range_type = range<1>;
using linear_id_type = uint32_t;
static constexpr int dimensions = 1;
static constexpr sycl::memory_scope fence_scope =
sycl::memory_scope::sub_group;

/* --- common interface members --- */

Expand Down Expand Up @@ -699,6 +702,33 @@ struct sub_group {
#endif
}

linear_id_type get_group_linear_range() const {
#ifdef __SYCL_DEVICE_ONLY__
return static_cast<linear_id_type>(get_group_range()[0]);
#else
throw runtime_error("Sub-groups are not supported on host device.",
PI_INVALID_DEVICE);
#endif
}

linear_id_type get_local_linear_range() const {
#ifdef __SYCL_DEVICE_ONLY__
return static_cast<linear_id_type>(get_local_range()[0]);
#else
throw runtime_error("Sub-groups are not supported on host device.",
PI_INVALID_DEVICE);
#endif
}

bool leader() const {
#ifdef __SYCL_DEVICE_ONLY__
return get_local_linear_id() == 0;
#else
throw runtime_error("Sub-groups are not supported on host device.",
PI_INVALID_DEVICE);
#endif
}

protected:
template <int dimensions> friend class cl::sycl::nd_item;
friend sub_group this_sub_group();
Expand Down
4 changes: 4 additions & 0 deletions sycl/include/CL/sycl/group.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,6 +17,7 @@
#include <CL/sycl/device_event.hpp>
#include <CL/sycl/h_item.hpp>
#include <CL/sycl/id.hpp>
#include <CL/sycl/memory_enums.hpp>
#include <CL/sycl/pointers.hpp>
#include <CL/sycl/range.hpp>
#include <stdexcept>
Expand Down Expand Up @@ -431,5 +432,8 @@ template <int Dims> group<Dims> this_group() {
#endif
}

template <typename Group>
void group_barrier(Group G, memory_scope FenceScope = Group::fence_scope);

} // namespace sycl
} // __SYCL_INLINE_NAMESPACE(cl)
25 changes: 25 additions & 0 deletions sycl/include/CL/sycl/memory_enums.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,25 @@
//==-------------- memory_enums.hpp --- SYCL enums -------------------------==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//

#pragma once

#include <CL/sycl/ONEAPI/atomic_enums.hpp>

__SYCL_INLINE_NAMESPACE(cl) {
namespace sycl {
using ONEAPI::memory_scope;

#if __cplusplus >= 201703L
inline constexpr auto memory_scope_work_item = memory_scope::work_item;
inline constexpr auto memory_scope_sub_group = memory_scope::sub_group;
inline constexpr auto memory_scope_work_group = memory_scope::work_group;
inline constexpr auto memory_scope_device = memory_scope::device;
inline constexpr auto memory_scope_system = memory_scope::system;
#endif
} // namespace sycl
} // __SYCL_INLINE_NAMESPACE(cl)
4 changes: 2 additions & 2 deletions sycl/include/CL/sycl/nd_item.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,6 @@
#pragma once

#include <CL/__spirv/spirv_ops.hpp>
#include <CL/sycl/ONEAPI/sub_group.hpp>
#include <CL/sycl/access/access.hpp>
#include <CL/sycl/detail/defines.hpp>
#include <CL/sycl/detail/helpers.hpp>
Expand All @@ -18,6 +17,7 @@
#include <CL/sycl/item.hpp>
#include <CL/sycl/nd_range.hpp>
#include <CL/sycl/range.hpp>
#include <CL/sycl/sub_group.hpp>

#include <cstddef>
#include <stdexcept>
Expand Down Expand Up @@ -67,7 +67,7 @@ template <int dimensions = 1> class nd_item {

group<dimensions> get_group() const { return Group; }

ONEAPI::sub_group get_sub_group() const { return ONEAPI::sub_group(); }
sub_group get_sub_group() const { return sub_group(); }

size_t __SYCL_ALWAYS_INLINE get_group(int dimension) const {
size_t Size = Group[dimension];
Expand Down
37 changes: 37 additions & 0 deletions sycl/include/CL/sycl/sub_group.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,37 @@
//==----------- sub_group.hpp --- SYCL sub-group ---------------------------==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//

#pragma once

#include <CL/sycl/ONEAPI/sub_group.hpp>
#include <CL/sycl/group.hpp>

__SYCL_INLINE_NAMESPACE(cl) {
namespace sycl {
using ONEAPI::sub_group;
// TODO move the entire sub_group class implementation to this file once
Copy link
Contributor

Choose a reason for hiding this comment

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

I was planning on moving the sub_group stuff out of ONEAPI namespace for the SYCL2020 sub-group support where I'm renaming the routines. Does the existing ONEAPI/sub_group API have to be left intact?

Copy link
Contributor

Choose a reason for hiding this comment

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

Apparently yes. No problem.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

That's right. nd_item returns ONEAPI::sub_group. To comply with the spec we'd need to either change the return type (which is a breaking change), or leave it as is and just add new members.

Copy link
Contributor

Choose a reason for hiding this comment

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

Apologies for commenting on a merged PR, but @Pennycook and I have a concern about adding an alias from ONEAPI::sub_group into the sycl:: namespace. There are two problems we see:

  1. We will probably want to deprecate the ONEAPI version of the this API at some point by adding the [[deprecated]] attribute. However, it's not clear how we can do this if sycl::sub_group is an alias. By adding [[deprecated]] to the ONEAPI version we would also deprecate the sycl version.

  2. The ONEAPI::sub_group contains many member functions that are not part of the standard sycl::sub_group type. By adding an alias, we silently add all of these non-standard member functions into the standard type. This is a violation of the extension policy.

Point (2) is a big concern to us because we do not want applications to mistakenly use non-standard APIs. The extension convention was created, in part, to make it obvious when an application uses an extended API. For example, consider the following code:

sycl::sub_group sg = it.get_sub_group(); // create a standard SYCL 2020 sub-group
auto x = sg.load(ptr); // This uses a non-standard member function from sycl::ONEAPI::sub_group, but it's not obvious

Application code like that above may not realize that sg.load() is non-standard, and they will be surprised when it doesn't work on another SYCL implementation. We had customer feedback during the design of the extension policy that customers want extensions to be named in way that makes it clear that they are extensions.

I think the right way to add support for SYCL 2020 sub_group would be to add a new class sycl::sub_group that is not an alias to the ONEAPI version. The new class would either:

  • Implement only the standard SYCL 2020 member functions, or
  • Implement the standard member functions and add extended ones using the extension naming convention (which means they start with the prefix ext_oneapi_).

Can we reconsider the way this is implemented?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

@gmlueck I see one big problem with that code. Consider I'm a big fan of always auto. If I write code like this:

auto sg = it.get_sub_group();

What kind of behavior should I expect? If I previously used ONEAPI extension, and we update the return type, my code is broken. If I write new code and we do not update the return type, it will also be broken (although in a slightly less painful way). Either way, the fact that nd_item::get_sub_group() must return different types breaks some use cases.

Application code like that above may not realize that sg.load() is non-standard

We can add more warnings for this case.

Copy link
Contributor

Choose a reason for hiding this comment

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

I acknowledge that whatever we do here we'll probably break somebody's code. Unfortunately I think it's unavoidable, and it's a side-effect of the sub_group class predating the extension mechanism.

If we'd had the forethought to use it.ext_oneapi_get_sub_group() we'd have no problems. Instead, we're in a situation where SYCL 1.2.1 + SYCL_INTEL_sub_group defines a function with the same name as SYCL 2020.

Can we guard things based on SYCL language version, maybe? If somebody compiles with SYCL 1.2.1 + DPC++ extensions, it.get_sub_group() could continue to return sycl::ONEAPI::sub_group as before. But if somebody opts-in to compiling with SYCL 2020, the only reasonable expectation is that it.get_sub_group() returns a sycl::sub_group.

Copy link
Contributor

Choose a reason for hiding this comment

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

Yes, I agree that things would be better if we had named this it.ext_oneapi_get_sub_group(). Actually, this is a good illustration of the problems that can occur when we don't follow the extension mechanism!

I like @Pennycook's suggestion about defining it.get_sub_group() differently depending on whether the compiler is in SYCL 1.2.1 mode or SYCL 2020 mode. In addition, I'd suggest adding a new member function it.ext_intel_get_sub_group(). That member function would always return the ONEAPI::sub_group type. This provides a quick fix for any applications that use the old sub-group extension; they can just change their call from it.get_sub_group() to it.ext_intel_get_sub_group().

// breaking changes are allowed.

template <>
inline void group_barrier<sub_group>(sub_group Group, memory_scope FenceScope) {
(void)Group;
(void)FenceScope;
#ifdef __SYCL_DEVICE_ONLY__
__spirv_ControlBarrier(__spv::Scope::Subgroup,
detail::spirv::getScope(FenceScope),
__spv::MemorySemanticsMask::AcquireRelease |
__spv::MemorySemanticsMask::SubgroupMemory |
__spv::MemorySemanticsMask::WorkgroupMemory |
__spv::MemorySemanticsMask::CrossWorkgroupMemory);
#else
throw runtime_error("Sub-groups are not supported on host device.",
PI_INVALID_DEVICE);
#endif
}
} // namespace sycl
} // __SYCL_INLINE_NAMESPACE(cl)