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 <sycl/exception.hpp> // for make_error_code, errc, exception
12 
13 #include <stddef.h> // for size_t
14 #include <stdint.h> // for int32_t
15 
16 namespace sycl {
17 inline namespace _V1 {
18 namespace ext::intel {
19 
20 template <class _name, class _dataT, int32_t _min_capacity = 0> class pipe {
21 public:
22  using value_type = _dataT;
23  static constexpr int32_t min_capacity = _min_capacity;
24  // Non-blocking pipes
25  // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
26  // friendly LLVM IR.
27  static _dataT read(bool &_Success) {
28 #ifdef __SYCL_DEVICE_ONLY__
29  __ocl_RPipeTy<_dataT> _RPipe =
30  __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
31  _dataT TempData;
32  _Success = !static_cast<bool>(
33  __spirv_ReadPipe(_RPipe, &TempData, m_Size, m_Alignment));
34  return TempData;
35 #else
36  (void)_Success;
37  throw sycl::exception(
38  sycl::make_error_code(sycl::errc::feature_not_supported),
39  "Device-side API are not supported on a host device. Please use "
40  "host-side API instead.");
41 #endif // __SYCL_DEVICE_ONLY__
42  }
43 
44  // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
45  // friendly LLVM IR.
46  static void write(const _dataT &_Data, bool &_Success) {
47 #ifdef __SYCL_DEVICE_ONLY__
48  __ocl_WPipeTy<_dataT> _WPipe =
49  __spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
50  _Success = !static_cast<bool>(
51  __spirv_WritePipe(_WPipe, &_Data, m_Size, m_Alignment));
52 #else
53  (void)_Success;
54  (void)_Data;
55  throw sycl::exception(
56  sycl::make_error_code(sycl::errc::feature_not_supported),
57  "Device-side API are not supported on a host device. Please use "
58  "host-side API instead.");
59 #endif // __SYCL_DEVICE_ONLY__
60  }
61 
62  // Blocking pipes
63  // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
64  // friendly LLVM IR.
65  static _dataT read() {
66 #ifdef __SYCL_DEVICE_ONLY__
67  __ocl_RPipeTy<_dataT> _RPipe =
68  __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
69  _dataT TempData;
70  __spirv_ReadPipeBlockingINTEL(_RPipe, &TempData, m_Size, m_Alignment);
71  return TempData;
72 #else
73  throw sycl::exception(
74  sycl::make_error_code(sycl::errc::feature_not_supported),
75  "Device-side API are not supported on a host device. Please use "
76  "host-side API instead..");
77 #endif // __SYCL_DEVICE_ONLY__
78  }
79 
80  // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
81  // friendly LLVM IR.
82  static void write(const _dataT &_Data) {
83 #ifdef __SYCL_DEVICE_ONLY__
84  __ocl_WPipeTy<_dataT> _WPipe =
85  __spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
86  __spirv_WritePipeBlockingINTEL(_WPipe, &_Data, m_Size, m_Alignment);
87 #else
88  (void)_Data;
89  throw sycl::exception(
90  sycl::make_error_code(sycl::errc::feature_not_supported),
91  "Device-side API are not supported on a host device. Please use "
92  "host-side API instead.");
93 #endif // __SYCL_DEVICE_ONLY__
94  }
95 
96 private:
97  static constexpr int32_t m_Size = sizeof(_dataT);
98  static constexpr int32_t m_Alignment = alignof(_dataT);
99 #ifdef __SYCL_DEVICE_ONLY__
100  static constexpr struct ConstantPipeStorage m_Storage = {m_Size, m_Alignment,
101  min_capacity};
102 #endif // __SYCL_DEVICE_ONLY__
103 };
104 
105 // IO pipes that provide interface to connect with hardware peripheral.
106 // Their name aliases are defined in vendor-provided header, below you can see
107 // an example of this header. There are defined aliases to ethernet_read_pipe
108 // and ethernet_write_pipe that users can use in their code to connect with
109 // HW peripheral.
110 /* namespace intelfpga {
111 template <int32_t ID>
112 struct ethernet_pipe_id {
113  static constexpr int32_t id = ID;
114 };
115 
116 template <class _dataT, size_t _min_capacity>
117 using ethernet_read_pipe =
118  kernel_readable_io_pipe<ethernet_pipe_id<0>, _dataT, _min_capacity>;
119 
120 template <class _dataT, size_t _min_capacity>
121 using ethernet_write_pipe =
122  kernel_writeable_io_pipe<ethernet_pipe_id<1>, _dataT, _min_capacity>;
123 } // namespace intelfpga */
124 
125 template <class _name, class _dataT, size_t _min_capacity = 0>
126 class kernel_readable_io_pipe {
127 public:
128  using value_type = _dataT;
129  static constexpr int32_t min_capacity = _min_capacity;
130  // Non-blocking pipes
131  // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
132  // friendly LLVM IR.
133  static _dataT read(bool &_Success) {
134 #ifdef __SYCL_DEVICE_ONLY__
135  __ocl_RPipeTy<_dataT> _RPipe =
136  __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
137  _dataT TempData;
138  _Success = !static_cast<bool>(
139  __spirv_ReadPipe(_RPipe, &TempData, m_Size, m_Alignment));
140  return TempData;
141 #else
142  (void)_Success;
143  throw sycl::exception(
144  sycl::make_error_code(sycl::errc::feature_not_supported),
145  "Device-side API are not supported on a host device. Please use "
146  "host-side API instead.");
147 #endif // __SYCL_DEVICE_ONLY__
148  }
149 
150  // Blocking pipes
151  // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
152  // friendly LLVM IR.
153  static _dataT read() {
154 #ifdef __SYCL_DEVICE_ONLY__
155  __ocl_RPipeTy<_dataT> _RPipe =
156  __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
157  _dataT TempData;
158  __spirv_ReadPipeBlockingINTEL(_RPipe, &TempData, m_Size, m_Alignment);
159  return TempData;
160 #else
161  throw sycl::exception(
162  sycl::make_error_code(sycl::errc::feature_not_supported),
163  "Device-side API are not supported on a host device. Please use "
164  "host-side API instead.");
165 #endif // __SYCL_DEVICE_ONLY__
166  }
167 
168 private:
169  static constexpr int32_t m_Size = sizeof(_dataT);
170  static constexpr int32_t m_Alignment = alignof(_dataT);
171  static constexpr int32_t ID = _name::id;
172 #ifdef __SYCL_DEVICE_ONLY__
173  static constexpr struct ConstantPipeStorage m_Storage
174  __attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, min_capacity};
175 #endif // __SYCL_DEVICE_ONLY__
176 };
177 
178 template <class _name, class _dataT, size_t _min_capacity = 0>
179 class kernel_writeable_io_pipe {
180 public:
181  using value_type = _dataT;
182  static constexpr int32_t min_capacity = _min_capacity;
183  // Non-blocking pipes
184  // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
185  // friendly LLVM IR.
186  static void write(const _dataT &_Data, bool &_Success) {
187 #ifdef __SYCL_DEVICE_ONLY__
188  __ocl_WPipeTy<_dataT> _WPipe =
189  __spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
190  _Success = !static_cast<bool>(
191  __spirv_WritePipe(_WPipe, &_Data, m_Size, m_Alignment));
192 #else
193  (void)_Data;
194  (void)_Success;
195  throw sycl::exception(
196  sycl::make_error_code(sycl::errc::feature_not_supported),
197  "Device-side API are not supported on a host device. Please use "
198  "host-side API instead.");
199 #endif // __SYCL_DEVICE_ONLY__
200  }
201 
202  // Blocking pipes
203  // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
204  // friendly LLVM IR.
205  static void write(const _dataT &_Data) {
206 #ifdef __SYCL_DEVICE_ONLY__
207  __ocl_WPipeTy<_dataT> _WPipe =
208  __spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
209  __spirv_WritePipeBlockingINTEL(_WPipe, &_Data, m_Size, m_Alignment);
210 #else
211  (void)_Data;
212  throw sycl::exception(
213  sycl::make_error_code(sycl::errc::feature_not_supported),
214  "Device-side API are not supported on a host device. Please use "
215  "host-side API instead.");
216 #endif // __SYCL_DEVICE_ONLY__
217  }
218 
219 private:
220  static constexpr int32_t m_Size = sizeof(_dataT);
221  static constexpr int32_t m_Alignment = alignof(_dataT);
222  static constexpr int32_t ID = _name::id;
223 #ifdef __SYCL_DEVICE_ONLY__
224  static constexpr struct ConstantPipeStorage m_Storage
225  __attribute__((io_pipe_id(ID))) = {m_Size, m_Alignment, min_capacity};
226 #endif // __SYCL_DEVICE_ONLY__
227 };
228 
229 } // namespace ext::intel
230 
231 } // namespace _V1
232 } // namespace sycl
sycl::_V1::ext::intel::kernel_writeable_io_pipe::write
static void write(const _dataT &_Data, bool &_Success)
Definition: pipes.hpp:186
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:94
sycl::_V1::ext::intel::kernel_readable_io_pipe
Definition: pipes.hpp:126
sycl::_V1::ext::intel::pipe::min_capacity
static constexpr int32_t min_capacity
Definition: pipes.hpp:23
sycl::_V1::ext::intel::kernel_writeable_io_pipe
Definition: pipes.hpp:179
sycl
Definition: access.hpp:18
sycl::_V1::ext::intel::pipe::read
static _dataT read()
Definition: pipes.hpp:65
sycl::_V1::ext::intel::kernel_readable_io_pipe::read
static _dataT read(bool &_Success)
Definition: pipes.hpp:133
sycl::_V1::ext::intel::pipe
Definition: pipes.hpp:20
sycl::_V1::ext::intel::pipe::value_type
_dataT value_type
Definition: pipes.hpp:22
sycl::_V1::exception
Definition: exception.hpp:68
sycl::_V1::ext::intel::kernel_writeable_io_pipe::value_type
_dataT value_type
Definition: pipes.hpp:181
sycl::_V1::ext::intel::pipe::write
static void write(const _dataT &_Data, bool &_Success)
Definition: pipes.hpp:46
sycl::_V1::ext::intel::kernel_readable_io_pipe::value_type
_dataT value_type
Definition: pipes.hpp:128
sycl::_V1::ext::intel::kernel_writeable_io_pipe::write
static void write(const _dataT &_Data)
Definition: pipes.hpp:205
exception.hpp
sycl::_V1::ext::intel::pipe::write
static void write(const _dataT &_Data)
Definition: pipes.hpp:82
sycl::_V1::ext::intel::kernel_readable_io_pipe::min_capacity
static constexpr int32_t min_capacity
Definition: pipes.hpp:129
sycl::_V1::ext::intel::kernel_writeable_io_pipe::min_capacity
static constexpr int32_t min_capacity
Definition: pipes.hpp:182
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:460
sycl::_V1::ext::intel::kernel_readable_io_pipe::read
static _dataT read()
Definition: pipes.hpp:153
sycl::_V1::ext::intel::pipe::read
static _dataT read(bool &_Success)
Definition: pipes.hpp:27