From 7efff1f584039bb377d190ced530888d082f9d02 Mon Sep 17 00:00:00 2001 From: rho180 <84344325+rho180@users.noreply.github.com> Date: Mon, 7 Mar 2022 13:27:37 -0500 Subject: [PATCH 01/19] Create HostPipes.md --- sycl/doc/design/HostPipes.md | 167 +++++++++++++++++++++++++++++++++++ 1 file changed, 167 insertions(+) create mode 100644 sycl/doc/design/HostPipes.md diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md new file mode 100644 index 0000000000000..b6c77b24b3b7a --- /dev/null +++ b/sycl/doc/design/HostPipes.md @@ -0,0 +1,167 @@ +# Implementation design for "HostPipes" + +This document describes the implementation design for the host pipes section +of the DPC++ extension [SYCL_INTEL_data_flow_pipes]. 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. + +## Requirements + +The extension specification document referenced above contains the full set of +requirements for this feature, but some requirements that are particularly +relevant to the design are called out here. + +The first 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 variable +instance in device code with the variable 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. + +## Design + +### Changes to DPC++ headers + +#### Partial specialization + +#### Attributes attached to the class + +The `pipe` class declaration borrows the C++ attribute `sycl-host-access` from +`device_global` to convey name information to the FPGA backend. Since this +is only needed for naming, we will set value of the property to `readwrite`. +As this attribute is also only needed for the device compiler, the `#ifdef __SYCL_DEVICE_ONLY__` +allows the customer to ue 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 { +ifdef __SYCL_DEVICE_ONLY__ + [[__sycl_detail__::add_ir_global_variable_attributes( + "sycl-host-access", + "readwrite" + )]] +#endif + ... +} +``` +The `[[__sycl_detail__::add_ir_global_variable_attributes()]]` attribute is +described more fully by the [compile-time properties][1] design +document. This attribute is also used for other classes that have properties, +so it is not specific to the `pipe` class. + +[1]: + +### 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 + +using a_pipe = pipe; +using b_pipe = pipe; + +``` + +The front-end will generate a 'const char * ' for each pipe class + +``` +const char *a_pipe_var; +const char *b_pipe_var; +``` + +The corresponding integration header defines a namespace scope variable of type +`__sycl_host_pipe_registration` 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 device +global variable that is defined in the translation unit: + +* The (host) address of the variable. +* 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(&::a_pipe_var, + /* same string returned from __builtin_sycl_unique_pipe_id(::a_pipe_var) */); + host_pipe_map::add(&::b_pipe_var, + /* same string returned from __builtin_sycl_unique_pipe_id(::b_pipe_var) */); +} + +} // namespace (unnamed) +} // namespace sycl::detail +``` + +Further details on adherence to C++ rules for unconstructed objects can be found +in the [device_global][2] design. + +[2]: + +Generating a unique pipe id is addressed in Open Questions below. + +### 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::device_global_map::add()`. + The runtime defines this function and maintains information about all the + device global variables in the application. This information includes: + + - The host address of the variable. + - The string which uniquely identifies the variable. + + We refer to this information as the "device global database" below. + +* The runtime implements the `read` and `write` functions of the pipe + class. These will use this [host pipe API][3]. + +[3] https://github.com/intel-sandbox/ip-authoring-specs/blob/MJ_ChangeDocs4/Pipe/Spec/cl_intel_host_pipe_symbol.asciidocdefined + +### Open Questions + +The 'unique pipe id' must be globally unique. Since all global variables in +the LLVM IR must have such a unique names, it is our intention to use this +naming. Is this possible? We would also need to define a builtin to return +this string (see 'builtin_sycl_unique_id' in the headers and footers section). + + From f17ef0f10d33486c7c9d0c2ca3bcf6386ea9e832 Mon Sep 17 00:00:00 2001 From: rho180 <84344325+rho180@users.noreply.github.com> Date: Mon, 7 Mar 2022 13:56:48 -0500 Subject: [PATCH 02/19] Update HostPipes.md --- sycl/doc/design/HostPipes.md | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md index b6c77b24b3b7a..9defe62becc23 100644 --- a/sycl/doc/design/HostPipes.md +++ b/sycl/doc/design/HostPipes.md @@ -10,10 +10,11 @@ links a device kernel with a host program. ## Requirements The extension specification document referenced above contains the full set of -requirements for this feature, but some requirements that are particularly -relevant to the design are called out here. +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. -The first issue relates to the mechanism for integrating host and device code. +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 variable instance in device code with the variable instance in host code. We will use @@ -24,8 +25,6 @@ database in the integration headers and footers. ### Changes to DPC++ headers -#### Partial specialization - #### Attributes attached to the class The `pipe` class declaration borrows the C++ attribute `sycl-host-access` from From df88193e34f64926024b2220cc3a60767cf967db Mon Sep 17 00:00:00 2001 From: rho180 <84344325+rho180@users.noreply.github.com> Date: Mon, 7 Mar 2022 13:58:07 -0500 Subject: [PATCH 03/19] Update HostPipes.md --- sycl/doc/design/HostPipes.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md index 9defe62becc23..3f8a47e610f60 100644 --- a/sycl/doc/design/HostPipes.md +++ b/sycl/doc/design/HostPipes.md @@ -154,7 +154,7 @@ Several changes are needed to the DPC++ runtime * The runtime implements the `read` and `write` functions of the pipe class. These will use this [host pipe API][3]. -[3] https://github.com/intel-sandbox/ip-authoring-specs/blob/MJ_ChangeDocs4/Pipe/Spec/cl_intel_host_pipe_symbol.asciidocdefined +[3] https://github.com/intel-sandbox/ip-authoring-specs/blob/MJ_ChangeDocs4/Pipe/Spec/cl_intel_host_pipe_symbol.asciidoc ### Open Questions From 35a5171b8b0c7b4cb64208e89fbd8fb6dbe6be4a Mon Sep 17 00:00:00 2001 From: rho180 <84344325+rho180@users.noreply.github.com> Date: Mon, 7 Mar 2022 14:03:16 -0500 Subject: [PATCH 04/19] Update HostPipes.md --- sycl/doc/design/HostPipes.md | 20 ++++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md index 3f8a47e610f60..35302fe642a96 100644 --- a/sycl/doc/design/HostPipes.md +++ b/sycl/doc/design/HostPipes.md @@ -1,12 +1,14 @@ -# Implementation design for "HostPipes" +# 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]. Pipes are a FIFO construct +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]: https://github.com/intel-sandbox/ip-authoring-specs/blob/main/Pipe/Spec/data_flow_pipes.asciidoc + ## Requirements The extension specification document referenced above contains the full set of @@ -51,11 +53,11 @@ ifdef __SYCL_DEVICE_ONLY__ } ``` The `[[__sycl_detail__::add_ir_global_variable_attributes()]]` attribute is -described more fully by the [compile-time properties][1] design +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. -[1]: +[3]: ### Changes to the DPC++ front-end @@ -133,7 +135,7 @@ __sycl_host_pipe_registration::__sycl_host_pipe_registration() noexcept { Further details on adherence to C++ rules for unconstructed objects can be found in the [device_global][2] design. -[2]: +[3]: Generating a unique pipe id is addressed in Open Questions below. @@ -142,19 +144,17 @@ Generating a unique pipe id is addressed in Open Questions below. 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::device_global_map::add()`. + footer which calls the function `sycl::detail::host_pipe_map::add()`. The runtime defines this function and maintains information about all the device global variables in the application. This information includes: - The host address of the variable. - The string which uniquely identifies the variable. - We refer to this information as the "device global database" below. - * The runtime implements the `read` and `write` functions of the pipe - class. These will use this [host pipe API][3]. + class. These will use this [host pipe API][4]. -[3] https://github.com/intel-sandbox/ip-authoring-specs/blob/MJ_ChangeDocs4/Pipe/Spec/cl_intel_host_pipe_symbol.asciidoc +[4]: https://github.com/intel-sandbox/ip-authoring-specs/blob/MJ_ChangeDocs4/Pipe/Spec/cl_intel_host_pipe_symbol.asciidoc ### Open Questions From 4baec326b9b3dd6e690386178b030f133e050fd5 Mon Sep 17 00:00:00 2001 From: rho180 <84344325+rho180@users.noreply.github.com> Date: Wed, 9 Mar 2022 09:31:19 -0500 Subject: [PATCH 05/19] Update HostPipes.md --- sycl/doc/design/HostPipes.md | 48 +++++++++++++++++++----------------- 1 file changed, 26 insertions(+), 22 deletions(-) diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md index 35302fe642a96..404de46018728 100644 --- a/sycl/doc/design/HostPipes.md +++ b/sycl/doc/design/HostPipes.md @@ -18,11 +18,13 @@ 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 variable -instance in device code with the variable instance in host code. We will use +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 @@ -42,21 +44,27 @@ support them as general attributes that customer code can use. template > -class pipe { +class pipe ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_global_variable_attributes( "sycl-host-access", "readwrite" )]] #endif +{ + static const char __pipe; ... } ``` -The `[[__sycl_detail__::add_ir_global_variable_attributes()]]` attribute is +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 char` 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 @@ -85,16 +93,9 @@ using b_pipe = pipe; ``` -The front-end will generate a 'const char * ' for each pipe class - -``` -const char *a_pipe_var; -const char *b_pipe_var; -``` - The corresponding integration header defines a namespace scope variable of type -`__sycl_host_pipe_registration` whose sole purpose is to run its -constructor before the application's main() function: +`__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 { @@ -111,10 +112,10 @@ __sycl_host_pipe_registration __sycl_host_pipe_registrar; ``` The integration footer contains the definition of the constructor, which calls -a function in the DPC++ runtime with the following information for each device -global variable that is defined in the translation unit: +a function in the DPC++ runtime with the following information for each host +pipe that is defined in the translation unit: -* The (host) address of the variable. +* The (host) address of the static member variable `__pipe`. * The variable's string from the `sycl-unique-id` attribute. ``` @@ -122,10 +123,10 @@ namespace sycl::detail { namespace { __sycl_host_pipe_registration::__sycl_host_pipe_registration() noexcept { - host_pipe_map::add(&::a_pipe_var, - /* same string returned from __builtin_sycl_unique_pipe_id(::a_pipe_var) */); - host_pipe_map::add(&::b_pipe_var, - /* same string returned from __builtin_sycl_unique_pipe_id(::b_pipe_var) */); + host_pipe_map::add(&a_pipe::__pipe, + /* same string returned from __builtin_sycl_unique_pipe_id(&a_pipe::__pipe) */); + host_pipe_map::add(&b_pipe::__pipe, + /* same string returned from __builtin_sycl_unique_pipe_id(&b_pipe::__pipe) */); } } // namespace (unnamed) @@ -133,7 +134,7 @@ __sycl_host_pipe_registration::__sycl_host_pipe_registration() noexcept { ``` Further details on adherence to C++ rules for unconstructed objects can be found -in the [device_global][2] design. +in the [device_global][3] design. [3]: @@ -152,7 +153,10 @@ Several changes are needed to the DPC++ runtime - 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]. + 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 From ce7bcef572a8f140b7a0afce4d666b8ab8a5971b Mon Sep 17 00:00:00 2001 From: rho180 <84344325+rho180@users.noreply.github.com> Date: Wed, 9 Mar 2022 09:36:23 -0500 Subject: [PATCH 06/19] Update HostPipes.md --- sycl/doc/design/HostPipes.md | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md index 404de46018728..26b6de2a0a429 100644 --- a/sycl/doc/design/HostPipes.md +++ b/sycl/doc/design/HostPipes.md @@ -44,13 +44,14 @@ support them as general attributes that customer code can use. template > -class pipe +class ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_global_variable_attributes( "sycl-host-access", "readwrite" )]] #endif +pipe { static const char __pipe; ... @@ -94,7 +95,7 @@ using b_pipe = pipe; ``` The corresponding integration header defines a namespace scope variable of type -`__sycl_host_pipe_registration` (referred to below as the __host pipe registrar__ +`__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: ``` From 3ce3d4a23b1a10962d717fa1c417cd427d17c494 Mon Sep 17 00:00:00 2001 From: rho180 <84344325+rho180@users.noreply.github.com> Date: Wed, 9 Mar 2022 13:21:57 -0500 Subject: [PATCH 07/19] Update HostPipes.md --- sycl/doc/design/HostPipes.md | 46 ++++++++++++++++++++++-------------- 1 file changed, 28 insertions(+), 18 deletions(-) diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md index 26b6de2a0a429..d9dc4f164c892 100644 --- a/sycl/doc/design/HostPipes.md +++ b/sycl/doc/design/HostPipes.md @@ -44,23 +44,25 @@ support them as general attributes that customer code can use. template > -class -ifdef __SYCL_DEVICE_ONLY__ +class pipe +{ + struct +#ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_global_variable_attributes( "sycl-host-access", "readwrite" )]] #endif -pipe -{ - static const char __pipe; + __pipeType { const char __p }; + + __pipeType __pipe; ... } ``` 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. +so it is not specific to the `pipe` class. The address of `static const char` 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 @@ -89,8 +91,15 @@ variable. To illustrate, consider a translation unit that defines two ``` #include -using a_pipe = pipe; -using b_pipe = pipe; +class some_pipe; +namespace inner { + class some_other_pipe; +} // namespace inner +... +pipe::write(...); // a usage of pipe +... +pipe::read(...); // a usage of pipe +... ``` @@ -114,7 +123,7 @@ __sycl_host_pipe_registration __sycl_host_pipe_registrar; 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 defined in the translation unit: +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. @@ -124,10 +133,10 @@ namespace sycl::detail { namespace { __sycl_host_pipe_registration::__sycl_host_pipe_registration() noexcept { - host_pipe_map::add(&a_pipe::__pipe, - /* same string returned from __builtin_sycl_unique_pipe_id(&a_pipe::__pipe) */); - host_pipe_map::add(&b_pipe::__pipe, - /* same string returned from __builtin_sycl_unique_pipe_id(&b_pipe::__pipe) */); + 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) @@ -163,9 +172,10 @@ Several changes are needed to the DPC++ runtime ### Open Questions -The 'unique pipe id' must be globally unique. Since all global variables in -the LLVM IR must have such a unique names, it is our intention to use this -naming. Is this possible? We would also need to define a builtin to return -this string (see 'builtin_sycl_unique_id' in the headers and footers section). - +* The 'unique pipe id' must be globally unique. Since all global variables in + the LLVM IR must have such a unique names, it is our intention to use this + naming. Is this possible? We would also need to define a builtin to return + this string (see 'builtin_sycl_unique_id' in the headers and footers section). +* Is the `[[__sycl_detail__::add_ir_attributes_global_variable()]]` usable on the + static class member `__pipe` as shown? From 6c5dd176977d99240b7c4bbdadf3dc45dd420dfa Mon Sep 17 00:00:00 2001 From: rho180 <84344325+rho180@users.noreply.github.com> Date: Thu, 10 Mar 2022 09:02:43 -0500 Subject: [PATCH 08/19] Update HostPipes.md --- sycl/doc/design/HostPipes.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md index d9dc4f164c892..03b6f9b52cb97 100644 --- a/sycl/doc/design/HostPipes.md +++ b/sycl/doc/design/HostPipes.md @@ -55,7 +55,7 @@ class pipe #endif __pipeType { const char __p }; - __pipeType __pipe; + static __pipeType __pipe; ... } ``` From 8f0bad9e15977e8ae7223c141e2ecb5568e22912 Mon Sep 17 00:00:00 2001 From: rho180 <84344325+rho180@users.noreply.github.com> Date: Thu, 10 Mar 2022 09:05:54 -0500 Subject: [PATCH 09/19] Update HostPipes.md --- sycl/doc/design/HostPipes.md | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md index 03b6f9b52cb97..101f4b0b107b5 100644 --- a/sycl/doc/design/HostPipes.md +++ b/sycl/doc/design/HostPipes.md @@ -53,11 +53,11 @@ class pipe "readwrite" )]] #endif - __pipeType { const char __p }; + __pipeType { const char __p; }; - static __pipeType __pipe; + static const __pipeType __pipe; ... -} +}; ``` The `[[__sycl_detail__::add_ir_attributes_global_variable()]]` attribute is described more fully by the [compile-time properties][3] design From e6f041636b31d7f9dbd6d992c79b4ae512539da6 Mon Sep 17 00:00:00 2001 From: rho180 <84344325+rho180@users.noreply.github.com> Date: Thu, 10 Mar 2022 09:06:31 -0500 Subject: [PATCH 10/19] Update HostPipes.md --- sycl/doc/design/HostPipes.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md index 101f4b0b107b5..62465fc739110 100644 --- a/sycl/doc/design/HostPipes.md +++ b/sycl/doc/design/HostPipes.md @@ -64,7 +64,7 @@ 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 char` member `__pipe` will be used to identify the pipe +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). From aa65324638ac682c940343519b537bfc3ff79ec9 Mon Sep 17 00:00:00 2001 From: rho180 <84344325+rho180@users.noreply.github.com> Date: Mon, 21 Mar 2022 13:50:43 -0400 Subject: [PATCH 11/19] Update HostPipes.md --- sycl/doc/design/HostPipes.md | 37 +++++++++++++++++++++--------------- 1 file changed, 22 insertions(+), 15 deletions(-) diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md index 62465fc739110..c8dbc961e5d66 100644 --- a/sycl/doc/design/HostPipes.md +++ b/sycl/doc/design/HostPipes.md @@ -31,10 +31,15 @@ database in the integration headers and footers. #### Attributes attached to the class -The `pipe` class declaration borrows the C++ attribute `sycl-host-access` from -`device_global` to convey name information to the FPGA backend. Since this -is only needed for naming, we will set value of the property to `readwrite`. -As this attribute is also only needed for the device compiler, the `#ifdef __SYCL_DEVICE_ONLY__` +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 ue 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 @@ -43,15 +48,17 @@ support them as general attributes that customer code can use. ``` template > + typename propertiesT = properties<>> class pipe { struct #ifdef __SYCL_DEVICE_ONLY__ [[__sycl_detail__::add_ir_global_variable_attributes( - "sycl-host-access", - "readwrite" + "sycl-host-pipe", + nullptr )]] + [[__sycl_detail__::host_pipe]] + [[__sycl_detail__::global_variable_allowed]] // may not be needed #endif __pipeType { const char __p; }; @@ -170,12 +177,12 @@ Several changes are needed to the DPC++ runtime [4]: https://github.com/intel-sandbox/ip-authoring-specs/blob/MJ_ChangeDocs4/Pipe/Spec/cl_intel_host_pipe_symbol.asciidoc -### Open Questions +### Changes to the sycl-post-link tool -* The 'unique pipe id' must be globally unique. Since all global variables in - the LLVM IR must have such a unique names, it is our intention to use this - naming. Is this possible? We would also need to define a builtin to return - this string (see 'builtin_sycl_unique_id' in the headers and footers section). - -* Is the `[[__sycl_detail__::add_ir_attributes_global_variable()]]` usable on the - static class member `__pipe` as shown? +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. From b4e42b9665b1828f522f7a666147b7ede574e2c3 Mon Sep 17 00:00:00 2001 From: rho180 <84344325+rho180@users.noreply.github.com> Date: Tue, 22 Mar 2022 11:23:51 -0400 Subject: [PATCH 12/19] Update HostPipes.md Add constexpr definition for __pipe in pipe class --- sycl/doc/design/HostPipes.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md index c8dbc961e5d66..143747014091c 100644 --- a/sycl/doc/design/HostPipes.md +++ b/sycl/doc/design/HostPipes.md @@ -62,7 +62,7 @@ class pipe #endif __pipeType { const char __p; }; - static const __pipeType __pipe; + static constexpr __pipeType __pipe = {0}; ... }; ``` From a30f73b0e2ece6fc06caea8376a5eedd1b579840 Mon Sep 17 00:00:00 2001 From: rho180 <84344325+rho180@users.noreply.github.com> Date: Thu, 18 Aug 2022 10:41:34 -0400 Subject: [PATCH 13/19] Update sycl/doc/design/HostPipes.md Co-authored-by: Steffen Larsen --- sycl/doc/design/HostPipes.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md index 143747014091c..1a3de9b75eb96 100644 --- a/sycl/doc/design/HostPipes.md +++ b/sycl/doc/design/HostPipes.md @@ -40,7 +40,7 @@ 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 ue another host compiler, even if it does not recognize these attributes. +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. From fe8c7e91cbdb7f51fee212298e471ab682657685 Mon Sep 17 00:00:00 2001 From: rho180 <84344325+rho180@users.noreply.github.com> Date: Thu, 18 Aug 2022 10:42:48 -0400 Subject: [PATCH 14/19] Update HostPipes.md --- sycl/doc/design/HostPipes.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md index 1a3de9b75eb96..8d6a99f61d513 100644 --- a/sycl/doc/design/HostPipes.md +++ b/sycl/doc/design/HostPipes.md @@ -53,7 +53,7 @@ class pipe { struct #ifdef __SYCL_DEVICE_ONLY__ - [[__sycl_detail__::add_ir_global_variable_attributes( + [[__sycl_detail__::add_ir_attributes_global_variable( "sycl-host-pipe", nullptr )]] From 632871bfe3488327f237b11639d44d13db7dc483 Mon Sep 17 00:00:00 2001 From: rho180 <84344325+rho180@users.noreply.github.com> Date: Thu, 18 Aug 2022 13:45:46 -0400 Subject: [PATCH 15/19] Update HostPipes.md --- sycl/doc/design/HostPipes.md | 13 +++++++++---- 1 file changed, 9 insertions(+), 4 deletions(-) diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md index 8d6a99f61d513..7af5b92fe66b8 100644 --- a/sycl/doc/design/HostPipes.md +++ b/sycl/doc/design/HostPipes.md @@ -46,16 +46,21 @@ they are considered implementation details of DPC++. We do not intend to support them as general attributes that customer code can use. ``` -template > +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", - nullptr + Props::meta_name..., + nullptr, + Props::meta_value... )]] [[__sycl_detail__::host_pipe]] [[__sycl_detail__::global_variable_allowed]] // may not be needed From cc6083b4ebd34e0248c7cf21d3815f7bea5c4560 Mon Sep 17 00:00:00 2001 From: rho180 <84344325+rho180@users.noreply.github.com> Date: Fri, 19 Aug 2022 09:03:50 -0400 Subject: [PATCH 16/19] Apply suggestions from code review Co-authored-by: Steffen Larsen --- sycl/doc/design/HostPipes.md | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md index 7af5b92fe66b8..ecf1d3b316e1e 100644 --- a/sycl/doc/design/HostPipes.md +++ b/sycl/doc/design/HostPipes.md @@ -7,7 +7,7 @@ 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]: https://github.com/intel-sandbox/ip-authoring-specs/blob/main/Pipe/Spec/data_flow_pipes.asciidoc +[1]: <../extensions/supported/sycl_ext_intel_dataflow_pipes.asciidoc> ## Requirements @@ -51,7 +51,7 @@ class pipe {/*...*/}; // Partial specialization to make propertiesT visible as a parameter pack // of properties. -template +template class pipe { struct @@ -169,7 +169,7 @@ 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 - device global variables in the application. This information includes: + host pipe variables in the application. This information includes: - The host address of the variable. - The string which uniquely identifies the variable. From 317b3890708322dbfc13389b295e72db51874824 Mon Sep 17 00:00:00 2001 From: Robert Ho <84344325+rho180@users.noreply.github.com> Date: Mon, 3 Oct 2022 11:24:00 -0400 Subject: [PATCH 17/19] Update HostPipes.md --- sycl/doc/design/HostPipes.md | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md index ecf1d3b316e1e..38fadbfee322f 100644 --- a/sycl/doc/design/HostPipes.md +++ b/sycl/doc/design/HostPipes.md @@ -156,9 +156,7 @@ __sycl_host_pipe_registration::__sycl_host_pipe_registration() noexcept { ``` Further details on adherence to C++ rules for unconstructed objects can be found -in the [device_global][3] design. - -[3]: +in the [device_global][2] design. Generating a unique pipe id is addressed in Open Questions below. From 060249e063f1d4cbc2bcf64dd008fbadadc9c7b5 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Mon, 9 Jan 2023 17:08:33 -0800 Subject: [PATCH 18/19] Add HostPipes document to index. --- sycl/doc/index.rst | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/doc/index.rst b/sycl/doc/index.rst index f67434e7faaa4..9799b29fc1002 100644 --- a/sycl/doc/index.rst +++ b/sycl/doc/index.rst @@ -41,6 +41,7 @@ Design Documents for the oneAPI DPC++ Compiler design/ITTAnnotations design/DeviceGlobal design/CompileTimeProperties + design/HostPipes New OpenCL Extensions New SPIR-V Extensions From b4f3f67da8407255f34c10e11d54b7f768f22ca3 Mon Sep 17 00:00:00 2001 From: Robert Ho <84344325+rho180@users.noreply.github.com> Date: Fri, 13 Jan 2023 09:08:28 -0500 Subject: [PATCH 19/19] Update HostPipes.md --- sycl/doc/design/HostPipes.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/design/HostPipes.md b/sycl/doc/design/HostPipes.md index 38fadbfee322f..6e3391496a7ec 100644 --- a/sycl/doc/design/HostPipes.md +++ b/sycl/doc/design/HostPipes.md @@ -158,7 +158,7 @@ __sycl_host_pipe_registration::__sycl_host_pipe_registration() noexcept { Further details on adherence to C++ rules for unconstructed objects can be found in the [device_global][2] design. -Generating a unique pipe id is addressed in Open Questions below. +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