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 #define SYCL_EXT_ONEAPI_CUDA_ASYNC_BARRIER 1
12 
13 #include <CL/__spirv/spirv_ops.hpp>
14 #include <cstddef>
15 
17 namespace sycl {
18 namespace ext {
19 namespace oneapi {
20 namespace experimental {
21 namespace cuda {
22 
23 class barrier {
24  int64_t state;
25 
26 public:
28 
29  // barriers cannot be moved or copied
30  barrier(const barrier &other) = delete;
31  barrier(barrier &&other) noexcept = delete;
32  barrier &operator=(const barrier &other) = delete;
33  barrier &operator=(barrier &&other) noexcept = delete;
34 
35  void initialize(uint32_t expected_count) {
36 #ifdef __SYCL_DEVICE_ONLY__
37  __clc_BarrierInitialize(&state, expected_count);
38 #else
39  (void)state;
40  (void)expected_count;
41  throw runtime_error("Barrier is not supported on host device.",
42  PI_ERROR_INVALID_DEVICE);
43 #endif
44  }
45 
46  void invalidate() {
47 #ifdef __SYCL_DEVICE_ONLY__
48  __clc_BarrierInvalidate(&state);
49 #else
50  throw runtime_error("Barrier is not supported on host device.",
51  PI_ERROR_INVALID_DEVICE);
52 #endif
53  }
54 
56 #ifdef __SYCL_DEVICE_ONLY__
57  return __clc_BarrierArrive(&state);
58 #else
59  throw runtime_error("Barrier is not supported on host device.",
60  PI_ERROR_INVALID_DEVICE);
61 #endif
62  }
63 
65 #ifdef __SYCL_DEVICE_ONLY__
66  return __clc_BarrierArriveAndDrop(&state);
67 #else
68  throw runtime_error("Barrier is not supported on host device.",
69  PI_ERROR_INVALID_DEVICE);
70 #endif
71  }
72 
74 #ifdef __SYCL_DEVICE_ONLY__
75  return __clc_BarrierArriveNoComplete(&state, count);
76 #else
77  (void)count;
78  throw runtime_error("Barrier is not supported on host device.",
79  PI_ERROR_INVALID_DEVICE);
80 #endif
81  }
82 
84 #ifdef __SYCL_DEVICE_ONLY__
85  return __clc_BarrierArriveAndDropNoComplete(&state, count);
86 #else
87  (void)count;
88  throw runtime_error("Barrier is not supported on host device.",
89  PI_ERROR_INVALID_DEVICE);
90 #endif
91  }
92 
94 #ifdef __SYCL_DEVICE_ONLY__
95  __clc_BarrierCopyAsyncArrive(&state);
96 #else
97  throw runtime_error("Barrier is not supported on host device.",
98  PI_ERROR_INVALID_DEVICE);
99 #endif
100  }
101 
103 #ifdef __SYCL_DEVICE_ONLY__
104  __clc_BarrierCopyAsyncArriveNoInc(&state);
105 #else
106  throw runtime_error("Barrier is not supported on host device.",
107  PI_ERROR_INVALID_DEVICE);
108 #endif
109  }
110 
111  void wait(arrival_token arrival) {
112 #ifdef __SYCL_DEVICE_ONLY__
113  __clc_BarrierWait(&state, arrival);
114 #else
115  (void)arrival;
116  throw runtime_error("Barrier is not supported on host device.",
117  PI_ERROR_INVALID_DEVICE);
118 #endif
119  }
120 
121  bool test_wait(arrival_token arrival) {
122 #ifdef __SYCL_DEVICE_ONLY__
123  return __clc_BarrierTestWait(&state, arrival);
124 #else
125  (void)arrival;
126  throw runtime_error("Barrier is not supported on host device.",
127  PI_ERROR_INVALID_DEVICE);
128 #endif
129  }
130 
132 #ifdef __SYCL_DEVICE_ONLY__
133  __clc_BarrierArriveAndWait(&state);
134 #else
135  throw runtime_error("Barrier is not supported on host device.",
136  PI_ERROR_INVALID_DEVICE);
137 #endif
138  }
139 
140 // On Windows certain headers define macros min/max
141 #pragma push_macro("max")
142 #ifdef max
143 #undef max
144 #endif
145  static constexpr uint64_t max() { return (1 << 20) - 1; }
146 #pragma pop_macro("max")
147 };
148 } // namespace cuda
149 } // namespace experimental
150 } // namespace oneapi
151 } // namespace ext
152 } // namespace sycl
153 } // __SYCL_INLINE_NAMESPACE(cl)
spirv_ops.hpp
cl::sycl::ext::oneapi::experimental::cuda::barrier::arrive_copy_async
void arrive_copy_async()
Definition: barrier.hpp:93
cl::sycl::ext::oneapi::experimental::cuda::barrier::arrival_token
int64_t arrival_token
Definition: barrier.hpp:27
cl::sycl::ext::oneapi::experimental::cuda::barrier::test_wait
bool test_wait(arrival_token arrival)
Definition: barrier.hpp:121
sycl
Definition: invoke_simd.hpp:68
cl::sycl::ext::oneapi::experimental::cuda::barrier::arrive_no_complete
arrival_token arrive_no_complete(int32_t count)
Definition: barrier.hpp:73
cl::sycl::ext::oneapi::experimental::cuda::barrier::initialize
void initialize(uint32_t expected_count)
Definition: barrier.hpp:35
cl::sycl::detail::int64_t
conditional< sizeof(long)==8, long, long long >::type int64_t
Definition: kernel_desc.hpp:34
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::ext::oneapi::experimental::cuda::barrier::invalidate
void invalidate()
Definition: barrier.hpp:46
cl::sycl::ext::oneapi::experimental::cuda::barrier::arrive
arrival_token arrive()
Definition: barrier.hpp:55
cl::sycl::ext::intel::esimd::barrier
__ESIMD_API void barrier()
Generic work-group barrier.
Definition: memory.hpp:919
cl::sycl::ext::oneapi::experimental::cuda::barrier
Definition: barrier.hpp:23
cl::sycl::ext::oneapi::experimental::cuda::barrier::arrive_copy_async_no_inc
void arrive_copy_async_no_inc()
Definition: barrier.hpp:102
cl::sycl::ext::oneapi::experimental::cuda::barrier::max
static constexpr uint64_t max()
Definition: barrier.hpp:145
cl::sycl::ext::oneapi::experimental::cuda::barrier::arrive_and_drop_no_complete
arrival_token arrive_and_drop_no_complete(int32_t count)
Definition: barrier.hpp:83
cl::sycl::ext::oneapi::experimental::cuda::barrier::wait
void wait(arrival_token arrival)
Definition: barrier.hpp:111
cl::sycl::ext::oneapi::experimental::cuda::barrier::arrive_and_drop
arrival_token arrive_and_drop()
Definition: barrier.hpp:64
cl::sycl::ext::oneapi::experimental::cuda::barrier::arrive_and_wait
void arrive_and_wait()
Definition: barrier.hpp:131
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:11