[SYCL][DOC] Add extension for annotated_arg class and its properties (#6024)

- The `sycl_ext_oneapi_annotated_arg` extension introduces the class template `annotated_ptr<T, properties<...>>`. It is meant to be used for applying properties on kernel arguments.

- The `sycl_ext_intel_fpga_annotated_arg_properties`  extension defines additional supported properties.

- Note that the class `annotated_ptr` is also defined for similar purposes in a separate extension (#5755). That class applies properties on *pointer* kernel arguments and retains some of those properties on the pointer read/write sites in device code.

Author: tiwaria1

sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_arg_properties.asciidoc

@@ -0,0 +1,280 @@
+= sycl_ext_intel_fpga_annotated_arg_properties
+
+:source-highlighter: coderay
+:coderay-linenums-mode: table
+
+// This section needs to be after the document title.
+:doctype: book
+:toc2:
+:toc: left
+:encoding: utf-8
+:lang: en
+
+:blank: pass:[ +]
+
+// Set the default source code type in this document to C++,
+// for syntax highlighting purposes.  This is needed because
+// docbook uses c++ and html5 uses cpp.
+:language: {basebackend@docbook:c++:cpp}
+
+// This is necessary for asciidoc, but not for asciidoctor
+:cpp: C++
+:dpcpp: DPC++
+
+== Notice
+
+Copyright (c) 2021 Intel Corporation.  All rights reserved.
+
+NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are
+trademarks of The Khronos Group Inc.  OpenCL(TM) is a trademark of Apple Inc.
+used by permission by Khronos.
+
+== Contact
+
+To report problems with this extension, please open a new issue at:
+
+https://github.com/intel/llvm/issues
+
+== Contributors
+
+Abhishek Tiwari, Intel +
+Joseph Garvey, Intel
+
+
+
+
+== Dependencies
+
+This extension is written against the SYCL 2020 specification, revision 4.
+
+It depends on the following extensions:
+
+ - link:../experimental/sycl_ext_oneapi_properties.asciidoc[sycl_ext_oneapi_properties]
+ - link:sycl_ext_oneapi_annotated_arg.asciidoc[sycl_ext_oneapi_annotated_arg]
+
+== Status
+
+This is a proposed extension specification, intended to gather community
+feedback.  Interfaces defined in this specification may not be implemented yet
+or may be in a preliminary state.  The specification itself may also change in
+incompatible ways before it is finalized.  *Shipping software products should
+not rely on APIs defined in this specification.*
+
+== Overview
+
+This extension introduces properties for the class
+`sycl::ext::oneapi::annotated_arg`. The properties will influence the kernel
+argument interfaces for FPGA kernels and can be ignored for other devices.
+
+Some examples of the syntax are shown below.
+
+.Example 1
+[source,c++]
+----
+annotated_arg<MyType, properties<register_map>> ptr_a;
+----
+
+.Example 2
+[source,c++]
+----
+auto data = ...
+auto arg = annotated_arg(data, properties{register_map});
+----
+
+== Specification
+
+=== Feature test macro
+
+This extension provides a feature-test macro as described in the core SYCL
+specification.  An implementation supporting this extension must predefine the
+macro `SYCL_EXT_INTEL_FPGA_ANNOTATED_ARG_PROPERTIES` to one of the values
+defined in the table below.  Applications can test for the existence of this
+macro to determine if the implementation supports this feature, or applications
+can test the macro's value to determine which of the extension's features the
+implementation supports.
+
+[%header,cols="1,5"]
+|===
+|Value
+|Description
+
+|1
+|Initial version of this extension.
+|===
+
+=== `annotated_arg` Properties
+
+Below is a list of compile-time constant properties supported with
+`annotated_arg`. These properties control the kernel argument interface on FPGA
+devices.
+
+```c++
+namespace sycl::ext::oneapi::experimental {
+struct register_map_key {
+  using value_t = property_value<register_map_key>;
+};
+
+inline constexpr register_map_key::value_t register_map;
+
+template<> struct is_property_key<register_map_key> : std::true_type {};
+
+template <typename T, typename PropertyListT>
+struct is_property_key_of<register_map_key,
+  annotated_arg<T, PropertyListT>> : std::true_type {};
+
+struct conduit_key {
+  using value_t = property_value<conduit_key>;
+};
+
+inline constexpr conduit_key::value_t conduit;
+
+template<> struct is_property_key<conduit_key> : std::true_type {};
+
+template <typename T, typename PropertyListT>
+struct is_property_key_of<conduit_key,
+  annotated_arg<T, PropertyListT>> : std::true_type {};
+
+struct stable_key {
+  using value_t = property_value<stable_key>;
+};
+
+inline constexpr stable_key::value_t stable;
+
+template<> struct is_property_key<stable_key> : std::true_type {};
+
+template <typename T, typename PropertyListT>
+struct is_property_key_of<stable_key,
+  annotated_arg<T, PropertyListT>> : std::true_type {};
+} // namespace experimental::oneapi::ext::sycl
+```
+--
+
+[frame="topbot",options="header"]
+|===
+|Property |Description
+
+a|
+[source,c++]
+----
+conduit
+----
+a|
+Directs the compiler to create a dedicated input port on the kernel for the
+input data.
+
+a|
+[source,c++]
+----
+register_map
+----
+a|
+Directs the compiler to create a register to store the base address of the
+of the pointer interface as opposed to creating a dedicated input port on the
+kernel for supplying the pointer base address.
+
+a|
+[source,c++]
+----
+stable
+----
+a|
+Specifies that the input pointer address to the kernel will not change during
+the execution of the kernel. The input can still change after all active
+kernel invocations have finished.
+
+If the input is changed while the kernel is executing, the behavior is
+undefined.
+
+|===
+--
+
+=== Aliases provided for convenience
+
+[source,c++]
+----
+namespace sycl::ext::oneapi::experimental{
+  template <typename T, typename PropertyListT>
+  using register_map = annotated_arg<T, properties{
+    register_map, PropertyListT}>;
+
+  template <typename T, typename PropertyListT>
+  using conduit = annotated_arg<T, properties{
+    conduit, PropertyListT}>;
+}; // namespace sycl::ext::oneapi::experimental
+----
+
+=== Usage Examples
+
+The examples below show a simple kernel with two integer arguments marked with
+`register_map` and `stable` properties.
+
+.Usage example with a SYCL functor
+```c++
+using sycl::ext::oneapi::experimental;
+struct MyKernel {
+  using RegisterMapArg = annotated_arg<int, properties<register_map, stable>>;
+  RegisterMapArg a;
+  RegisterMapArg b;
+  ...
+  void operator()() const {
+    ... = a * b;
+  }
+};
+
+int main () {
+  sycl::queue q;
+  int data_a = ...
+  int data_b = ...
+  
+  MyKernel my_k;
+  my_k.a = data_a;
+  my_k.a = data_b;
+  ...
+  q.single_task(my_k).wait();
+  ...
+}
+```
+
+.Usage example with a SYCL lambda
+```c++
+using sycl::ext::oneapi::experimental;
+
+int main () {
+  sycl::queue q;
+  int data_a = ...
+  int data_b = ...
+  auto a = annotated_arg(data_a, properties{register_map, stable});
+  auto b = annotated_arg(data_b, properties{register_map, stable});
+  ...
+  q.single_task([=] {
+    ... = a * b;
+  }).wait();
+  ...
+}
+```
+
+== Issues
+
+1. How to document the motivation for this without duplicating what we already
+wrote for the `annotated_ptr` extension? Is the duplication acceptable?
+
+2. TODO: Correct the syntax of the aliases provided in this document.
+
+== Revision History
+
+[cols="5,15,15,70"]
+[grid="rows"]
+[options="header"]
+|========================================
+|Rev|Date|Author|Changes
+|1|2022-04-13|Abhishek Tiwari|*Initial draft*
+|========================================
+
+//************************************************************************
+//Other formatting suggestions:
+//
+//* Use *bold* text for host APIs, or [source] syntax highlighting.
+//* Use +mono+ text for device APIs, or [source] syntax highlighting.
+//* Use +mono+ text for extension names, types, or enum values.
+//* Use _italics_ for parameters.
+//************************************************************************