DPC++ Runtime
Runtime libraries for oneAPI DPC++
barrier.hpp
Go to the documentation of this file.
1 //==--- barrier.hpp - SYCL_ONEAPI_BARRIER ---------------------------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #pragma once
10 
11 #include <CL/__spirv/spirv_ops.hpp> // for __clc_BarrierInitialize
12 #include <sycl/exception.hpp>
13 
14 #include <stdint.h> // for int32_t, int64_t, uint32_t, uint64_t
15 
16 #define SYCL_EXT_ONEAPI_CUDA_ASYNC_BARRIER 1
17 
18 namespace sycl {
19 inline namespace _V1 {
20 namespace ext {
21 namespace oneapi {
22 namespace experimental {
23 namespace cuda {
24 
25 class barrier {
26  int64_t state;
27 
28 public:
30 
31  // barriers cannot be moved or copied
32  barrier(const barrier &other) = delete;
33  barrier(barrier &&other) noexcept = delete;
34  barrier &operator=(const barrier &other) = delete;
35  barrier &operator=(barrier &&other) noexcept = delete;
36 
37  void initialize(uint32_t expected_count) {
38 #ifdef __SYCL_DEVICE_ONLY__
39  __clc_BarrierInitialize(&state, expected_count);
40 #else
41  (void)state;
42  (void)expected_count;
44  "Barrier is not supported on host.");
45 #endif
46  }
47 
48  void invalidate() {
49 #ifdef __SYCL_DEVICE_ONLY__
50  __clc_BarrierInvalidate(&state);
51 #else
53  "Barrier is not supported on host.");
54 #endif
55  }
56 
58 #ifdef __SYCL_DEVICE_ONLY__
59  return __clc_BarrierArrive(&state);
60 #else
62  "Barrier is not supported on host.");
63 #endif
64  }
65 
67 #ifdef __SYCL_DEVICE_ONLY__
68  return __clc_BarrierArriveAndDrop(&state);
69 #else
71  "Barrier is not supported on host.");
72 #endif
73  }
74 
76 #ifdef __SYCL_DEVICE_ONLY__
77  return __clc_BarrierArriveNoComplete(&state, count);
78 #else
79  (void)count;
81  "Barrier is not supported on host.");
82 #endif
83  }
84 
86 #ifdef __SYCL_DEVICE_ONLY__
87  return __clc_BarrierArriveAndDropNoComplete(&state, count);
88 #else
89  (void)count;
91  "Barrier is not supported on host.");
92 #endif
93  }
94 
96 #ifdef __SYCL_DEVICE_ONLY__
97  __clc_BarrierCopyAsyncArrive(&state);
98 #else
100  "Barrier is not supported on host.");
101 #endif
102  }
103 
105 #ifdef __SYCL_DEVICE_ONLY__
106  __clc_BarrierCopyAsyncArriveNoInc(&state);
107 #else
109  "Barrier is not supported on host.");
110 #endif
111  }
112 
113  void wait(arrival_token arrival) {
114 #ifdef __SYCL_DEVICE_ONLY__
115  __clc_BarrierWait(&state, arrival);
116 #else
117  (void)arrival;
119  "Barrier is not supported on host.");
120 #endif
121  }
122 
123  bool test_wait(arrival_token arrival) {
124 #ifdef __SYCL_DEVICE_ONLY__
125  return __clc_BarrierTestWait(&state, arrival);
126 #else
127  (void)arrival;
129  "Barrier is not supported on host.");
130 #endif
131  }
132 
134 #ifdef __SYCL_DEVICE_ONLY__
135  __clc_BarrierArriveAndWait(&state);
136 #else
138  "Barrier is not supported on host.");
139 #endif
140  }
141 
142 // On Windows certain headers define macros min/max
143 #pragma push_macro("max")
144 #ifdef max
145 #undef max
146 #endif
147  static constexpr uint64_t max() { return (1 << 20) - 1; }
148 #pragma pop_macro("max")
149 };
150 } // namespace cuda
151 } // namespace experimental
152 } // namespace oneapi
153 } // namespace ext
154 } // namespace _V1
155 } // namespace sycl
barrier & operator=(const barrier &other)=delete
arrival_token arrive_and_drop_no_complete(int32_t count)
Definition: barrier.hpp:85
void initialize(uint32_t expected_count)
Definition: barrier.hpp:37
barrier & operator=(barrier &&other) noexcept=delete
arrival_token arrive_no_complete(int32_t count)
Definition: barrier.hpp:75
barrier(barrier &&other) noexcept=delete
conditional< sizeof(long)==8, long, long long >::type int64_t
Definition: kernel_desc.hpp:51
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:65
Definition: access.hpp:18
_Abi const simd< _Tp, _Abi > & noexcept
Definition: simd.hpp:1324