Skip to content

Commit a982a12

Browse files
rho180steffenlarsenbader
authored
[SYCL][Doc] Host pipe design (#5850)
Co-authored-by: Steffen Larsen <[email protected]> Co-authored-by: Alexey Bader <[email protected]>
1 parent ded5db5 commit a982a12

File tree

2 files changed

+192
-0
lines changed

2 files changed

+192
-0
lines changed

‎sycl/doc/design/HostPipes.md

Lines changed: 191 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,191 @@
1+
# Implementation design for "Host Pipes"
2+
3+
This document describes the implementation design for the host pipes section
4+
of the DPC++ extension [SYCL_INTEL_data_flow_pipes][1]. Pipes are a FIFO construct
5+
that provide links between elements of a design that are accessed through read
6+
and write application programming interfaces (APIs), without the notion of a
7+
memory address/pointer to elements within the FIFO. A host pipe is a pipe that
8+
links a device kernel with a host program.
9+
10+
[1]: <../extensions/supported/sycl_ext_intel_dataflow_pipes.asciidoc>
11+
12+
## Requirements
13+
14+
The extension specification document referenced above contains the full set of
15+
requirements for this feature, but a requirement that is particularly
16+
relevant to the design, and similar in nature to one raised in the [device_global][2]
17+
design is called out here.
18+
19+
This issue relates to the mechanism for integrating host and device code.
20+
Like device global variables, host pipes are referenced in both
21+
host and device code, so they require some mechanism to correlate the pipe
22+
instance in device code with the pipe instance in host code. We will use
23+
a similar mechanism as the device global implementation that creates a map
24+
database in the integration headers and footers.
25+
26+
[2]: <DeviceGlobal.md>
27+
28+
## Design
29+
30+
### Changes to DPC++ headers
31+
32+
#### Attributes attached to the class
33+
34+
The `pipe` class uses a new C++ attribute `[[__sycl_detail__::host_pipe]]` on the
35+
`pipe::__pipeType` type to identify the `static const __pipeType` member `__pipe`
36+
as a host pipe. Similar to `[[__sycl_detail__::device_global]]`, this will inform
37+
the front end to generate a `sycl-unique-id` for each `__pipe`. The `pipe` class
38+
also introduces the global variable attribute `sycl-host-pipe` attribute to inform the sycl-post-link tool
39+
to generate the SPIR-V decoration `HostAccessINTEL` for each `__pipe` using the
40+
`sycl-unique-id` generated.
41+
42+
As these attributes are only needed for the device compiler, the `#ifdef __SYCL_DEVICE_ONLY__`
43+
allows the customer to use another host compiler, even if it does not recognize these attributes.
44+
Also note that these attributes are all in the `__sycl_detail__` namespace, so
45+
they are considered implementation details of DPC++. We do not intend to
46+
support them as general attributes that customer code can use.
47+
48+
```
49+
template <typename name, typename dataT, typename propertiesT = properties<>>
50+
class pipe {/*...*/};
51+
52+
// Partial specialization to make propertiesT visible as a parameter pack
53+
// of properties.
54+
template <typename Name, typename DataT, typename ...Props>
55+
class pipe
56+
{
57+
struct
58+
#ifdef __SYCL_DEVICE_ONLY__
59+
[[__sycl_detail__::add_ir_attributes_global_variable(
60+
"sycl-host-pipe",
61+
Props::meta_name...,
62+
nullptr,
63+
Props::meta_value...
64+
)]]
65+
[[__sycl_detail__::host_pipe]]
66+
[[__sycl_detail__::global_variable_allowed]] // may not be needed
67+
#endif
68+
__pipeType { const char __p; };
69+
70+
static constexpr __pipeType __pipe = {0};
71+
...
72+
};
73+
```
74+
The `[[__sycl_detail__::add_ir_attributes_global_variable()]]` attribute is
75+
described more fully by the [compile-time properties][3] design
76+
document. This attribute is also used for other classes that have properties,
77+
so it is not specific to the `pipe` class.
78+
79+
The address of `static const __pipeType` member `__pipe` will be used to identify the pipe
80+
in host code, and provide one half of the host-to-device mapping of the pipe
81+
(see the section on __New content in the integration header and footer__ below).
82+
83+
[3]: <CompileTimeProperties.md>
84+
85+
### Changes to the DPC++ front-end
86+
87+
There are several changes to the device compiler front-end:
88+
89+
* The front-end adds a new LLVM IR attribute `sycl-unique-id` to the definition
90+
of each `pipe` variable, which provides a unique string identifier
91+
for each.
92+
93+
* The front-end generates new content in both the integration header and the
94+
integration footer, which is described in more detail below.
95+
96+
#### New content in the integration header and footer
97+
98+
New content in the integration header and footer provides a mapping from the
99+
host address of each pipe variable to the unique string for that
100+
variable. To illustrate, consider a translation unit that defines two
101+
`pipe` classes:
102+
103+
```
104+
#include <sycl/sycl.hpp>
105+
106+
class some_pipe;
107+
namespace inner {
108+
class some_other_pipe;
109+
} // namespace inner
110+
...
111+
pipe<class some_pipe, ...>::write(...); // a usage of pipe<class some_pipe, ...>
112+
...
113+
pipe<class some_other_pipe, ...>::read(...); // a usage of pipe<class some_other_pipe, ...>
114+
...
115+
116+
```
117+
118+
The corresponding integration header defines a namespace scope variable of type
119+
`__sycl_host_pipe_registration` (referred to below as the __host pipe registrar__)
120+
whose sole purpose is to run its constructor before the application's main() function:
121+
122+
```
123+
namespace sycl::detail {
124+
namespace {
125+
126+
class __sycl_host_pipe_registration {
127+
public:
128+
__sycl_host_pipe_registration() noexcept;
129+
};
130+
__sycl_host_pipe_registration __sycl_host_pipe_registrar;
131+
132+
} // namespace (unnamed)
133+
} // namespace sycl::detail
134+
```
135+
136+
The integration footer contains the definition of the constructor, which calls
137+
a function in the DPC++ runtime with the following information for each host
138+
pipe that is used in the translation unit:
139+
140+
* The (host) address of the static member variable `__pipe`.
141+
* The variable's string from the `sycl-unique-id` attribute.
142+
143+
```
144+
namespace sycl::detail {
145+
namespace {
146+
147+
__sycl_host_pipe_registration::__sycl_host_pipe_registration() noexcept {
148+
host_pipe_map::add(&pipe<some_pipe, ...>::__pipe,
149+
/* same string returned from __builtin_sycl_unique_pipe_id(pipe<some_pipe, ...>::__pipe) */);
150+
host_pipe_map::add(&inner::pipe<some_other_pipe>::__pipe,
151+
/* same string returned from __builtin_sycl_unique_pipe_id(pipe<some_other_pipe, ...>::__pipe) */);
152+
}
153+
154+
} // namespace (unnamed)
155+
} // namespace sycl::detail
156+
```
157+
158+
Further details on adherence to C++ rules for unconstructed objects can be found
159+
in the [device_global][2] design.
160+
161+
Unique pipe ids will be generated by the same method as [device_global][2] uses to generate `sycl-unique-id`s.
162+
163+
### Changes to the DPC++ runtime
164+
165+
Several changes are needed to the DPC++ runtime
166+
167+
* As we noted above, the front-end generates new content in the integration
168+
footer which calls the function `sycl::detail::host_pipe_map::add()`.
169+
The runtime defines this function and maintains information about all the
170+
host pipe variables in the application. This information includes:
171+
172+
- The host address of the variable.
173+
- The string which uniquely identifies the variable.
174+
175+
* The runtime implements the `read` and `write` functions of the pipe
176+
class. These will use this [host pipe API][4]. These functions will
177+
need to retrieve the mapping added to the __host pipe registrar__
178+
for the pipe being read or written to, and pass it to the corresponding
179+
underlying OpenCL API call
180+
181+
[4]: https://github.com/intel-sandbox/ip-authoring-specs/blob/MJ_ChangeDocs4/Pipe/Spec/cl_intel_host_pipe_symbol.asciidoc
182+
183+
### Changes to the sycl-post-link tool
184+
185+
As mentioned in the __Attributes attached to the class__ section, the sycl-post-link tool
186+
will generate the `HostAccessINTEL` decoration for each variable declared of a
187+
type marked with the global variable attribute `sycl-host-pipe`. The name operand
188+
should be filled with the id generated by the front end when the `host-pipe` attribute
189+
is encountered. Since there is no current use for specific host access information,
190+
the access field can be set to `1` (read/write). If a use for this information
191+
is found, this can be changed in the future.

‎sycl/doc/index.rst

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -42,6 +42,7 @@ Design Documents for the oneAPI DPC++ Compiler
4242
design/ITTAnnotations
4343
design/DeviceGlobal
4444
design/CompileTimeProperties
45+
design/HostPipes
4546
design/ESIMDStatelesAccessors
4647
design/ESIMDDesignNotes
4748
design/DeviceIf

0 commit comments

Comments
 (0)