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