[SYCL][Doc] Update SYCL_INTEL_data_flow_pipes extension for FPGA host pipe support

Add a memory order parameter to device-side read/write members and
default to `sycl::memory_order::seq_cst`.

Replace `min_capacity` property with compile-time properties list for
use with `SYCL_INTEL_FPGA_data_flow_pipes_properties` extension.

Add host pipe read/write members with additional `sycl::queue` parameter.

Signed-off-by: Peter Colberg <peter.colberg@intel.com>

Author: pcolberg

sycl/doc/extensions/proposed/sycl_ext_intel_dataflow_pipes.asciidoc

@@ -35,7 +35,7 @@ This document describes an extension that adds pipes to SYCL.  Pipes are first i
 
 == Notice
 
-Copyright (c) 2019-2021 Intel Corporation.  All rights reserved.
+Copyright (c) 2019-2022 Intel Corporation.  All rights reserved.
 
 == Status
 
@@ -53,10 +53,24 @@ Revision: 3
 == Contact
 Michael Kinsner, Intel (michael 'dot' kinsner 'at' intel 'dot' com)
 
+== Contributors
+
+Michael Kinsner, Intel +
+Shuo Niu, Intel +
+Bo Lei, Intel +
+Marco Jacques, Intel +
+Joe Garvey, Intel +
+Aditi Kumaraswamy, Intel +
+Robert Ho, Intel +
+Sherry Yuan, Intel +
+Peter Colberg, Intel 
+
 == Dependencies
 
 This extension is written against the SYCL 2020 specification, Revision 3.
 
+It also depends on the `sycl_ext_oneapi_properties` extension.
+
 The use of blocking pipe reads or writes requires support for https://github.com/KhronosGroup/SPIRV-Registry/blob/master/extensions/INTEL/SPV_INTEL_blocking_pipes.asciidoc[SPV_INTEL_blocking_pipes] if SPIR-V is used by an implementation.
 
 == Overview
@@ -116,7 +130,7 @@ A pipe type is a specialization of the pipe class:
 ----
 template <typename name,
           typename dataT,
-          size_t min_capacity = 0>
+          typename propertiesT = properties<>>
 class pipe;
 ----
 
@@ -129,7 +143,7 @@ A difference in any of the three template parameters identifies a different pipe
 using pipe<class foo, int>;
 using pipe<class bar, int>;
 using pipe<class bar, float>;
-using pipe<class bar, float, 5>;
+using pipe<class bar, float, decltype(properties{min_capacity<5>})>;
 ----
 
 
@@ -174,19 +188,21 @@ The pipe class exposes static member functions for writing a data word to a pipe
 
 Blocking and non-blocking forms of the read and write members are defined, with the form chosen based on overload resolution.
 
+The `sycl::memory_order` parameter of read/write functions controls how other memory accesses, including regular, non-atomic memory accesses, are to be ordered around the pipe read/write operation.  The default memory order is `sycl::memory_order::seq_cst`.
+
 [source,c++,Read write members,linenums]
 ----
 template <typename name,
           typename dataT,
-          size_t min_capacity = 0>
+          typename propertiesT = properties<>>
 class pipe {
   // Blocking
-  static dataT read();
-  static void write( const dataT &data );
+  static dataT read( memory_order order = memory_order::seq_cst );
+  static void write( const dataT &data, memory_order order = memory_order::seq_cst );
 
   // Non-blocking
-  static dataT read( bool &success_code );
-  static void write( const dataT &data, bool &success_code );
+  static dataT read( bool &success_code, memory_order order = memory_order::seq_cst );
+  static void write( const dataT &data, bool &success_code, memory_order order = memory_order::seq_cst );
 }
 ----
 
@@ -196,7 +212,7 @@ The template parameters of the device type are defined as:
 
 * `name`: Type that is the basis of pipe identification.  Typically a user-defined class, in a user namespace.  Forward declaration of the type is sufficient, and the type does not need to be defined.
 * `dataT`: The type of data word/packet contained within a pipe.  This is the data type that is read during a successful `pipe::read` operation, or written during a successful `pipe::write` operation.  The type must be standard layout and trivially copyable.
-* `min_capacity`: User defined minimum number of words in units of `dataT` that the pipe must be able to store without any being read out.  A minimum capacity is required in some algorithms to avoid deadlock, or for performance tuning.  An implementation can include more capacity than this parameter, but not less.
+* `propertiesT`: The list of properties that are associated with the pipe.
 
 == Pipe types and {cpp} scope
 
@@ -254,6 +270,26 @@ Type aliases in {cpp} through the `using` mechanism do not change the type of a
   pipe<type_alias, int>::write(0);
 ----
 
