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/detail/export.hpp> // for __SYCL_EX...
12 #include <sycl/device.hpp> // for device
13 #include <sycl/event.hpp> // for event
14 #include <sycl/exception.hpp> // for make_erro...
15 #include <sycl/ext/intel/experimental/pipe_properties.hpp> // for protocol_...
16 #include <sycl/ext/oneapi/properties/properties.hpp> // for ValueOrDe...
17 #include <sycl/handler.hpp> // for handler
18 #include <sycl/info/info_desc.hpp> // for event_com...
19 #include <sycl/memory_enums.hpp> // for memory_order
20 #include <sycl/queue.hpp> // for queue
21 
22 #ifdef __SYCL_DEVICE_ONLY__
25 #endif
26 
27 #ifdef XPTI_ENABLE_INSTRUMENTATION
28 #include <xpti/xpti_data_types.h>
29 #include <xpti/xpti_trace_framework.hpp>
30 #endif
31 
32 #include <stdint.h> // for int32_t
33 #include <string> // for string
34 #include <tuple> // for _Swallow_...
35 
36 namespace sycl {
37 inline namespace _V1 {
38 namespace ext {
39 namespace intel {
40 namespace experimental {
41 
42 // A helper templateless base class to get the host_pipe name.
43 class pipe_base {
44 
45 protected:
48 
49  __SYCL_EXPORT static std::string get_pipe_name(const void *HostPipePtr);
50 };
51 
52 template <class _name, class _dataT, int32_t _min_capacity = 0,
53  class _propertiesT = decltype(oneapi::experimental::properties{}),
54  class = void>
55 class pipe : public pipe_base {
56 public:
57  struct
58 #ifdef __SYCL_DEVICE_ONLY__
59  [[__sycl_detail__::add_ir_attributes_global_variable(
60  "sycl-host-pipe", "sycl-host-pipe-size", nullptr,
61  sizeof(_dataT))]] [[__sycl_detail__::sycl_type(host_pipe)]]
62 #endif // __SYCL_DEVICE_ONLY___
64 #ifdef __SYCL_DEVICE_ONLY__
65  : ConstantPipeStorage
66 #endif // __SYCL_DEVICE_ONLY___
67  {
68  int32_t _ReadyLatency;
69  int32_t _BitsPerSymbol;
70  bool _UsesValid;
73  };
74 
75  // Non-blocking pipes
76 
77  // Host API
78  static _dataT read(queue &Q, bool &Success,
80  // Order is currently unused.
81  std::ignore = Order;
82 
83  const device Dev = Q.get_device();
84  bool IsPipeSupported =
85  Dev.has_extension("cl_intel_program_scope_host_pipe");
86  if (!IsPipeSupported) {
87  return _dataT();
88  }
89  _dataT Data;
90  void *DataPtr = &Data;
91  const void *HostPipePtr = &m_Storage;
92  const std::string PipeName = pipe_base::get_pipe_name(HostPipePtr);
93 
94  event E = Q.submit([=](handler &CGH) {
95  CGH.ext_intel_read_host_pipe(PipeName, DataPtr,
96  sizeof(_dataT) /* non-blocking */);
97  });
98  E.wait();
99  if (E.get_info<sycl::info::event::command_execution_status>() ==
100  sycl::info::event_command_status::complete) {
101  Success = true;
102  return *(_dataT *)DataPtr;
103  } else {
104  Success = false;
105  return _dataT();
106  }
107  }
108 
109  static void write(queue &Q, const _dataT &Data, bool &Success,
111  // Order is currently unused.
112  std::ignore = Order;
113 
114  const device Dev = Q.get_device();
115  bool IsPipeSupported =
116  Dev.has_extension("cl_intel_program_scope_host_pipe");
117  if (!IsPipeSupported) {
118  return;
119  }
120 
121  const void *HostPipePtr = &m_Storage;
122  const std::string PipeName = pipe_base::get_pipe_name(HostPipePtr);
123  void *DataPtr = const_cast<_dataT *>(&Data);
124 
125  event E = Q.submit([=](handler &CGH) {
126  CGH.ext_intel_write_host_pipe(PipeName, DataPtr,
127  sizeof(_dataT) /* non-blocking */);
128  });
129  E.wait();
130  Success = E.get_info<sycl::info::event::command_execution_status>() ==
131  sycl::info::event_command_status::complete;
132  }
133 
134  // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
135  // friendly LLVM IR.
136  template <typename _functionPropertiesT>
137  static _dataT read(bool &Success, _functionPropertiesT) {
138 #ifdef __SYCL_DEVICE_ONLY__
139  // Get latency control properties
140  using _latency_anchor_id_prop = typename detail::GetOrDefaultValT<
141  _functionPropertiesT, latency_anchor_id_key,
143  using _latency_constraint_prop = typename detail::GetOrDefaultValT<
144  _functionPropertiesT, latency_constraint_key,
146 
147  // Get latency control property values
148  static constexpr int32_t _anchor_id = _latency_anchor_id_prop::value;
149  static constexpr int32_t _target_anchor = _latency_constraint_prop::target;
150  static constexpr latency_control_type _control_type =
151  _latency_constraint_prop::type;
152  static constexpr int32_t _relative_cycle = _latency_constraint_prop::cycle;
153 
154  int32_t _control_type_code = 0; // latency_control_type::none is default
155  if constexpr (_control_type == latency_control_type::exact) {
156  _control_type_code = 1;
157  } else if constexpr (_control_type == latency_control_type::max) {
158  _control_type_code = 2;
159  } else if constexpr (_control_type == latency_control_type::min) {
160  _control_type_code = 3;
161  }
162 
163  __ocl_RPipeTy<_dataT> _RPipe =
164  __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
165  _dataT TempData;
166  Success = !static_cast<bool>(__latency_control_nb_read_wrapper(
167  _RPipe, &TempData, _anchor_id, _target_anchor, _control_type_code,
168  _relative_cycle));
169  return TempData;
170 #else
171  (void)Success;
172  throw sycl::exception(
173  sycl::make_error_code(sycl::errc::feature_not_supported),
174  "Device-side API are not supported on a host device. Please use "
175  "host-side API instead.");
176 #endif // __SYCL_DEVICE_ONLY__
177  }
178 
179  static _dataT read(bool &Success) {
180  return read(Success, oneapi::experimental::properties{});
181  }
182 
183  // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
184  // friendly LLVM IR.
185  template <typename _functionPropertiesT>
186  static void write(const _dataT &Data, bool &Success, _functionPropertiesT) {
187 #ifdef __SYCL_DEVICE_ONLY__
188  // Get latency control properties
189  using _latency_anchor_id_prop = typename detail::GetOrDefaultValT<
190  _functionPropertiesT, latency_anchor_id_key,
192  using _latency_constraint_prop = typename detail::GetOrDefaultValT<
193  _functionPropertiesT, latency_constraint_key,
195 
196  // Get latency control property values
197  static constexpr int32_t _anchor_id = _latency_anchor_id_prop::value;
198  static constexpr int32_t _target_anchor = _latency_constraint_prop::target;
199  static constexpr latency_control_type _control_type =
200  _latency_constraint_prop::type;
201  static constexpr int32_t _relative_cycle = _latency_constraint_prop::cycle;
202 
203  int32_t _control_type_code = 0; // latency_control_type::none is default
204  if constexpr (_control_type == latency_control_type::exact) {
205  _control_type_code = 1;
206  } else if constexpr (_control_type == latency_control_type::max) {
207  _control_type_code = 2;
208  } else if constexpr (_control_type == latency_control_type::min) {
209  _control_type_code = 3;
210  }
211 
212  __ocl_WPipeTy<_dataT> _WPipe =
213  __spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
214  Success = !static_cast<bool>(__latency_control_nb_write_wrapper(
215  _WPipe, &Data, _anchor_id, _target_anchor, _control_type_code,
216  _relative_cycle));
217 #else
218  (void)Success;
219  (void)Data;
220  throw sycl::exception(
221  sycl::make_error_code(sycl::errc::feature_not_supported),
222  "Device-side API are not supported on a host device. Please use "
223  "host-side API instead.");
224 #endif // __SYCL_DEVICE_ONLY__
225  }
226 
227  static void write(const _dataT &Data, bool &Success) {
228  write(Data, Success, oneapi::experimental::properties{});
229  }
230 
231  static const void *get_host_ptr() { return &m_Storage; }
232 
233  // Blocking pipes
234 
235  // Host API
236  static _dataT read(queue &Q, memory_order Order = memory_order::seq_cst) {
237  // Order is currently unused.
238  std::ignore = Order;
239 
240  const device Dev = Q.get_device();
241  bool IsPipeSupported =
242  Dev.has_extension("cl_intel_program_scope_host_pipe");
243  if (!IsPipeSupported) {
244  return _dataT();
245  }
246  _dataT Data;
247  void *DataPtr = &Data;
248  const void *HostPipePtr = &m_Storage;
249  const std::string PipeName = pipe_base::get_pipe_name(HostPipePtr);
250  event E = Q.submit([=](handler &CGH) {
251  CGH.ext_intel_read_host_pipe(PipeName, DataPtr, sizeof(_dataT),
252  true /*blocking*/);
253  });
254  E.wait();
255  return *(_dataT *)DataPtr;
256  }
257 
258  static void write(queue &Q, const _dataT &Data,
260  // Order is currently unused.
261  std::ignore = Order;
262 
263  const device Dev = Q.get_device();
264  bool IsPipeSupported =
265  Dev.has_extension("cl_intel_program_scope_host_pipe");
266  if (!IsPipeSupported) {
267  return;
268  }
269  const void *HostPipePtr = &m_Storage;
270  const std::string PipeName = pipe_base::get_pipe_name(HostPipePtr);
271  void *DataPtr = const_cast<_dataT *>(&Data);
272  event E = Q.submit([=](handler &CGH) {
273  CGH.ext_intel_write_host_pipe(PipeName, DataPtr, sizeof(_dataT),
274  true /*blocking */);
275  });
276  E.wait();
277  }
278 
279  // Reading from pipe is lowered to SPIR-V instruction OpReadPipe via SPIR-V
280  // friendly LLVM IR.
281  template <typename _functionPropertiesT>
282  static _dataT read(_functionPropertiesT) {
283 #ifdef __SYCL_DEVICE_ONLY__
284  // Get latency control properties
285  using _latency_anchor_id_prop = typename detail::GetOrDefaultValT<
286  _functionPropertiesT, latency_anchor_id_key,
288  using _latency_constraint_prop = typename detail::GetOrDefaultValT<
289  _functionPropertiesT, latency_constraint_key,
291 
292  // Get latency control property values
293  static constexpr int32_t _anchor_id = _latency_anchor_id_prop::value;
294  static constexpr int32_t _target_anchor = _latency_constraint_prop::target;
295  static constexpr latency_control_type _control_type =
296  _latency_constraint_prop::type;
297  static constexpr int32_t _relative_cycle = _latency_constraint_prop::cycle;
298 
299  int32_t _control_type_code = 0; // latency_control_type::none is default
300  if constexpr (_control_type == latency_control_type::exact) {
301  _control_type_code = 1;
302  } else if constexpr (_control_type == latency_control_type::max) {
303  _control_type_code = 2;
304  } else if constexpr (_control_type == latency_control_type::min) {
305  _control_type_code = 3;
306  }
307 
308  __ocl_RPipeTy<_dataT> _RPipe =
309  __spirv_CreatePipeFromPipeStorage_read<_dataT>(&m_Storage);
310  _dataT TempData;
311  __latency_control_bl_read_wrapper(_RPipe, &TempData, _anchor_id,
312  _target_anchor, _control_type_code,
313  _relative_cycle);
314  return TempData;
315 #else
316  throw sycl::exception(
317  sycl::make_error_code(sycl::errc::feature_not_supported),
318  "Device-side API are not supported on a host device. Please use "
319  "host-side API instead.");
320 #endif // __SYCL_DEVICE_ONLY__
321  }
322 
323  static _dataT read() { return read(oneapi::experimental::properties{}); }
324 
325  // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
326  // friendly LLVM IR.
327  template <typename _functionPropertiesT>
328  static void write(const _dataT &Data, _functionPropertiesT) {
329 #ifdef __SYCL_DEVICE_ONLY__
330  // Get latency control properties
331  using _latency_anchor_id_prop = typename detail::GetOrDefaultValT<
332  _functionPropertiesT, latency_anchor_id_key,
334  using _latency_constraint_prop = typename detail::GetOrDefaultValT<
335  _functionPropertiesT, latency_constraint_key,
337 
338  // Get latency control property values
339  static constexpr int32_t _anchor_id = _latency_anchor_id_prop::value;
340  static constexpr int32_t _target_anchor = _latency_constraint_prop::target;
341  static constexpr latency_control_type _control_type =
342  _latency_constraint_prop::type;
343  static constexpr int32_t _relative_cycle = _latency_constraint_prop::cycle;
344 
345  int32_t _control_type_code = 0; // latency_control_type::none is default
346  if constexpr (_control_type == latency_control_type::exact) {
347  _control_type_code = 1;
348  } else if constexpr (_control_type == latency_control_type::max) {
349  _control_type_code = 2;
350  } else if constexpr (_control_type == latency_control_type::min) {
351  _control_type_code = 3;
352  }
353 
354  __ocl_WPipeTy<_dataT> _WPipe =
355  __spirv_CreatePipeFromPipeStorage_write<_dataT>(&m_Storage);
356  __latency_control_bl_write_wrapper(_WPipe, &Data, _anchor_id,
357  _target_anchor, _control_type_code,
358  _relative_cycle);
359 #else
360  (void)Data;
361  throw sycl::exception(
362  sycl::make_error_code(sycl::errc::feature_not_supported),
363  "Device-side API are not supported on a host device. Please use "
364  "host-side API instead.");
365 #endif // __SYCL_DEVICE_ONLY__
366  }
367 
368  static void write(const _dataT &Data) {
369  write(Data, oneapi::experimental::properties{});
370  }
371 
372 private:
373  static constexpr int32_t m_Size = sizeof(_dataT);
374  static constexpr int32_t m_Alignment = alignof(_dataT);
375  static constexpr int32_t m_Capacity = _min_capacity;
376 
377  static constexpr int32_t m_ready_latency =
379  _propertiesT, ready_latency_key>::template get<int32_t>(0);
380  static constexpr int32_t m_bits_per_symbol =
382  _propertiesT, bits_per_symbol_key>::template get<int32_t>(8);
383  static constexpr bool m_uses_valid =
385  _propertiesT, uses_valid_key>::template get<bool>(true);
386  static constexpr bool m_first_symbol_in_high_order_bits =
388  _propertiesT,
389  first_symbol_in_high_order_bits_key>::template get<int32_t>(0);
390  static constexpr protocol_name m_protocol = oneapi::experimental::detail::
391  ValueOrDefault<_propertiesT, protocol_key>::template get<protocol_name>(
393 
394 public:
395  static constexpr struct ConstantPipeStorageExp m_Storage = {
396 #ifdef __SYCL_DEVICE_ONLY__
397  {m_Size, m_Alignment, m_Capacity},
398 #endif // __SYCL_DEVICE_ONLY___
399  m_ready_latency,
400  m_bits_per_symbol,
401  m_uses_valid,
402  m_first_symbol_in_high_order_bits,
403  m_protocol};
404 
405 #ifdef __SYCL_DEVICE_ONLY__
406 private:
407  // FPGA BE will recognize this function and extract its arguments.
408  // TODO: Pass latency control parameters via the __spirv_* builtin when ready.
409  template <typename _T>
410  static int32_t __latency_control_nb_read_wrapper(
411  __ocl_RPipeTy<_T> Pipe, _T *Data, int32_t /* AnchorID */,
412  int32_t /* TargetAnchor */, int32_t /* Type */, int32_t /* Cycle */) {
413  return __spirv_ReadPipe(Pipe, Data, m_Size, m_Alignment);
414  }
415 
416  // FPGA BE will recognize this function and extract its arguments.
417  // TODO: Pass latency control parameters via the __spirv_* builtin when ready.
418  template <typename _T>
419  static int32_t __latency_control_nb_write_wrapper(
420  __ocl_WPipeTy<_T> Pipe, const _T *Data, int32_t /* AnchorID */,
421  int32_t /* TargetAnchor */, int32_t /* Type */, int32_t /* Cycle */) {
422  return __spirv_WritePipe(Pipe, Data, m_Size, m_Alignment);
423  }
424 
425  // FPGA BE will recognize this function and extract its arguments.
426  // TODO: Pass latency control parameters via the __spirv_* builtin when ready.
427  template <typename _T>
428  static void __latency_control_bl_read_wrapper(
429  __ocl_RPipeTy<_T> Pipe, _T *Data, int32_t /* AnchorID */,
430  int32_t /* TargetAnchor */, int32_t /* Type */, int32_t /* Cycle */) {
431  return __spirv_ReadPipeBlockingINTEL(Pipe, Data, m_Size, m_Alignment);
432  }
433 
434  // FPGA BE will recognize this function and extract its arguments.
435  // TODO: Pass latency control parameters via the __spirv_* builtin when ready.
436  template <typename _T>
437  static void __latency_control_bl_write_wrapper(
438  __ocl_WPipeTy<_T> Pipe, const _T *Data, int32_t /* AnchorID*/,
439  int32_t /* TargetAnchor */, int32_t /* Type */, int32_t /* Cycle */) {
440  return __spirv_WritePipeBlockingINTEL(Pipe, Data, m_Size, m_Alignment);
441  }
442 #endif // __SYCL_DEVICE_ONLY__
443 };
444 
445 } // namespace experimental
446 } // namespace intel
447 } // namespace ext
448 } // namespace _V1
449 } // namespace sycl
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Definition: device.hpp:59
bool has_extension(const std::string &extension_name) const
Check SYCL extension support by device.
Definition: device.cpp:133
An event object can be used to synchronize memory transfers, enqueues of kernels and signaling barrie...
Definition: event.hpp:44
void wait()
Wait for the event.
Definition: event.cpp:47
detail::is_event_info_desc< Param >::return_type get_info() const
Queries this SYCL event for information.
Definition: event.cpp:77
static std::string get_pipe_name(const void *HostPipePtr)
Definition: pipes.cpp:19
static _dataT read(queue &Q, bool &Success, memory_order Order=memory_order::seq_cst)
Definition: pipes.hpp:78
static void write(const _dataT &Data, bool &Success, _functionPropertiesT)
Definition: pipes.hpp:186
static void write(const _dataT &Data)
Definition: pipes.hpp:368
static _dataT read(bool &Success)
Definition: pipes.hpp:179
static void write(queue &Q, const _dataT &Data, bool &Success, memory_order Order=memory_order::seq_cst)
Definition: pipes.hpp:109
static void write(const _dataT &Data, bool &Success)
Definition: pipes.hpp:227
static _dataT read(bool &Success, _functionPropertiesT)
Definition: pipes.hpp:137
static void write(const _dataT &Data, _functionPropertiesT)
Definition: pipes.hpp:328
static const void * get_host_ptr()
Definition: pipes.hpp:231
static _dataT read(_functionPropertiesT)
Definition: pipes.hpp:282
static void write(queue &Q, const _dataT &Data, memory_order Order=memory_order::seq_cst)
Definition: pipes.hpp:258
static _dataT read(queue &Q, memory_order Order=memory_order::seq_cst)
Definition: pipes.hpp:236
Command group handler class.
Definition: handler.hpp:453
Encapsulates a single SYCL queue which schedules kernels on a SYCL device.
Definition: queue.hpp:119
device get_device() const
Definition: queue.cpp:77
std::enable_if_t< std::is_invocable_r_v< void, T, handler & >, event > submit(T CGF, const detail::code_location &CodeLoc=detail::code_location::current())
Submits a command group function object to the queue, in order to be scheduled for execution on the d...
Definition: queue.hpp:347
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
Definition: exception.cpp:94
Definition: access.hpp:18