diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md new file mode 100644 index 0000000000000..6e3391496a7ec --- /dev/null +++ b/sycl/doc/design/HostPipes.md @@ -0,0 +1,191 @@ +# Implementation design for "Host Pipes" + +This document describes the implementation design for the host pipes section +of the DPC++ extension [SYCL_INTEL_data_flow_pipes][1]. Pipes are a FIFO construct +that provide links between elements of a design that are accessed through read +and write application programming interfaces (APIs), without the notion of a +memory address/pointer to elements within the FIFO. A host pipe is a pipe that +links a device kernel with a host program. + +[1]: <../extensions/supported/sycl_ext_intel_dataflow_pipes.asciidoc> + +## Requirements + +The extension specification document referenced above contains the full set of +requirements for this feature, but a requirement that is particularly +relevant to the design, and similar in nature to one raised in the [device_global][2] +design is called out here. + +This issue relates to the mechanism for integrating host and device code. +Like device global variables, host pipes are referenced in both +host and device code, so they require some mechanism to correlate the pipe +instance in device code with the pipe instance in host code. We will use +a similar mechanism as the device global implementation that creates a map +database in the integration headers and footers. + +[2]: + +## Design + +### Changes to DPC++ headers + +#### Attributes attached to the class + +The `pipe` class uses a new C++ attribute `[[__sycl_detail__::host_pipe]]` on the +`pipe::__pipeType` type to identify the `static const __pipeType` member `__pipe` +as a host pipe. Similar to `[[__sycl_detail__::device_global]]`, this will inform +the front end to generate a `sycl-unique-id` for each `__pipe`. The `pipe` class +also introduces the global variable attribute `sycl-host-pipe` attribute to inform the sycl-post-link tool +to generate the SPIR-V decoration `HostAccessINTEL` for each `__pipe` using the +`sycl-unique-id` generated. + +As these attributes are only needed for the device compiler, the `#ifdef __SYCL_DEVICE_ONLY__` +allows the customer to use another host compiler, even if it does not recognize these attributes. +Also note that these attributes are all in the `__sycl_detail__` namespace, so +they are considered implementation details of DPC++. We do not intend to +support them as general attributes that customer code can use. + +``` +template > +class pipe {/*...*/}; + +// Partial specialization to make propertiesT visible as a parameter pack +// of properties. +template +class pipe +{ + struct +#ifdef __SYCL_DEVICE_ONLY__ + [[__sycl_detail__::add_ir_attributes_global_variable( + "sycl-host-pipe", + Props::meta_name..., + nullptr, + Props::meta_value... + )]] + [[__sycl_detail__::host_pipe]] + [[__sycl_detail__::global_variable_allowed]] // may not be needed +#endif + __pipeType { const char __p; }; + + static constexpr __pipeType __pipe = {0}; + ... +}; +``` +The `[[__sycl_detail__::add_ir_attributes_global_variable()]]` attribute is +described more fully by the [compile-time properties][3] design +document. This attribute is also used for other classes that have properties, +so it is not specific to the `pipe` class. + +The address of `static const __pipeType` member `__pipe` will be used to identify the pipe +in host code, and provide one half of the host-to-device mapping of the pipe +(see the section on __New content in the integration header and footer__ below). + +[3]: + +### Changes to the DPC++ front-end + +There are several changes to the device compiler front-end: + +* The front-end adds a new LLVM IR attribute `sycl-unique-id` to the definition + of each `pipe` variable, which provides a unique string identifier + for each. + +* The front-end generates new content in both the integration header and the + integration footer, which is described in more detail below. + +#### New content in the integration header and footer + +New content in the integration header and footer provides a mapping from the +host address of each pipe variable to the unique string for that +variable. To illustrate, consider a translation unit that defines two +`pipe` classes: + +``` +#include + +class some_pipe; +namespace inner { + class some_other_pipe; +} // namespace inner +... +pipe::write(...); // a usage of pipe +... +pipe::read(...); // a usage of pipe +... + +``` + +The corresponding integration header defines a namespace scope variable of type +`__sycl_host_pipe_registration` (referred to below as the __host pipe registrar__) +whose sole purpose is to run its constructor before the application's main() function: + +``` +namespace sycl::detail { +namespace { + +class __sycl_host_pipe_registration { + public: + __sycl_host_pipe_registration() noexcept; +}; +__sycl_host_pipe_registration __sycl_host_pipe_registrar; + +} // namespace (unnamed) +} // namespace sycl::detail +``` + +The integration footer contains the definition of the constructor, which calls +a function in the DPC++ runtime with the following information for each host +pipe that is used in the translation unit: + +* The (host) address of the static member variable `__pipe`. +* The variable's string from the `sycl-unique-id` attribute. + +``` +namespace sycl::detail { +namespace { + +__sycl_host_pipe_registration::__sycl_host_pipe_registration() noexcept { + host_pipe_map::add(&pipe::__pipe, + /* same string returned from __builtin_sycl_unique_pipe_id(pipe::__pipe) */); + host_pipe_map::add(&inner::pipe::__pipe, + /* same string returned from __builtin_sycl_unique_pipe_id(pipe::__pipe) */); +} + +} // namespace (unnamed) +} // namespace sycl::detail +``` + +Further details on adherence to C++ rules for unconstructed objects can be found +in the [device_global][2] design. + +Unique pipe ids will be generated by the same method as [device_global][2] uses to generate `sycl-unique-id`s. + +### Changes to the DPC++ runtime + +Several changes are needed to the DPC++ runtime + +* As we noted above, the front-end generates new content in the integration + footer which calls the function `sycl::detail::host_pipe_map::add()`. + The runtime defines this function and maintains information about all the + host pipe variables in the application. This information includes: + + - The host address of the variable. + - The string which uniquely identifies the variable. + +* The runtime implements the `read` and `write` functions of the pipe + class. These will use this [host pipe API][4]. These functions will + need to retrieve the mapping added to the __host pipe registrar__ + for the pipe being read or written to, and pass it to the corresponding + underlying OpenCL API call + +[4]: https://github.com/intel-sandbox/ip-authoring-specs/blob/MJ_ChangeDocs4/Pipe/Spec/cl_intel_host_pipe_symbol.asciidoc + +### Changes to the sycl-post-link tool + +As mentioned in the __Attributes attached to the class__ section, the sycl-post-link tool +will generate the `HostAccessINTEL` decoration for each variable declared of a +type marked with the global variable attribute `sycl-host-pipe`. The name operand +should be filled with the id generated by the front end when the `host-pipe` attribute +is encountered. Since there is no current use for specific host access information, +the access field can be set to `1` (read/write). If a use for this information +is found, this can be changed in the future. diff --git a/sycl/doc/index.rst b/sycl/doc/index.rst index 15122f0dfd767..0432177e09007 100644 --- a/sycl/doc/index.rst +++ b/sycl/doc/index.rst @@ -42,6 +42,7 @@ Design Documents for the oneAPI DPC++ Compiler design/ITTAnnotations design/DeviceGlobal design/CompileTimeProperties + design/HostPipes design/ESIMDStatelesAccessors design/ESIMDDesignNotes design/DeviceIf