From b8ea144a15eaec9634db5477c824965fcf7e2d6f Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Mon, 17 May 2021 09:34:46 +0300 Subject: [PATCH 1/7] [SYCL] Implement SYCL2020 sub_group class --- sycl/include/CL/sycl/detail/spirv.hpp | 9 ++ sycl/include/CL/sycl/sub_group.hpp | 128 ++++++++++++++++++++++++++ 2 files changed, 137 insertions(+) create mode 100644 sycl/include/CL/sycl/sub_group.hpp diff --git a/sycl/include/CL/sycl/detail/spirv.hpp b/sycl/include/CL/sycl/detail/spirv.hpp index cc89053783485..961aa8332b90c 100644 --- a/sycl/include/CL/sycl/detail/spirv.hpp +++ b/sycl/include/CL/sycl/detail/spirv.hpp @@ -19,6 +19,7 @@ #ifdef __SYCL_DEVICE_ONLY__ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { +class sub_group; namespace ONEAPI { struct sub_group; } // namespace ONEAPI @@ -35,6 +36,10 @@ template <> struct group_scope<::cl::sycl::ONEAPI::sub_group> { static constexpr __spv::Scope::Flag value = __spv::Scope::Flag::Subgroup; }; +template <> struct group_scope<::cl::sycl::sub_group> { + static constexpr __spv::Scope::Flag value = __spv::Scope::Flag::Subgroup; +}; + // Generic shuffles and broadcasts may require multiple calls to // intrinsics, and should use the fewest broadcasts possible // - Loop over chunks until remaining bytes < chunk size @@ -735,6 +740,10 @@ EnableIfGenericShuffle SubgroupShuffleUp(T x, id<1> local_id) { return Result; } +template bool GroupNonUniformElect() { + return __spirv_GroupNonUniformElect(group_scope::value); +} + } // namespace spirv } // namespace detail } // namespace sycl diff --git a/sycl/include/CL/sycl/sub_group.hpp b/sycl/include/CL/sycl/sub_group.hpp new file mode 100644 index 0000000000000..dc1511e2b9218 --- /dev/null +++ b/sycl/include/CL/sycl/sub_group.hpp @@ -0,0 +1,128 @@ +//==----------- 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 +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +class sub_group { +public: + using id_type = id<1>; + using range_type = range<1>; + using linear_id_type = uint32_t; + static constexpr int dimensions = 1; + static constexpr memory_scope fence_scope = memory_scope::sub_group; + + id_type get_group_id() const { +#ifdef __SYCL_DEVICE_ONLY__ + return __spirv_SubgroupId(); +#else + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif + } + + id_type get_local_id() const { +#ifdef __SYCL_DEVICE_ONLY__ + return __spirv_SubgroupLocalInvocationId(); +#else + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif + } + + range_type get_local_range() const { +#ifdef __SYCL_DEVICE_ONLY__ + return __spirv_SubgroupSize(); +#else + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif + } + + range_type get_group_range() const { +#ifdef __SYCL_DEVICE_ONLY__ + return __spirv_NumSubgroups(); +#else + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif + } + + range_type get_max_local_range() const { +#ifdef __SYCL_DEVICE_ONLY__ + return __spirv_SubgroupMaxSize(); +#else + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif + } + + linear_id_type get_group_linear_id() const { +#ifdef __SYCL_DEVICE_ONLY__ + return static_cast(get_group_id()[0]); +#else + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif + } + + linear_id_type get_local_linear_id() const { +#ifdef __SYCL_DEVICE_ONLY__ + return static_cast(get_local_id()[0]); +#else + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif + } + + linear_id_type get_group_linear_range() const { +#ifdef __SYCL_DEVICE_ONLY__ + return static_cast(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(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__ + using namespace sycl::detail::spirv; + return GroupNonUniformElect>(); +#else + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif + } +}; +} +} From 2d386e9cc6534f736c6c6a6e099e9807a179b919 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Mon, 17 May 2021 09:59:33 +0300 Subject: [PATCH 2/7] Revert some changes to avoid ABI breakage --- sycl/include/CL/sycl.hpp | 1 + sycl/include/CL/sycl/ONEAPI/sub_group.hpp | 40 +++++-- sycl/include/CL/sycl/nd_item.hpp | 4 +- sycl/include/CL/sycl/sub_group.hpp | 121 ++-------------------- 4 files changed, 43 insertions(+), 123 deletions(-) diff --git a/sycl/include/CL/sycl.hpp b/sycl/include/CL/sycl.hpp index 83590e0928235..5f0cad5364997 100644 --- a/sycl/include/CL/sycl.hpp +++ b/sycl/include/CL/sycl.hpp @@ -52,6 +52,7 @@ #include #include #include +#include #include #include #include diff --git a/sycl/include/CL/sycl/ONEAPI/sub_group.hpp b/sycl/include/CL/sycl/ONEAPI/sub_group.hpp index 69d4676f471c9..bf90e925f1b0a 100644 --- a/sycl/include/CL/sycl/ONEAPI/sub_group.hpp +++ b/sycl/include/CL/sycl/ONEAPI/sub_group.hpp @@ -104,6 +104,7 @@ struct sub_group { using range_type = range<1>; using linear_id_type = uint32_t; static constexpr int dimensions = 1; + static constexpr memory_scope fence_scope = memory_scope::sub_group; /* --- common interface members --- */ @@ -699,13 +700,40 @@ struct sub_group { #endif } -protected: - template friend class cl::sycl::nd_item; - friend sub_group this_sub_group(); - sub_group() = default; -}; + linear_id_type get_group_linear_range() const { +#ifdef __SYCL_DEVICE_ONLY__ + return static_cast(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(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__ + using namespace sycl::detail::spirv; + return GroupNonUniformElect>(); +#else + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); +#endif + + protected: + template friend class cl::sycl::nd_item; + friend sub_group this_sub_group(); + sub_group() = default; + }; -inline sub_group this_sub_group() { + inline sub_group this_sub_group() { #ifdef __SYCL_DEVICE_ONLY__ return sub_group(); #else diff --git a/sycl/include/CL/sycl/nd_item.hpp b/sycl/include/CL/sycl/nd_item.hpp index 7e9bdfa4ff354..fae0b7a4202e1 100644 --- a/sycl/include/CL/sycl/nd_item.hpp +++ b/sycl/include/CL/sycl/nd_item.hpp @@ -9,7 +9,6 @@ #pragma once #include -#include #include #include #include @@ -18,6 +17,7 @@ #include #include #include +#include #include #include @@ -67,7 +67,7 @@ template class nd_item { group 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]; diff --git a/sycl/include/CL/sycl/sub_group.hpp b/sycl/include/CL/sycl/sub_group.hpp index dc1511e2b9218..220790fd821f9 100644 --- a/sycl/include/CL/sycl/sub_group.hpp +++ b/sycl/include/CL/sycl/sub_group.hpp @@ -8,121 +8,12 @@ #pragma once -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include +#include __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { -class sub_group { -public: - using id_type = id<1>; - using range_type = range<1>; - using linear_id_type = uint32_t; - static constexpr int dimensions = 1; - static constexpr memory_scope fence_scope = memory_scope::sub_group; - - id_type get_group_id() const { -#ifdef __SYCL_DEVICE_ONLY__ - return __spirv_SubgroupId(); -#else - throw runtime_error("Sub-groups are not supported on host device.", - PI_INVALID_DEVICE); -#endif - } - - id_type get_local_id() const { -#ifdef __SYCL_DEVICE_ONLY__ - return __spirv_SubgroupLocalInvocationId(); -#else - throw runtime_error("Sub-groups are not supported on host device.", - PI_INVALID_DEVICE); -#endif - } - - range_type get_local_range() const { -#ifdef __SYCL_DEVICE_ONLY__ - return __spirv_SubgroupSize(); -#else - throw runtime_error("Sub-groups are not supported on host device.", - PI_INVALID_DEVICE); -#endif - } - - range_type get_group_range() const { -#ifdef __SYCL_DEVICE_ONLY__ - return __spirv_NumSubgroups(); -#else - throw runtime_error("Sub-groups are not supported on host device.", - PI_INVALID_DEVICE); -#endif - } - - range_type get_max_local_range() const { -#ifdef __SYCL_DEVICE_ONLY__ - return __spirv_SubgroupMaxSize(); -#else - throw runtime_error("Sub-groups are not supported on host device.", - PI_INVALID_DEVICE); -#endif - } - - linear_id_type get_group_linear_id() const { -#ifdef __SYCL_DEVICE_ONLY__ - return static_cast(get_group_id()[0]); -#else - throw runtime_error("Sub-groups are not supported on host device.", - PI_INVALID_DEVICE); -#endif - } - - linear_id_type get_local_linear_id() const { -#ifdef __SYCL_DEVICE_ONLY__ - return static_cast(get_local_id()[0]); -#else - throw runtime_error("Sub-groups are not supported on host device.", - PI_INVALID_DEVICE); -#endif - } - - linear_id_type get_group_linear_range() const { -#ifdef __SYCL_DEVICE_ONLY__ - return static_cast(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(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__ - using namespace sycl::detail::spirv; - return GroupNonUniformElect>(); -#else - throw runtime_error("Sub-groups are not supported on host device.", - PI_INVALID_DEVICE); -#endif - } -}; -} -} +using ONEAPI::sub_group; +// TODO move the entire sub_group class implementation to this file once +// breaking changes are allowed. +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) From 85c3c55c84550d4733dda54ce80279b92e8a0c30 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Tue, 18 May 2021 16:10:02 +0300 Subject: [PATCH 3/7] fix build --- sycl/include/CL/__spirv/spirv_ops.hpp | 4 +++- sycl/include/CL/sycl/ONEAPI/sub_group.hpp | 4 ++-- sycl/include/CL/sycl/detail/spirv.hpp | 5 ----- 3 files changed, 5 insertions(+), 8 deletions(-) diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp index 34829ea7892af..9db8768556e28 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -567,6 +567,9 @@ extern SYCL_EXTERNAL void __spirv_ocl_prefetch(const __attribute__((opencl_global)) char *Ptr, size_t NumBytes) noexcept; +__SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT bool + __spirv_GroupNonUniformElect(__spv::Scope::Flag) noexcept; + #else // if !__SYCL_DEVICE_ONLY__ template @@ -606,5 +609,4 @@ __spirv_MemoryBarrier(__spv::Scope Memory, uint32_t Semantics) noexcept; __SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT void __spirv_GroupWaitEvents(__spv::Scope Execution, uint32_t NumEvents, __ocl_event_t *WaitEvents) noexcept; - #endif // !__SYCL_DEVICE_ONLY__ diff --git a/sycl/include/CL/sycl/ONEAPI/sub_group.hpp b/sycl/include/CL/sycl/ONEAPI/sub_group.hpp index bf90e925f1b0a..1c7c46871feb7 100644 --- a/sycl/include/CL/sycl/ONEAPI/sub_group.hpp +++ b/sycl/include/CL/sycl/ONEAPI/sub_group.hpp @@ -720,12 +720,12 @@ struct sub_group { bool leader() const { #ifdef __SYCL_DEVICE_ONLY__ - using namespace sycl::detail::spirv; - return GroupNonUniformElect>(); + return detail::spirv::GroupNonUniformElect(); #else throw runtime_error("Sub-groups are not supported on host device.", PI_INVALID_DEVICE); #endif + } protected: template friend class cl::sycl::nd_item; diff --git a/sycl/include/CL/sycl/detail/spirv.hpp b/sycl/include/CL/sycl/detail/spirv.hpp index 961aa8332b90c..d2f9fdcff8b13 100644 --- a/sycl/include/CL/sycl/detail/spirv.hpp +++ b/sycl/include/CL/sycl/detail/spirv.hpp @@ -19,7 +19,6 @@ #ifdef __SYCL_DEVICE_ONLY__ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { -class sub_group; namespace ONEAPI { struct sub_group; } // namespace ONEAPI @@ -36,10 +35,6 @@ template <> struct group_scope<::cl::sycl::ONEAPI::sub_group> { static constexpr __spv::Scope::Flag value = __spv::Scope::Flag::Subgroup; }; -template <> struct group_scope<::cl::sycl::sub_group> { - static constexpr __spv::Scope::Flag value = __spv::Scope::Flag::Subgroup; -}; - // Generic shuffles and broadcasts may require multiple calls to // intrinsics, and should use the fewest broadcasts possible // - Loop over chunks until remaining bytes < chunk size From b1a8a111108630a12b709e2abc85943308271a3e Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Wed, 19 May 2021 11:17:49 +0300 Subject: [PATCH 4/7] Address some feedback, implement barrier --- sycl/include/CL/__spirv/spirv_ops.hpp | 3 --- sycl/include/CL/sycl/ONEAPI/sub_group.hpp | 5 ++-- sycl/include/CL/sycl/detail/spirv.hpp | 21 +++++++++++++--- sycl/include/CL/sycl/enums.hpp | 30 +++++++++++++++++++++++ sycl/include/CL/sycl/group.hpp | 4 +++ sycl/include/CL/sycl/sub_group.hpp | 17 +++++++++++++ 6 files changed, 71 insertions(+), 9 deletions(-) create mode 100644 sycl/include/CL/sycl/enums.hpp diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp index 9db8768556e28..456e6ab84cf61 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -567,9 +567,6 @@ extern SYCL_EXTERNAL void __spirv_ocl_prefetch(const __attribute__((opencl_global)) char *Ptr, size_t NumBytes) noexcept; -__SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT bool - __spirv_GroupNonUniformElect(__spv::Scope::Flag) noexcept; - #else // if !__SYCL_DEVICE_ONLY__ template diff --git a/sycl/include/CL/sycl/ONEAPI/sub_group.hpp b/sycl/include/CL/sycl/ONEAPI/sub_group.hpp index 1c7c46871feb7..b5403b9af5e71 100644 --- a/sycl/include/CL/sycl/ONEAPI/sub_group.hpp +++ b/sycl/include/CL/sycl/ONEAPI/sub_group.hpp @@ -20,6 +20,7 @@ #include #include #include +#include #include @@ -104,7 +105,7 @@ struct sub_group { using range_type = range<1>; using linear_id_type = uint32_t; static constexpr int dimensions = 1; - static constexpr memory_scope fence_scope = memory_scope::sub_group; + static constexpr sycl::memory_scope fence_scope = sycl::memory_scope::sub_group; /* --- common interface members --- */ @@ -720,7 +721,7 @@ struct sub_group { bool leader() const { #ifdef __SYCL_DEVICE_ONLY__ - return detail::spirv::GroupNonUniformElect(); + return get_local_linear_id() == 0; #else throw runtime_error("Sub-groups are not supported on host device.", PI_INVALID_DEVICE); diff --git a/sycl/include/CL/sycl/detail/spirv.hpp b/sycl/include/CL/sycl/detail/spirv.hpp index d2f9fdcff8b13..89ba6429597fc 100644 --- a/sycl/include/CL/sycl/detail/spirv.hpp +++ b/sycl/include/CL/sycl/detail/spirv.hpp @@ -15,6 +15,7 @@ #include #include #include +#include #ifdef __SYCL_DEVICE_ONLY__ __SYCL_INLINE_NAMESPACE(cl) { @@ -259,6 +260,22 @@ getScope(ONEAPI::memory_scope Scope) { } } +constexpr __spv::Scope::Flag +getScope(memory_scope Scope) { + switch (Scope) { + case memory_scope::work_item: + return __spv::Scope::Invocation; + case memory_scope::sub_group: + return __spv::Scope::Subgroup; + case memory_scope::work_group: + return __spv::Scope::Workgroup; + case memory_scope::device: + return __spv::Scope::Device; + case memory_scope::system: + return __spv::Scope::CrossDevice; + } +} + template inline typename detail::enable_if_t::value, T> AtomicCompareExchange(multi_ptr MPtr, @@ -735,10 +752,6 @@ EnableIfGenericShuffle SubgroupShuffleUp(T x, id<1> local_id) { return Result; } -template bool GroupNonUniformElect() { - return __spirv_GroupNonUniformElect(group_scope::value); -} - } // namespace spirv } // namespace detail } // namespace sycl diff --git a/sycl/include/CL/sycl/enums.hpp b/sycl/include/CL/sycl/enums.hpp new file mode 100644 index 0000000000000..0c14e05ef30d2 --- /dev/null +++ b/sycl/include/CL/sycl/enums.hpp @@ -0,0 +1,30 @@ +//==-------------- 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 + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +enum class memory_scope : unsigned char { + work_item = 0, + sub_group = 1, + work_group = 2, + device = 3, + system = 4 +}; + +#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 +} +} + diff --git a/sycl/include/CL/sycl/group.hpp b/sycl/include/CL/sycl/group.hpp index 51b7ebe6ce76a..671b217704f81 100644 --- a/sycl/include/CL/sycl/group.hpp +++ b/sycl/include/CL/sycl/group.hpp @@ -19,6 +19,7 @@ #include #include #include +#include #include #include @@ -431,5 +432,8 @@ template group this_group() { #endif } +template +void group_barrier(Group G, memory_scope FenceScope = Group::fence_scope); + } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/include/CL/sycl/sub_group.hpp b/sycl/include/CL/sycl/sub_group.hpp index 220790fd821f9..071a9559ea0cf 100644 --- a/sycl/include/CL/sycl/sub_group.hpp +++ b/sycl/include/CL/sycl/sub_group.hpp @@ -8,6 +8,7 @@ #pragma once +#include #include __SYCL_INLINE_NAMESPACE(cl) { @@ -15,5 +16,21 @@ namespace sycl { using ONEAPI::sub_group; // TODO move the entire sub_group class implementation to this file once // breaking changes are allowed. + +template <> inline void group_barrier(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) From 54e68a587e73b9f35d6d23fb5232c3e3266ce347 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Wed, 19 May 2021 11:18:18 +0300 Subject: [PATCH 5/7] clang-format --- sycl/include/CL/sycl/ONEAPI/sub_group.hpp | 17 +++++++++-------- sycl/include/CL/sycl/detail/spirv.hpp | 5 ++--- sycl/include/CL/sycl/enums.hpp | 15 +++++++-------- sycl/include/CL/sycl/group.hpp | 2 +- sycl/include/CL/sycl/sub_group.hpp | 21 +++++++++++---------- 5 files changed, 30 insertions(+), 30 deletions(-) diff --git a/sycl/include/CL/sycl/ONEAPI/sub_group.hpp b/sycl/include/CL/sycl/ONEAPI/sub_group.hpp index b5403b9af5e71..863c58ee2ea05 100644 --- a/sycl/include/CL/sycl/ONEAPI/sub_group.hpp +++ b/sycl/include/CL/sycl/ONEAPI/sub_group.hpp @@ -17,10 +17,10 @@ #include #include #include +#include #include #include #include -#include #include @@ -105,7 +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; + static constexpr sycl::memory_scope fence_scope = + sycl::memory_scope::sub_group; /* --- common interface members --- */ @@ -728,13 +729,13 @@ struct sub_group { #endif } - protected: - template friend class cl::sycl::nd_item; - friend sub_group this_sub_group(); - sub_group() = default; - }; +protected: + template friend class cl::sycl::nd_item; + friend sub_group this_sub_group(); + sub_group() = default; +}; - inline sub_group this_sub_group() { +inline sub_group this_sub_group() { #ifdef __SYCL_DEVICE_ONLY__ return sub_group(); #else diff --git a/sycl/include/CL/sycl/detail/spirv.hpp b/sycl/include/CL/sycl/detail/spirv.hpp index 89ba6429597fc..a593d9968e455 100644 --- a/sycl/include/CL/sycl/detail/spirv.hpp +++ b/sycl/include/CL/sycl/detail/spirv.hpp @@ -14,8 +14,8 @@ #include #include #include -#include #include +#include #ifdef __SYCL_DEVICE_ONLY__ __SYCL_INLINE_NAMESPACE(cl) { @@ -260,8 +260,7 @@ getScope(ONEAPI::memory_scope Scope) { } } -constexpr __spv::Scope::Flag -getScope(memory_scope Scope) { +constexpr __spv::Scope::Flag getScope(memory_scope Scope) { switch (Scope) { case memory_scope::work_item: return __spv::Scope::Invocation; diff --git a/sycl/include/CL/sycl/enums.hpp b/sycl/include/CL/sycl/enums.hpp index 0c14e05ef30d2..16b5f89f64652 100644 --- a/sycl/include/CL/sycl/enums.hpp +++ b/sycl/include/CL/sycl/enums.hpp @@ -11,11 +11,11 @@ __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { enum class memory_scope : unsigned char { - work_item = 0, - sub_group = 1, - work_group = 2, - device = 3, - system = 4 + work_item = 0, + sub_group = 1, + work_group = 2, + device = 3, + system = 4 }; #if __cplusplus >= 201703L @@ -25,6 +25,5 @@ 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) diff --git a/sycl/include/CL/sycl/group.hpp b/sycl/include/CL/sycl/group.hpp index 671b217704f81..a001b4fadfbda 100644 --- a/sycl/include/CL/sycl/group.hpp +++ b/sycl/include/CL/sycl/group.hpp @@ -15,11 +15,11 @@ #include #include #include +#include #include #include #include #include -#include #include #include diff --git a/sycl/include/CL/sycl/sub_group.hpp b/sycl/include/CL/sycl/sub_group.hpp index 071a9559ea0cf..f65cf80419a4f 100644 --- a/sycl/include/CL/sycl/sub_group.hpp +++ b/sycl/include/CL/sycl/sub_group.hpp @@ -8,8 +8,8 @@ #pragma once -#include #include +#include __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { @@ -17,19 +17,20 @@ using ONEAPI::sub_group; // TODO move the entire sub_group class implementation to this file once // breaking changes are allowed. -template <> inline void group_barrier(sub_group Group, memory_scope FenceScope) { +template <> +inline void group_barrier(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); + __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); + throw runtime_error("Sub-groups are not supported on host device.", + PI_INVALID_DEVICE); #endif } } // namespace sycl From 3507d53001c700a8d7da57db38fbecb0ac5bca60 Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Fri, 21 May 2021 18:03:31 +0300 Subject: [PATCH 6/7] address feedback --- sycl/include/CL/__spirv/spirv_ops.hpp | 1 + sycl/include/CL/sycl/ONEAPI/sub_group.hpp | 2 +- sycl/include/CL/sycl/detail/spirv.hpp | 16 ---------------- sycl/include/CL/sycl/group.hpp | 2 +- .../CL/sycl/{enums.hpp => memory_enums.hpp} | 10 +++------- 5 files changed, 6 insertions(+), 25 deletions(-) rename sycl/include/CL/sycl/{enums.hpp => memory_enums.hpp} (87%) diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp index 456e6ab84cf61..34829ea7892af 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -606,4 +606,5 @@ __spirv_MemoryBarrier(__spv::Scope Memory, uint32_t Semantics) noexcept; __SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT void __spirv_GroupWaitEvents(__spv::Scope Execution, uint32_t NumEvents, __ocl_event_t *WaitEvents) noexcept; + #endif // !__SYCL_DEVICE_ONLY__ diff --git a/sycl/include/CL/sycl/ONEAPI/sub_group.hpp b/sycl/include/CL/sycl/ONEAPI/sub_group.hpp index 863c58ee2ea05..e667075181825 100644 --- a/sycl/include/CL/sycl/ONEAPI/sub_group.hpp +++ b/sycl/include/CL/sycl/ONEAPI/sub_group.hpp @@ -17,8 +17,8 @@ #include #include #include -#include #include +#include #include #include diff --git a/sycl/include/CL/sycl/detail/spirv.hpp b/sycl/include/CL/sycl/detail/spirv.hpp index a593d9968e455..cc89053783485 100644 --- a/sycl/include/CL/sycl/detail/spirv.hpp +++ b/sycl/include/CL/sycl/detail/spirv.hpp @@ -14,7 +14,6 @@ #include #include #include -#include #include #ifdef __SYCL_DEVICE_ONLY__ @@ -260,21 +259,6 @@ getScope(ONEAPI::memory_scope Scope) { } } -constexpr __spv::Scope::Flag getScope(memory_scope Scope) { - switch (Scope) { - case memory_scope::work_item: - return __spv::Scope::Invocation; - case memory_scope::sub_group: - return __spv::Scope::Subgroup; - case memory_scope::work_group: - return __spv::Scope::Workgroup; - case memory_scope::device: - return __spv::Scope::Device; - case memory_scope::system: - return __spv::Scope::CrossDevice; - } -} - template inline typename detail::enable_if_t::value, T> AtomicCompareExchange(multi_ptr MPtr, diff --git a/sycl/include/CL/sycl/group.hpp b/sycl/include/CL/sycl/group.hpp index a001b4fadfbda..3ace710a2aea0 100644 --- a/sycl/include/CL/sycl/group.hpp +++ b/sycl/include/CL/sycl/group.hpp @@ -15,9 +15,9 @@ #include #include #include -#include #include #include +#include #include #include #include diff --git a/sycl/include/CL/sycl/enums.hpp b/sycl/include/CL/sycl/memory_enums.hpp similarity index 87% rename from sycl/include/CL/sycl/enums.hpp rename to sycl/include/CL/sycl/memory_enums.hpp index 16b5f89f64652..a4e5123b50b06 100644 --- a/sycl/include/CL/sycl/enums.hpp +++ b/sycl/include/CL/sycl/memory_enums.hpp @@ -8,15 +8,11 @@ #pragma once +#include + __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { -enum class memory_scope : unsigned char { - work_item = 0, - sub_group = 1, - work_group = 2, - device = 3, - system = 4 -}; +using ONEAPI::memory_scope; #if __cplusplus >= 201703L inline constexpr auto memory_scope_work_item = memory_scope::work_item; From 9d3fbdfcc17cc0b68a28805580eef8592d46e30b Mon Sep 17 00:00:00 2001 From: Alexander Batashev Date: Fri, 21 May 2021 18:28:37 +0300 Subject: [PATCH 7/7] Update comment --- sycl/include/CL/sycl/memory_enums.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/memory_enums.hpp b/sycl/include/CL/sycl/memory_enums.hpp index a4e5123b50b06..0e918f92b1fe8 100644 --- a/sycl/include/CL/sycl/memory_enums.hpp +++ b/sycl/include/CL/sycl/memory_enums.hpp @@ -1,4 +1,4 @@ -//==-------------- enums.hpp --- SYCL enums --------------------------------==// +//==-------------- 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.