[SYCL][Doc] Host pipe design

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]: <DeviceGlobal.md>
+
+## 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 <typename name, typename dataT, typename propertiesT = properties<>>
+class pipe {/*...*/};
+
+// Partial specialization to make propertiesT visible as a parameter pack
+// of properties.
+template <typename Name, typename DataT, typename ...Props>
+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]: <CompileTimeProperties.md>
+
+### 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 <sycl/sycl.hpp>
+
+class some_pipe;
+namespace inner {
+  class some_other_pipe;
+} // namespace inner
+...
+pipe<class some_pipe, ...>::write(...); // a usage of pipe<class some_pipe, ...>
+...
+pipe<class some_other_pipe, ...>::read(...); // a usage of pipe<class some_other_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<some_pipe, ...>::__pipe,
+    /* same string returned from __builtin_sycl_unique_pipe_id(pipe<some_pipe, ...>::__pipe) */);
+  host_pipe_map::add(&inner::pipe<some_other_pipe>::__pipe,
+    /* same string returned from __builtin_sycl_unique_pipe_id(pipe<some_other_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.

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