17 namespace ext::intel {
19 template <
class _name,
class _dataT,
int32_t _min_capacity = 0>
class pipe {
22 static constexpr int32_t min_capacity = _min_capacity;
26 static _dataT
read(
bool &_Success) {
27 #ifdef __SYCL_DEVICE_ONLY__
28 __ocl_RPipeTy<_dataT> _RPipe =
29 __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
31 _Success = !
static_cast<bool>(
32 __spirv_ReadPipe(_RPipe, &TempData, m_Size, m_Alignment));
36 throw sycl::exception(
38 "Pipes are not supported on a host device.");
39 #endif // __SYCL_DEVICE_ONLY__
44 static void write(
const _dataT &_Data,
bool &_Success) {
45 #ifdef __SYCL_DEVICE_ONLY__
46 __ocl_WPipeTy<_dataT> _WPipe =
47 __spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
48 _Success = !
static_cast<bool>(
49 __spirv_WritePipe(_WPipe, &_Data, m_Size, m_Alignment));
53 throw sycl::exception(
55 "Pipes are not supported on a host device.");
56 #endif // __SYCL_DEVICE_ONLY__
63 #ifdef __SYCL_DEVICE_ONLY__
64 __ocl_RPipeTy<_dataT> _RPipe =
65 __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
67 __spirv_ReadPipeBlockingINTEL(_RPipe, &TempData, m_Size, m_Alignment);
70 throw sycl::exception(
72 "Pipes are not supported on a host device.");
73 #endif // __SYCL_DEVICE_ONLY__
78 static void write(
const _dataT &_Data) {
79 #ifdef __SYCL_DEVICE_ONLY__
80 __ocl_WPipeTy<_dataT> _WPipe =
81 __spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
82 __spirv_WritePipeBlockingINTEL(_WPipe, &_Data, m_Size, m_Alignment);
85 throw sycl::exception(
87 "Pipes are not supported on a host device.");
88 #endif // __SYCL_DEVICE_ONLY__
92 static constexpr int32_t m_Size =
sizeof(_dataT);
93 static constexpr int32_t m_Alignment =
alignof(_dataT);
94 #ifdef __SYCL_DEVICE_ONLY__
95 static constexpr
struct ConstantPipeStorage m_Storage = {m_Size, m_Alignment,
97 #endif // __SYCL_DEVICE_ONLY__
120 template <
class _name,
class _dataT,
size_t _min_capacity = 0>
124 static constexpr int32_t min_capacity = _min_capacity;
128 static _dataT
read(
bool &_Success) {
129 #ifdef __SYCL_DEVICE_ONLY__
130 __ocl_RPipeTy<_dataT> _RPipe =
131 __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
133 _Success = !
static_cast<bool>(
134 __spirv_ReadPipe(_RPipe, &TempData, m_Size, m_Alignment));
138 throw sycl::exception(
140 "Pipes are not supported on a host device.");
141 #endif // __SYCL_DEVICE_ONLY__
148 #ifdef __SYCL_DEVICE_ONLY__
149 __ocl_RPipeTy<_dataT> _RPipe =
150 __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
152 __spirv_ReadPipeBlockingINTEL(_RPipe, &TempData, m_Size, m_Alignment);
155 throw sycl::exception(
157 "Pipes are not supported on a host device.");
158 #endif // __SYCL_DEVICE_ONLY__
162 static constexpr int32_t m_Size =
sizeof(_dataT);
163 static constexpr int32_t m_Alignment =
alignof(_dataT);
164 static constexpr int32_t ID = _name::id;
165 #ifdef __SYCL_DEVICE_ONLY__
166 static constexpr
struct ConstantPipeStorage m_Storage
167 __attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, min_capacity};
168 #endif // __SYCL_DEVICE_ONLY__
171 template <
class _name,
class _dataT,
size_t _min_capacity = 0>
175 static constexpr int32_t min_capacity = _min_capacity;
179 static void write(
const _dataT &_Data,
bool &_Success) {
180 #ifdef __SYCL_DEVICE_ONLY__
181 __ocl_WPipeTy<_dataT> _WPipe =
182 __spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
183 _Success = !
static_cast<bool>(
184 __spirv_WritePipe(_WPipe, &_Data, m_Size, m_Alignment));
188 throw sycl::exception(
190 "Pipes are not supported on a host device.");
191 #endif // __SYCL_DEVICE_ONLY__
197 static void write(
const _dataT &_Data) {
198 #ifdef __SYCL_DEVICE_ONLY__
199 __ocl_WPipeTy<_dataT> _WPipe =
200 __spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
201 __spirv_WritePipeBlockingINTEL(_WPipe, &_Data, m_Size, m_Alignment);
204 throw sycl::exception(
206 "Pipes are not supported on a host device.");
207 #endif // __SYCL_DEVICE_ONLY__
211 static constexpr int32_t m_Size =
sizeof(_dataT);
212 static constexpr int32_t m_Alignment =
alignof(_dataT);
213 static constexpr int32_t ID = _name::id;
214 #ifdef __SYCL_DEVICE_ONLY__
215 static constexpr
struct ConstantPipeStorage m_Storage
216 __attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, min_capacity};
217 #endif // __SYCL_DEVICE_ONLY__