DPC++ Runtime
Runtime libraries for oneAPI DPC++
group_barrier.hpp
Go to the documentation of this file.
1 
2 //==------------------------- group_barrier.hpp ----------------------------==//
3 //
4 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
5 // See https://llvm.org/LICENSE.txt for license information.
6 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
7 //
8 //===----------------------------------------------------------------------===//
9 
10 #pragma once
11 
12 #include <CL/__spirv/spirv_ops.hpp>
15 #include <sycl/detail/spirv.hpp>
17 #include <sycl/group.hpp>
18 #include <sycl/sub_group.hpp>
19 
20 namespace sycl {
22 
23 namespace detail {
24 template <typename G> struct group_barrier_scope {};
25 template <> struct group_barrier_scope<sycl::sub_group> {
26  constexpr static auto Scope = __spv::Scope::Subgroup;
27 };
28 template <int D> struct group_barrier_scope<sycl::group<D>> {
29  constexpr static auto Scope = __spv::Scope::Workgroup;
30 };
31 } // namespace detail
32 
33 template <typename Group>
34 typename std::enable_if<is_group_v<Group>>::type
35 group_barrier(Group, memory_scope FenceScope = Group::fence_scope) {
36  (void)FenceScope;
37 #ifdef __SYCL_DEVICE_ONLY__
38  // Per SYCL spec, group_barrier must perform both control barrier and memory
39  // fence operations. All work-items execute a release fence prior to
40  // barrier and acquire fence afterwards. The rest of semantics flags specify
41  // which type of memory this behavior is applied to.
43  sycl::detail::spirv::getScope(FenceScope),
48 #else
49  throw sycl::runtime_error("Barriers are not supported on host device",
50  PI_ERROR_INVALID_DEVICE);
51 #endif
52 }
53 
54 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
55 } // namespace sycl
#define __SYCL_INLINE_VER_NAMESPACE(X)
std::enable_if< is_group_v< Group > >::type group_barrier(Group, memory_scope FenceScope=Group::fence_scope)
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
__SYCL_CONVERGENT__ SYCL_EXTERNAL void __spirv_ControlBarrier(__spv::Scope Execution, __spv::Scope Memory, uint32_t Semantics) noexcept
Definition: spirv_ops.cpp:26