+== Host pipe read/write
+
+The read/write member functions of a host pipe have different signatures when they are called from the host side, in which case a `sycl::queue` is added to the parameters.
+
+[source,c++,Host pipe read write members,linenums]
+----
+template <typename name,
+          typename dataT,
+          typename propertiesT = properties<>>
+class pipe {
+  // Blocking
+  static dataT read( const queue &q, memory_order order = memory_order::seq_cst );
+  static void write( const queue &q, const dataT &data, memory_order order = memory_order::seq_cst );
+
+  // Non-blocking
+  static dataT read( const queue &q, bool &success_code, memory_order order = memory_order::seq_cst );
+  static void write( const queue &q, const dataT &data, bool &success_code, memory_order order = memory_order::seq_cst );
+}
+----
+
 == Host pipe map/unmap
 
 Pipes expose two additional static member functions that are available within host code, and which map to the OpenCL C host pipe extension map/unmap interface.  These member functions provide higher bandwidth or otherwise more efficient communication on some platforms, by allowing block transfers of larger data sets.
@@ -262,7 +298,7 @@ Pipes expose two additional static member functions that are available within ho
 ----
 template <typename name,
           typename dataT,
-          size_t min_capacity = 0>
+          typename propertiesT = properties<>>
 class pipe {
   template <pipe_property::writeable host_writeable>
     static dataT* map(size_t requested_size, size_t &mapped_size);
@@ -303,7 +339,7 @@ Multiple reads or multiple writes to the same pipe from more than one kernel are
 When there are accesses to a pipe from different work-items or host threads, the order of data written to or read from the pipe is not defined.  Specifically, regarding multiple accesses to the same pipe:
 
 1. *Accesses to a single pipe within a single work-item of a kernel or thread of the host program:* Operations on the same pipe occur in program order with respect to the work-item or host thread.  No "concurrent" accesses or reordering of accesses are observable from the perspective of the single pipe.  If there are multiple pipe access operations to the same pipe within a single kernel, they execute in program order from the perspective of a single work-item.
-2. *Accesses to multiple pipes within a single work-item of a kernel or thread of the host program:*  Different pipes are treated in the same way as non-aliased memory, in that accesses to one pipe may be reordered relative to accesses to another pipe.  There is no expectation of program ordering of pipe operations across different pipes, only for a single pipe.  If a happens-before relationship across pipes is required, synchronization mechanisms such as atomics or barriers must be used.
+2. *Accesses to multiple pipes within a single work-item of a kernel or thread of the host program:*  Different pipes are treated in the same way as non-aliased memory, in that accesses to one pipe may be reordered relative to accesses to another pipe.  There is no expectation of program ordering of pipe operations across different pipes, only for a single pipe unless a memory order stronger than `memory_order_relaxed` or some other synchronization mechanism, such as a barrier, is used.
 3. *Accesses to a single pipe within two work-items of the same kernel (same or different invocations of a single kernel), and/or threads of the host program:* No ordering guarantees are made on the order of pipe operations across device work-items or host threads.  For example, if two work-items executing a kernel write to a pipe, there are no guarantees that the work-item with lower _id_ (for any definition of _id_) executes before the pipe write from a higher _id_.  The execution order of work-items executing a kernel are not defined by SYCL, may be dynamically reordered, and may not be deterministic.  If ordering guarantees are required across work-items and/or host threads, synchronization mechanisms such as atomics or barriers must be used.
 
 === Restrictions on pipes accessed by both kernels and the host program
@@ -428,8 +464,6 @@ Pipes in the context of this extension step outside the OpenCL and SYCL memory m
 
 . There is no implicit synchronizes-with relationship between different pipes and/or with non-pipe memory in a named address space (e.g. global, local, private).  Specifically, there is no implicit global or local release of side effects through a pipe access, and observation of data or control information on one pipe does not imply any knowledge through happens-before relationship with a different pipe or with memory not associated with the pipe.
 
-. Pipe read and write operations behave as if they are SYCL relaxed atomic load and store operations.  When paired with sycl::atomic_fences to establish a sychronizes-with relationship, pipe operations can provide guarantees on side effect visibility in memory, as defined by the SYCL memory model.
-
 . At a work-group barrier, there is an implicit acquire and release of side effects for any pipes operated on within the kernel, either before or after the barrier.  This occurs without an explicit memory fence being applied to or around the barrier.
 
 . There are no guarantees on pipe operation side effect latency.  Writes to a pipe will eventually be visible to read operations on the pipe, without a synchronization point, but that visibility is not guaranteed to be by the time that the next instruction is executed by a writing work-item, for example.  There may be arbitrary latency between a write to a pipe and visibility of the data on a read endpoint of the pipe.  Likewise, there may be arbitrary latency between a read from a pipe, and visibility at a write endpoint that there is capacity available to write to (assuming that capacity was full prior to the read).
@@ -752,6 +786,7 @@ extension's APIs the implementation supports.
 |2|2019-11-13|Michael Kinsner|Incorporate feedback
 |3|2020-04-27|Michael Kinsner|Clarify that pipe operations behave as-if they are relaxed atomic operations.  Make SYCL2020 the baseline
 |4|2021-12-02|Shuo Niu|Add experimental latency control API
+|5|2022-03-18|Peter Colberg|Add memory order parameter and compile-time properties.  Add host pipe read/write functions.
 |========================================
 
 //************************************************************************