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