DPC++ Runtime
Runtime libraries for oneAPI DPC++
pipes.hpp
Go to the documentation of this file.
1 //==---------------- pipes.hpp - SYCL pipes ------------*- C++ -*-----------==//
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>
12 #include <CL/__spirv/spirv_types.hpp>
13 #include <sycl/stl.hpp>
14 
15 namespace sycl {
16 __SYCL_INLINE_VER_NAMESPACE(_V1) {
17 namespace ext::intel {
18 
19 template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
20 public:
21  using value_type = _dataT;
22  static constexpr int32_t min_capacity = _min_capacity;
23  // Non-blocking pipes
24  // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
25  // friendly LLVM IR.
26  static _dataT read(bool &_Success) {
27 #ifdef __SYCL_DEVICE_ONLY__
28  __ocl_RPipeTy<_dataT> _RPipe =
29  __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
30  _dataT TempData;
31  _Success = !static_cast<bool>(
32  __spirv_ReadPipe(_RPipe, &TempData, m_Size, m_Alignment));
33  return TempData;
34 #else
35  (void)_Success;
36  throw sycl::exception(
37  sycl::make_error_code(sycl::errc::feature_not_supported),
38  "Pipes are not supported on a host device.");
39 #endif // __SYCL_DEVICE_ONLY__
40  }
41 
42  // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
43  // friendly LLVM IR.
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));
50 #else
51  (void)_Success;
52  (void)_Data;
53  throw sycl::exception(
54  sycl::make_error_code(sycl::errc::feature_not_supported),
55  "Pipes are not supported on a host device.");
56 #endif // __SYCL_DEVICE_ONLY__
57  }
58 
59  // Blocking pipes
60  // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
61  // friendly LLVM IR.
62  static _dataT read() {
63 #ifdef __SYCL_DEVICE_ONLY__
64  __ocl_RPipeTy<_dataT> _RPipe =
65  __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
66  _dataT TempData;
67  __spirv_ReadPipeBlockingINTEL(_RPipe, &TempData, m_Size, m_Alignment);
68  return TempData;
69 #else
70  throw sycl::exception(
71  sycl::make_error_code(sycl::errc::feature_not_supported),
72  "Pipes are not supported on a host device.");
73 #endif // __SYCL_DEVICE_ONLY__
74  }
75 
76  // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
77  // friendly LLVM IR.
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);
83 #else
84  (void)_Data;
85  throw sycl::exception(
86  sycl::make_error_code(sycl::errc::feature_not_supported),
87  "Pipes are not supported on a host device.");
88 #endif // __SYCL_DEVICE_ONLY__
89  }
90 
91 private:
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,
96  min_capacity};
97 #endif // __SYCL_DEVICE_ONLY__
98 };
99 
100 // IO pipes that provide interface to connect with hardware peripheral.
101 // Their name aliases are defined in vendor-provided header, below you can see
102 // an example of this header. There are defined aliases to ethernet_read_pipe
103 // and ethernet_write_pipe that users can use in their code to connect with
104 // HW peripheral.
105 /* namespace intelfpga {
106 template <int32_t ID>
107 struct ethernet_pipe_id {
108  static constexpr int32_t id = ID;
109 };
110 
111 template <class _dataT, size_t _min_capacity>
112 using ethernet_read_pipe =
113  kernel_readable_io_pipe<ethernet_pipe_id<0>, _dataT, _min_capacity>;
114 
115 template <class _dataT, size_t _min_capacity>
116 using ethernet_write_pipe =
117  kernel_writeable_io_pipe<ethernet_pipe_id<1>, _dataT, _min_capacity>;
118 } // namespace intelfpga */
119 
120 template <class _name, class _dataT, size_t _min_capacity = 0>
121 class kernel_readable_io_pipe {
122 public:
123  using value_type = _dataT;
124  static constexpr int32_t min_capacity = _min_capacity;
125  // Non-blocking pipes
126  // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
127  // friendly LLVM IR.
128  static _dataT read(bool &_Success) {
129 #ifdef __SYCL_DEVICE_ONLY__
130  __ocl_RPipeTy<_dataT> _RPipe =
131  __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
132  _dataT TempData;
133  _Success = !static_cast<bool>(
134  __spirv_ReadPipe(_RPipe, &TempData, m_Size, m_Alignment));
135  return TempData;
136 #else
137  (void)_Success;
138  throw sycl::exception(
139  sycl::make_error_code(sycl::errc::feature_not_supported),
140  "Pipes are not supported on a host device.");
141 #endif // __SYCL_DEVICE_ONLY__
142  }
143 
144  // Blocking pipes
145  // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
146  // friendly LLVM IR.
147  static _dataT read() {
148 #ifdef __SYCL_DEVICE_ONLY__
149  __ocl_RPipeTy<_dataT> _RPipe =
150  __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
151  _dataT TempData;
152  __spirv_ReadPipeBlockingINTEL(_RPipe, &TempData, m_Size, m_Alignment);
153  return TempData;
154 #else
155  throw sycl::exception(
156  sycl::make_error_code(sycl::errc::feature_not_supported),
157  "Pipes are not supported on a host device.");
158 #endif // __SYCL_DEVICE_ONLY__
159  }
160 
161 private:
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__
169 };
170 
171 template <class _name, class _dataT, size_t _min_capacity = 0>
172 class kernel_writeable_io_pipe {
173 public:
174  using value_type = _dataT;
175  static constexpr int32_t min_capacity = _min_capacity;
176  // Non-blocking pipes
177  // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
178  // friendly LLVM IR.
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));
185 #else
186  (void)_Data;
187  (void)_Success;
188  throw sycl::exception(
189  sycl::make_error_code(sycl::errc::feature_not_supported),
190  "Pipes are not supported on a host device.");
191 #endif // __SYCL_DEVICE_ONLY__
192  }
193 
194  // Blocking pipes
195  // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
196  // friendly LLVM IR.
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);
202 #else
203  (void)_Data;
204  throw sycl::exception(
205  sycl::make_error_code(sycl::errc::feature_not_supported),
206  "Pipes are not supported on a host device.");
207 #endif // __SYCL_DEVICE_ONLY__
208  }
209 
210 private:
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__
218 };
219 
220 } // namespace ext::intel
221 
222 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
223 } // namespace sycl
spirv_ops.hpp
sycl::_V1::ext::intel::kernel_writeable_io_pipe::write
static void write(const _dataT &_Data, bool &_Success)
Definition: pipes.hpp:179
sycl::_V1::make_error_code
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:92
stl.hpp
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:11
sycl::_V1::ext::intel::kernel_readable_io_pipe
Definition: pipes.hpp:121
sycl::_V1::ext::intel::kernel_writeable_io_pipe
Definition: pipes.hpp:172
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
sycl::_V1::ext::intel::pipe::read
static _dataT read()
Definition: pipes.hpp:62
sycl::_V1::ext::intel::kernel_readable_io_pipe::read
static _dataT read(bool &_Success)
Definition: pipes.hpp:128
sycl::_V1::ext::intel::pipe
Definition: pipes.hpp:19
sycl::_V1::ext::intel::pipe::value_type
_dataT value_type
Definition: pipes.hpp:21
sycl::_V1::ext::intel::kernel_writeable_io_pipe::value_type
_dataT value_type
Definition: pipes.hpp:174
sycl::_V1::ext::intel::pipe::write
static void write(const _dataT &_Data, bool &_Success)
Definition: pipes.hpp:44
sycl::_V1::ext::intel::kernel_readable_io_pipe::value_type
_dataT value_type
Definition: pipes.hpp:123
sycl::_V1::ext::intel::kernel_writeable_io_pipe::write
static void write(const _dataT &_Data)
Definition: pipes.hpp:197
sycl::_V1::ext::intel::pipe::write
static void write(const _dataT &_Data)
Definition: pipes.hpp:78
sycl::_V1::ext::oneapi::experimental::__attribute__
__attribute__((always_inline)) auto invoke_simd(sycl
The invoke_simd free function invokes a SIMD function using all work-items in a sub_group.
Definition: invoke_simd.hpp:357
sycl::_V1::ext::intel::kernel_readable_io_pipe::read
static _dataT read()
Definition: pipes.hpp:147
sycl::_V1::ext::intel::pipe::read
static _dataT read(bool &_Success)
Definition: pipes.hpp:26
spirv_types.hpp