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 <sycl/detail/spirv.hpp> // for ControlBarrier
13 #include <sycl/detail/type_traits.hpp> // for is_group
14 #include <sycl/exception.hpp> // for make_error_code, errc, exception
15 #include <sycl/memory_enums.hpp> // for memory_scope
16 
17 #include <type_traits> // for enable_if_t
18 
19 namespace sycl {
20 inline namespace _V1 {
21 
22 template <typename Group>
23 std::enable_if_t<is_group_v<Group>>
25  // Per SYCL spec, group_barrier must perform both control barrier and memory
26  // fence operations. All work-items execute a release fence prior to
27  // barrier and acquire fence afterwards.
28 #ifdef __SYCL_DEVICE_ONLY__
29  detail::spirv::ControlBarrier(G, FenceScope, memory_order::seq_cst);
30 #else
31  (void)G;
32  (void)FenceScope;
34  "Barriers are not supported on host");
35 #endif
36 }
37 
38 } // namespace _V1
39 } // namespace sycl
fence_scope
The scope that fence() operation should apply to.
Definition: common.hpp:345
void group_barrier(ext::oneapi::experimental::root_group< dimensions > G, memory_scope FenceScope=decltype(G)::fence_scope)
Definition: root_group.hpp:100
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:64
Definition: access.hpp:18