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..e667075181825 100644 --- a/sycl/include/CL/sycl/ONEAPI/sub_group.hpp +++ b/sycl/include/CL/sycl/ONEAPI/sub_group.hpp @@ -18,6 +18,7 @@ #include #include #include +#include #include #include @@ -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 --- */ @@ -699,6 +702,33 @@ struct sub_group { #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__ + return get_local_linear_id() == 0; +#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(); diff --git a/sycl/include/CL/sycl/group.hpp b/sycl/include/CL/sycl/group.hpp index 51b7ebe6ce76a..3ace710a2aea0 100644 --- a/sycl/include/CL/sycl/group.hpp +++ b/sycl/include/CL/sycl/group.hpp @@ -17,6 +17,7 @@ #include #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/memory_enums.hpp b/sycl/include/CL/sycl/memory_enums.hpp new file mode 100644 index 0000000000000..0e918f92b1fe8 --- /dev/null +++ b/sycl/include/CL/sycl/memory_enums.hpp @@ -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 + +__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) 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 new file mode 100644 index 0000000000000..f65cf80419a4f --- /dev/null +++ b/sycl/include/CL/sycl/sub_group.hpp @@ -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 +#include + +__SYCL_INLINE_NAMESPACE(cl) { +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)