diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_arg_properties.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_arg_properties.asciidoc deleted file mode 100644 index 35e17671b40a7..0000000000000 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_arg_properties.asciidoc +++ /dev/null @@ -1,280 +0,0 @@ -= 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> 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; -}; - -inline constexpr register_map_key::value_t register_map; - -template<> struct is_property_key : std::true_type {}; - -template -struct is_property_key_of> : std::true_type {}; - -struct conduit_key { - using value_t = property_value; -}; - -inline constexpr conduit_key::value_t conduit; - -template<> struct is_property_key : std::true_type {}; - -template -struct is_property_key_of> : std::true_type {}; - -struct stable_key { - using value_t = property_value; -}; - -inline constexpr stable_key::value_t stable; - -template<> struct is_property_key : std::true_type {}; - -template -struct is_property_key_of> : 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 - using register_map = annotated_arg; - - template - using conduit = annotated_arg; -}; // 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>; - 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. -//************************************************************************ diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_arg_properties.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_arg_properties.asciidoc new file mode 100644 index 0000000000000..6852a3f4c4a26 --- /dev/null +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_arg_properties.asciidoc @@ -0,0 +1,547 @@ += sycl_ext_intel_fpga_kernel_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 + +[%hardbreaks] +Copyright (C) 2022-2022 Intel Corporation. All rights reserved. + +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 5. + +It depends on the following extensions: + + - link:../experimental/sycl_ext_oneapi_properties.asciidoc[sycl_ext_oneapi_properties] + - link:sycl_ext_oneapi_annotated_ptr.asciidoc[sycl_ext_oneapi_annotated_ptr] + - 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 classes +`sycl::ext::oneapi::experimental::annotated_ptr` and +`sycl::ext::oneapi::experimental::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. + +[source,c++] +---- +auto ptr = ... +annotated_ptr arg_a{ + ptr, properties{buffer_location<1>, awidth<32>, dwidth<64>}}; + +... + +auto ptr2 = ... +annotated_arg arg_b{ + ptr2, properties{buffer_location<2>, awidth<32>, dwidth<64>}}; + +... + +int val = 5; +annotated_arg arg_c{val, properties{register_map, stable}}; + +... + +annotated_ptr, awidth<32>, + dwidth<64>> arg_d; + +annotated_arg> arg_e; + +annotated_arg, awidth<32>, + dwidth<64>> arg_f; + + +---- + + +== 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_KERNEL_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_ptr` and `annotated_arg` Properties + +Below is a list of compile-time constant properties supported by +the `annotated_ptr` and `annotated_arg` classes. These properties control the +kernel argument interface on FPGA devices. The properties are allowed even on +kernels that are submitted to other devices, but they are silently ignored when +the kernel is submitted to a non-FPGA device. + +All of the properties defined in this extension are meaningful only on the +kernel argument and are not meaningful within the kernel body. + +The properties that are defined for pointers will be meaningful for +`annotated_arg` only when T is a pointer type. + +```c++ +namespace sycl::ext::intel::experimental { +struct conduit_key { + using value_t = sycl::ext::oneapi::experimental::properties::property_value< + conduit_key>; +}; + +struct register_map_key { + using value_t = sycl::ext::oneapi::experimental::properties::property_value< + register_map_key>; +}; + +struct stable_key { + using value_t = sycl::ext::oneapi::experimental::properties::property_value< + stable_key>; +}; + +struct buffer_location_key { + template + using value_t = sycl::ext::oneapi::experimental::properties::property_value< + buffer_location_key, std::integral_constant>; +}; + +struct awidth_key { + template + using value_t = sycl::ext::oneapi::experimental::properties::property_value< + awidth_key, std::integral_constant>; +}; + +struct dwidth_key { + template + using value_t = sycl::ext::oneapi::experimental::properties::property_value< + dwidth_key, std::integral_constant>; +}; + +enum class read_write_mode_enum { + read_write, + read, + write +}; + +struct read_write_mode_key { + template + using value_t = sycl::ext::oneapi::experimental::properties::property_value< + read_write_mode_key, std::integral_constant>; +}; + +struct latency_key { + template + using value_t = sycl::ext::oneapi::experimental::properties::property_value< + latency_key, std::integral_constant>; +}; + +struct maxburst_key { + template + using value_t = sycl::ext::oneapi::experimental::properties::property_value< + maxburst_key, std::integral_constant>; +}; + +struct wait_request_key { + template + using value_t = sycl::ext::oneapi::experimental::properties::property_value< + wait_request_key, std::integral_constant>; +}; + +inline constexpr conduit_key::value_t + conduit; +inline constexpr register_map_key::value_t + register_map; +inline constexpr stable_key::value_t stable; +template inline constexpr buffer_location_key::value_t + buffer_location; +template inline constexpr awidth_key::value_t + awidth; +template inline constexpr dwidth_key::value_t + dwidth; +template +inline constexpr read_write_mode_key::value_t + read_write_mode; +inline constexpr read_write_mode_key::value_t< + read_write_mode_enum::read> read_write_mode_read; +inline constexpr read_write_mode_key::value_t< + read_write_mode_enum::write> read_write_mode_write; +inline constexpr read_write_mode_key::value_t< + read_write_mode_enum::read_write> + read_write_mode_readwrite; +template inline constexpr latency_key::value_t + latency; +template inline constexpr maxburst_key::value_t + maxburst; +template inline constexpr wait_request_key::value_t + wait_request; +inline constexpr wait_request_key::value_t + wait_request_requested; +inline constexpr wait_request_key::value_t + wait_request_not_requested; +} // namespace sycl::ext::intel::experimental + +// Type trait specializations +namespace sycl::ext::oneapi::experimental { +template<> struct is_property_key< + sycl::ext::intel::experimental::conduit_key> : std::true_type {}; +template<> struct is_property_key< + sycl::ext::intel::experimental::register_map_key> : std::true_type {}; +template<> struct is_property_key< + sycl::ext::intel::experimental::stable_key> : std::true_type {}; +template<> struct is_property_key< + sycl::ext::intel::experimental::buffer_location_key> : std::true_type {}; +template<> struct is_property_key< + sycl::ext::intel::experimental::awidth_key> : std::true_type {}; +template<> struct is_property_key< + sycl::ext::intel::experimental::dwidth_key> : std::true_type {}; +template<> struct is_property_key< + sycl::ext::intel::experimental::read_write_mode_key> : std::true_type {}; +template<> struct is_property_key< + sycl::ext::intel::experimental::latency_key> : std::true_type {}; +template<> struct is_property_key< + sycl::ext::intel::experimental::maxburst_key> : std::true_type {}; +template<> struct is_property_key< + sycl::ext::intel::experimental::wait_request_key> : std::true_type {}; + +template +struct is_property_key_of< + sycl::ext::intel::experimental::conduit_key, + annotated_ptr> : std::true_type {}; +template +struct is_property_key_of< + sycl::ext::intel::experimental::register_map_key, + annotated_ptr> : std::true_type {}; +template +struct is_property_key_of< + sycl::ext::intel::experimental::stable_key, + annotated_ptr> : std::true_type {}; +template +struct is_property_key_of< + sycl::ext::intel::experimental::buffer_location_key, + annotated_ptr> : std::true_type {}; +template +struct is_property_key_of< + sycl::ext::intel::experimental::awidth_key, + annotated_ptr> : std::true_type {}; +template +struct is_property_key_of< + sycl::ext::intel::experimental::dwidth_key, + annotated_ptr> : std::true_type {}; +template +struct is_property_key_of< + sycl::ext::intel::experimental::read_write_mode_key, + annotated_ptr> : std::true_type {}; +template +struct is_property_key_of< + sycl::ext::intel::experimental::latency_key, + annotated_ptr> : std::true_type {}; +template +struct is_property_key_of< + sycl::ext::intel::experimental::maxburst_key, + annotated_ptr> : std::true_type {}; +template +struct is_property_key_of< + sycl::ext::intel::experimental::wait_request_key, + annotated_ptr> : std::true_type {}; + +template +struct is_property_key_of< + sycl::ext::intel::experimental::conduit_key, + annotated_arg> : std::true_type {}; +template +struct is_property_key_of< + sycl::ext::intel::experimental::register_map_key, + annotated_arg> : std::true_type {}; +template +struct is_property_key_of< + sycl::ext::intel::experimental::stable_key, + annotated_arg> : std::true_type {}; +template +struct is_property_key_of< + sycl::ext::intel::experimental::buffer_location_key, + annotated_arg> : std::true_type {}; +template +struct is_property_key_of< + sycl::ext::intel::experimental::awidth_key, + annotated_arg> : std::true_type {}; +template +struct is_property_key_of< + sycl::ext::intel::experimental::dwidth_key, + annotated_arg> : std::true_type {}; +template +struct is_property_key_of< + sycl::ext::intel::experimental::read_write_mode_key, + annotated_arg> : std::true_type {}; +template +struct is_property_key_of< + sycl::ext::intel::experimental::latency_key, + annotated_arg> : std::true_type {}; +template +struct is_property_key_of< + sycl::ext::intel::experimental::maxburst_key, + annotated_arg> : std::true_type {}; +template +struct is_property_key_of< + sycl::ext::intel::experimental::wait_request_key, + annotated_arg> : std::true_type {}; +} // namespace sycl::ext::oneapi::experimental +``` +-- + +[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. + +a| +[source,c++] +---- +register_map +---- +a| +Directs the compiler to create a register to store the input as opposed to +creating a dedicated input port on the kernel. + +a| +[source,c++] +---- +stable +---- +a| +While the SYCL software model makes kernel arguments read-only, the IP which is +output by the FPGA device compiler can be plugged into external systems where +kernel arguments can change while the kernel executes. + +This property specifies that the input to the kernel will not change between +pipelined invocations of the kernel. The input can still change after all active +kernel invocations have finished. + +If the input is changed while the pipelined kernel invocations are executing, +the behavior is undefined. + +a| +[source,c++] +---- +buffer_location +---- +a| +Specifies a global memory identifier for the pointer interface. + +This property is only meaningful on pointer kernel arguments. + +a| +[source,c++] +---- +awidth +---- +a| +Specifies the width of the memory-mapped address bus in bits. The default is +determined by the implementation. + +This property is only meaningful for pointer kernel arguments and only +when the `buffer_location` property is specified. + +a| +[source,c++] +---- +dwidth +---- +a| +Specifies the width of the memory-mapped data bus in bits. The default is set +to 64. + +This property is only meaningful for pointer kernel arguments and only +when the `buffer_location` property is specified. + +a| +[source,c++] +---- +read_write_mode +---- +a| +Specifies the port direction of the memory interface associated with the input +pointer. `mode` can be one of: + +`read_write` - Interface can be used for read and write operations. + +`read` - Interface can only be used for read operations. + +`write` - Interface can only be used for write operations. + +The default is set to `read_write`. + +For convenience, the following are provided: + + - read_write_mode_read + - read_write_mode_write + - read_write_mode_readwrite + +This property is only meaningful for pointer kernel arguments and only +when the `buffer_location` property is specified. + +a| +[source,c++] +---- +latency +---- +a| +Specifies the guaranteed latency in cycles, from when a read command exits +the kernel to when the external memory returns valid read data. The default +is set to 1. + +A value of 0 specifies a variable latency and a positive value specifies a +fixed latency. + +This property is only meaningful for pointer kernel arguments and only +when the `buffer_location` property is specified. + +a| +[source,c++] +---- +maxburst +---- +a| +Specifies the maximum number of data transfers that can be associated with a +read or write transaction. The default is set to 1. + +This property is only meaningful for pointer kernel arguments and only +when the `buffer_location` property is specified. + +a| +[source,c++] +---- +wait_request +---- +a| +Specifies whether the 'wait request' signal is generated or not. This signal is +asserted by the memory system when it is unable to respond to a read or write +request. The default is set to `false`. + +For convenience, the following are provided: + + - wait_request_requested + - wait_request_not_requested + +This property is only meaningful for pointer kernel arguments and only +when the `buffer_location` property is specified. +|=== +-- + +=== Usage Examples + +The example below shows a simple kernel with one `annotated_ptr` kernel +argument and one `annotated_arg` kernel argument. + +.Usage Example +```c++ +using sycl::ext::intel::experimental; +{ + sycl::queue q{...}; + + // Allocate memory + auto ptr_a = ... + constexpr int kN = 10; + + // Add properties + auto arg_a = annotated_ptr(ptr_a, properties{ + register_map, buffer_location<1>, awidth<18>, dwidth<64>}); + auto arg_n = annotated_arg(kN, properties{register_map, stable}); + + q.single_task([=] { + for (int i=0; i> a; - annotated_arg> b; + MyType arg_a; + MyType arg_b; ... void operator()() const { ... } }; ---- -When a nested structure is created by wrapping other types with `annotated_arg`, -only the top-most level `annotated_arg` properties will be applied to the -kernel arguments. - -In the example above, only properties `PropA` and `PropB` will be applied to the -kernel argument and the other properties on members of `MyType` will not be -processed. +It is illegal to apply `annotated_arg` to members of kernel arguments. In the +above example, encapsulating `annotated_arg` within `MyType` is illegal. === Representation of `annotated_arg` [source,c++] ---- -namespace sycl::ext::oneapi { +namespace sycl::ext::oneapi::experimental { template < typename T, typename PropertyListT = properties<>> class annotated_arg { ... @@ -173,163 +164,54 @@ class annotated_arg { `annotated_arg` is a class template, parameterized by the type of the underlying allocation `T` and a list of associated properties specified by `PropertyListT`. +`T` can be any type except the following types or a structure containing one of +the following types: + +* sycl::accessor +* sycl::stream +* sycl::local_accessor +* sycl::unsampled_image_accessor +* sycl::sampled_image_accessor +* sycl::half + The properties supported with `annotated_arg` may be defined in -separate extensions. +separate extensions. Please note that there cannot be duplicated property in a +properties list. Otherwise, a compiler time error is triggered. The section below describes the constructors and member functions for `annotated_arg`. +The template parameter `T` in the definition of `annotated_arg` template below +must be a legal parameter type as defined by the SYCL specification. Given `annotated_arg`, +`T` must be a device copy-able type. + [source,c++] ---- -namespace sycl::ext::oneapi { - template > - class annotated_arg { - T data; - +namespace sycl::ext::oneapi::experimental { +template > +class annotated_arg { public: - annotated_arg(); - annotated_arg(const T& v_); - - // Conversion operator to convert to the underlying type - operator T&() noexcept; - operator const T&() const noexcept; - - // Available if the operator+ is valid for objects of type T - T operator+(const T&) noexcept; - const T operator+(const T&) const noexcept; - - // Available if the operator+ is valid for objects of type T - T operator+() noexcept; - const T operator+() const noexcept; - - // Available if the operator- is valid for objects of type T - T operator-(const T&) noexcept; - const T operator-(const T&) const noexcept; - - // Available if the operator- is valid for objects of type T - T operator-() noexcept; - const T operator-() const noexcept; - - // Available if the operator* is valid for objects of type T - T operator*(const T&) noexcept; - const T operator*(const T&) const noexcept; - - // Available if the operator/ is valid for objects of type T - T operator/(const T&) noexcept; - const T operator/(const T&) const noexcept; - - // Available if the operator% is valid for objects of type T - T operator%(const T&) noexcept; - const T operator%(const T&) const noexcept; - - // Available if the operator^ is valid for objects of type T - T operator^(const T&) noexcept; - const T operator^(const T&) const noexcept; - - // Available if the operator| is valid for objects of type T - T operator|(const T&) noexcept; - const T operator|(const T&) const noexcept; - - // Available if the operator& is valid for objects of type T - T operator&(const T&) noexcept; - const T operator&(const T&) const noexcept; - - // Available if the operator~ is valid for objects of type T - T operator~() noexcept; - const T operator~() const noexcept; - - // Available if the operator! is valid for objects of type T - T operator!() noexcept; - const T operator!() const noexcept; - - // Assignment from underlying type - T& operator=(const T&) noexcept; - - // Available if the operator< is valid for objects of type T - bool operator<(const T&) const noexcept; - - // Available if the operator> is valid for objects of type T - bool operator>(const T&) const noexcept; - - // Available if the operator+= is valid for objects of type T - T& operator+=(const T&) noexcept; - - // Available if the operator-= is valid for objects of type T - T& operator-=(const T&) noexcept; - - // Available if the operator*= is valid for objects of type T - T& operator*=(const T&) noexcept; - - // Available if the operator/= is valid for objects of type T - T& operator/=(const T&) noexcept; - - // Available if the operator%= is valid for objects of type T - T& operator%=(const T&) noexcept; - - // Available if the operator^= is valid for objects of type T - T& operator^=(const T&) noexcept; + annotated_arg() noexcept; + annotated_arg(const T& v_, const PropertyListT &P = properties{}) noexcept; + template + annotated_arg(const T& v_, PropertyValueTs... props) noexcept; - // Available if the operator&= is valid for objects of type T - T& operator&=(const T&) noexcept; + template explicit annotated_arg( + const annotated_arg&) noexcept; + template + explicit annotated_arg(const annotated_arg&, + properties) noexcept; - // Available if the operator|= is valid for objects of type T - T& operator|=(const T&) noexcept; + annotated_arg(const annotated_arg&) = default; + annotated_arg& operator=(annotated_arg&) = default; - // Available if the operator<< is valid for objects of type T - T operator<<(const T&) noexcept; - const T operator<<(const T&) const noexcept; - - // Available if the operator>> is valid for objects of type T - T operator>>(const T&) noexcept; - const T operator>>(const T&) const noexcept; - - // Available if the operator>>= is valid for objects of type T - T& operator>>=(const T&) noexcept; - - // Available if the operator<<= is valid for objects of type T - T& operator<<=(const T&) noexcept; - - // Available if the operator== is valid for objects of type T - bool operator==(const T&) const noexcept; - - // Available if the operator!= is valid for objects of type T - bool operator!=(const T&) const noexcept; - - // Available if the operator<= is valid for objects of type T - bool operator<=(const T&) const noexcept; - - // Available if the operator>= is valid for objects of type T - bool operator>=(const T&) const noexcept; - - // Available if the operator&& is valid for objects of type T - bool operator&&(const T&) const noexcept; - - // Available if the operator|| is valid for objects of type T - bool operator||(const T&) const noexcept; - - // Available if the operator++ is valid for objects of type T - T& operator++() noexcept; - - // Available if the operator++ is valid for objects of type T - T operator++(int) noexcept; - - // Available if the operator-- is valid for objects of type T - T& operator--() noexcept; - - // Available if the operator-- is valid for objects of type T - T operator--(int) noexcept; - - // Available if the operator-> is valid for objects of type T - T& operator->() noexcept; - const T& operator->() const noexcept; - - // Available if the operator[] is valid for objects of type T - T& operator[](std::ptrdiff_t idx) noexcept; - const T& operator[](std::ptrdiff_t idx) const noexcept; + // Conversion operator to convert to the underlying type + operator T() noexcept; + operator T() const noexcept; - // Available if the operator() is valid for objects of type T - template auto operator()(Args... args) noexcept; - template auto operator()(Args... args) const noexcept; + // Available if the operator[] is valid for objects of type T, return + // type will match the return type of T::operator[](std::ptrdiff_t) + /* ... */ operator[](std::ptrdiff_t idx) const noexcept; template static constexpr bool has_property(); @@ -338,8 +220,17 @@ namespace sycl::ext::oneapi { // instances of propertyT template static constexpr /*unspecified*/ get_property(); - }; -}; // namespace sycl::ext::oneapi +}; + +//Deduction guides +template +annotated_arg(T, PropertyValueTs... values) -> + annotated_arg; + +template +annotated_arg(annotated_arg, PropertiesB>) -> + annotated_arg; +} // namespace sycl::ext::oneapi::experimental ---- [frame="topbot",options="header"] @@ -350,7 +241,7 @@ namespace sycl::ext::oneapi { a| [source,c++] ---- -annotated_arg(); +annotated_arg() noexcept; ---- | Not available in device code. Constructs an `annotated_arg` object which is default initialized. @@ -359,399 +250,92 @@ Constructs an `annotated_arg` object which is default initialized. a| [source,c++] ---- -annotated_arg(const T& v_); +annotated_arg(const T& v_, const PropertyListT &P = properties{}) noexcept; ---- | Not available in device code. Constructs an `annotated_arg` object from the input object `v_`. -// --- ROW BREAK --- -a| -[source,c++] ----- -operator T&() noexcept; -operator const T&() const noexcept; ----- -| Implicit conversion to a reference to the underlying type `T`. - -// --- ROW BREAK --- -a| -[source,c++] ----- -T operator+(const T&) noexcept; -const T operator+(const T&) const noexcept; ----- -| -Available if the `operator+(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T operator+() noexcept; -const T operator+() const noexcept; ----- -| -Available if the `operator+` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T operator-(const T&) noexcept; -const T operator-(const T&) const noexcept; ----- -| -Available if the `operator-(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T operator-() noexcept; -const T operator-() const noexcept; ----- -| -Available if the `operator-` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T operator*(const T&) noexcept; -const T operator*(const T&) const noexcept; ----- -| -Available if the `operator*(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T operator/(const T&) noexcept; -const T operator/(const T&) const noexcept; ----- -| -Available if the `operator/(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T operator%(const T&) noexcept; -const T operator%(const T&) const noexcept; ----- -| -Available if the `operator%(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T operator^(const T&) noexcept; -const T operator^(const T&) const noexcept; ----- -| -Available if the `operator^(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T operator\|(const T&) noexcept; -const T operator\|(const T&) const noexcept; ----- -| -Available if the `operator\|(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T operator&(const T&) noexcept; -const T operator&(const T&) const noexcept; ----- -| -Available if the `operator&(const T&)` is valid for objects of type `T` +The new property set `PropertyListT` contains all properties in `P`. // --- ROW BREAK --- a| [source,c++] ---- -T operator~() noexcept; -const T operator~() const noexcept; +template +annotated_arg(const T& v_, const PropertyValueTs... props) noexcept; ---- -| -Available if the `operator~` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T operator!() noexcept; -const T operator!() const noexcept; ----- -| -Available if the `operator!` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T& operator=(const T&) noexcept; ----- -| -Assignment from underlying type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -bool operator<(const T&) const noexcept; ----- -| -Available if the `operator<(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -bool operator>(const T&) const noexcept; ----- -| -Available if the `operator>(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T& operator+=(const T&) noexcept; ----- -| -Available if the `operator+=(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T& operator-=(const T&) noexcept; ----- -| -Available if the `operator-=(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T& operator*=(const T&) noexcept; ----- -| -Available if the `operator*=(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T& operator/=(const T&) noexcept; ----- -| -Available if the `operator/=(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T& operator%=(const T&) noexcept; ----- -| -Available if the `operator%=(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T& operator^=(const T&) noexcept; ----- -| -Available if the `operator^=(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T& operator&=(const T&) noexcept; ----- -| -Available if the `operator&=(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T& operator\|=(const T&) noexcept; ----- -| -Available if the `operator\|=(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T operator<<(const T&) noexcept; -const T operator<<(const T&) const noexcept; ----- -| -Available if the `operator<<(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T operator>>(const T&) noexcept; -const T operator>>(const T&) const noexcept; ----- -| -Available if the `operator>>(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T& operator>>=(const T&) noexcept; ----- -| -Available if the `operator>>=(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T& operator<<=(const T&) noexcept; ----- -| -Available if the `operator<<=(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -bool operator==(const T&) const noexcept; ----- -| -Available if the `operator==(const T&)` is valid for objects of type `T` +| Not available in device code. +Constructs an `annotated_arg` object from the input object `v_`. -// --- ROW BREAK --- -a| -[source,c++] ----- -bool operator!=(const T&) const noexcept; ----- -| -Available if the `operator!=(const T&)` is valid for objects of type `T` +The new property set `PropertyListT` contains all properties listed in `props`. // --- ROW BREAK --- a| [source,c++] ---- -bool operator<=(const T&) const noexcept; +template explicit annotated_arg( + const annotated_arg &ConvertFrom); ---- -| -Available if the `operator<=(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -bool operator>=(const T&) const noexcept; ----- -| -Available if the `operator>=(const T&)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -bool operator&&(const T&) const noexcept; ----- -| -Available if the `operator&&(const T&)` is valid for objects of type `T` +| Not available in device code. +Constructs the `annotated_arg` object from the `ConvertFrom` object if +the list of properties in `PropertyListT` is a superset of the list of +properties in `P`. -// --- ROW BREAK --- -a| -[source,c++] ----- -bool operator\|\|(const T&) const noexcept; ----- -| -Available if the `operator\|\|(const T&)` is valid for objects of type `T` +`T2` must be implicitly convertible to `T`. // --- ROW BREAK --- a| [source,c++] ---- -T& operator++() noexcept; +template +explicit annotated_arg(const annotated_arg& v_, + properties P) noexcept; ---- -| -Available if the `operator++` is valid for objects of type `T` +| Not available in device code. +Constructs the `annotated_arg` object from the input object `v_`. -// --- ROW BREAK --- -a| -[source,c++] ----- -T operator++(int) noexcept; ----- -| -Available if the `operator++(int)` is valid for objects of type `T` +The new `PropertyListT` is the union of all properties contained within +`PropertyListU` and `PropertyListV`. If there are any common properties in the +two lists with different values, a compile-time error is triggered. Common +properties with the same value (or no value) are allowed. -// --- ROW BREAK --- -a| -[source,c++] ----- -T& operator--() noexcept; ----- -| -Available if the `operator--` is valid for objects of type `T` +`T2` must be implicitly convertible to `T`. // --- ROW BREAK --- a| [source,c++] ---- -T operator--(int) noexcept; +annotated_arg(const annotated_arg&) = default; ---- -| -Available if the `operator--(int)` is valid for objects of type `T` +| Not available in device code. +Compiler generated copy constructor. // --- ROW BREAK --- a| [source,c++] ---- -T& operator->() noexcept; -const T& operator->() const noexcept; +annotated_arg& operator=(annotated_arg&) = default; ---- -| -Available if the `operator->` is valid for objects of type `T` - -Provides member access through `T` that is a pointer or a class which defines -`operator->`. +| Compiler generated assignment operator. // --- ROW BREAK --- a| [source,c++] ---- -T& operator[](std::ptrdiff_t idx) noexcept; -const T& operator[](std::ptrdiff_t idx) const noexcept; +operator T() noexcept; +operator const T() const noexcept; ---- -| -Available if the `operator[]` is valid for objects of type `T` +| Implicit conversion to a reference to the underlying type `T`. // --- ROW BREAK --- a| [source,c++] ---- -template auto operator()(Args... args) noexcept; -template auto operator()(Args... args) const noexcept; +/* ... */ operator[](std::ptrdiff_t idx) const noexcept; ---- | -Available if the `operator()` is valid for objects of type `T` +Available if the `operator[]` is valid for objects of type `T`. This function +will call the subscript operator defined for `T`. // --- ROW BREAK --- a| @@ -761,7 +345,7 @@ template static constexpr bool has_property(); ---- | -Returns true if the property list contains the property with property key class +Returns true if `PropertyListT` contains the property with property key class `propertyT`. Returns false if it does not. Available only when `propertyT` is a property key class. @@ -774,22 +358,13 @@ template static constexpr /* unspecified */ get_property(); ---- | -Returns a copy of the property value contained in the property list -`PropertyListT`. Must produce a compile error if `PropertyListT` does not -contain a property with the `propertyT` key. +Returns a copy of the property value contained in `PropertyListT`. Must produce +a compile error if `PropertyListT` does not contain a property with the +`propertyT` key. Available only if `propertyT` is the property key class of a compile-time constant property. -// --- ROW BREAK --- -a| -[source,c++] ----- -~annotated_arg(); ----- -| -Compiler supplied destructor function. - |=== == Issues @@ -805,12 +380,3 @@ None. |Rev|Date|Author|Changes |1|2022-03-09|Abhishek Tiwari|*Initial working 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. -//************************************************************************ diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc new file mode 100644 index 0000000000000..3fe61ef389234 --- /dev/null +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -0,0 +1,704 @@ += sycl_ext_oneapi_annotated_ptr + +: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 + +[%hardbreaks] +Copyright (C) 2022 Intel Corporation. All rights reserved. + +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 + +Gregory Lueck, Intel + +Jason Sewall, Intel + +Jessica Davies, Intel + +Joe Garvey, Intel + +John Pennycook, Intel + +Michael Kinsner, Intel + +Roland Schulz, Intel + +== Dependencies + +This extension is written against the SYCL 2020 revision 5 specification. All +references below to the "core SYCL specification" or to section numbers in the +SYCL specification refer to that revision. + +This extension also depends on the following other SYCL extensions: + +* link:../experimental/sycl_ext_oneapi_properties.asciidoc[sycl_ext_oneapi_properties] +extension. + +== 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 a pointer wrapper class that provides a mechanism to +attach compile-time constant information to a pointer in a manner that allows +the compiler to reliably maintain and analyze the information. + +[NOTE] +==== +In this document, we use the shortened form `annotated_ptr` to refer to the +proposed `sycl::ext::oneapi::experimental::annotated_ptr` class. +==== + +The purpose of this document is to clearly describe and specify `annotated_ptr` +and related concepts, types, and mechanisms, and to give examples and context +for their usage. + +== Specification + +=== Feature test macro + +This extension provides a feature-test macro as described in the core SYCL +specification, Section 6.3.3 "Feature test macros". Therefore, an +implementation supporting this extension must predefine the macro +`SYCL_EXT_ONEAPI_ANNOTATED_PTR` 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 +that the implementation supports. + +[%header,cols="1,5"] +|=== +|Value |Description +|1 |Initial extension version +|=== + +=== Motivation + +When compiling a program with pointer operations, it is often desirable to +attach certain information such as optimization hints or functional directives +onto the pointer itself. + +Any acceptable mechanism to attach such information should meet the following +requirements: + + 1. It should not be burdensome for users to use the mechanism. + 2. It should be reasonably scalable and not add significant burden on compiler + developers to maintain the mechanism. + 3. It should preserve the attached information through most uses of the + pointer. + +For example, consider the following case: We want to specify some compile-time +constant information `I` for a pointer. One possible method to do so is for the +compiler to provide a special function for reading to and writing from the +pointer and have this function convey `I` to the compiler at the call site where +the read or write occurs. + +That is, users would convert code that looks like: +```c++ +{ + SomeType* a; + ... + *a = some_val; + ... + read_val = *a +} +``` + +To code that looks like: +```c++ +{ + SomeType* a; + ... + __special_store_with_annotation(a, some_val, + I /*some compile time constant information*/); + ... + __special_load_with_annotation(a, &read_val, + I /*some compile time constant information*/); +} +``` + +This mechanism does not meet requirement (1) listed above because users have to +replace all their pointer read/write code with special function calls. + +Another mechanism could be that the compiler provides attributes which can be +applied to the pointer declaration to convey some compile-time constant +information. In this case users would change their code to: +```cpp +{ + // 'I' is the compiler time constant information being annotated on 'a' + [[special_annotations::use_case_x_attribute(I)]] SomeType* a; + ... + *a = some_val; + ... + read_val = *a +} +``` +However, this mechanism fails to meet requirements (2) and (3) listed above +since: + + * Adding new attributes can add non-trivial attribute creation overheads for + compiler developers. With growing set of attributes, there is an associated + burden of maintaining them and of preparing diagnostic checks. + * Compilers are not required to preserve information propagated as attributes. + Hence this method is only suitable for hints, and not functional directives. + Further, from a given load/store call a compiler may not be able to + successfully trace back to the declaration to deduce which annotation was + applied to it. + +The `annotated_ptr` class described in this document is a class template that +encapsulates a pointer. The template accepts a list of compile-time constant +properties. The implementation of the class defined here should preserve the +information provided as compile-time constant properties through all uses of the +pointer unless noted otherwise. + +.Unsupported Usage Example +[source,c++] +---- +using sycl::ext::oneapi::experimental; +struct MyType { + annotated_ptr> a; + annotated_ptr> b; +}; + +struct MyKernel { + MyType arg_a; + MyType arg_b; + ... + void operator()() const { + ... + } +}; +---- +It is illegal to apply `annotated_ptr` to members of kernel arguments. In the +above example, encapsulating `annotated_ptr` within `MyType` is illegal. + +=== Representation of `annotated_ptr` + +`annotated_ptr` is a class template, parameterized by the type of the underlying +allocation `T`, and a list of properties `PropertyListT`. + +[source,c++] +---- +namespace sycl::ext::oneapi::experimental { +template > +class annotated_ptr { + ... +---- + +`PropertyListT` enables properties to be associated with an `annotated_ptr`. +Properties may be specified for an `annotated_ptr` to provide semantic +modification or optimization hint information. Please note that there cannot +be duplicated property in a properties list. Otherwise, a compiler time error is triggered. + +Here's an example of how a property could be used: + +[source,c++] +---- +using namespace sycl::ext::oneapi::experimental; +{ + sycl::queue q; + int* ptr = ... + // alignment of the pointer in bytes specified using the property 'alignment' + auto arg_a = annotated_ptr(ptr, properties{alignment<4>}); + + // alignment in bytes and address bus width of the pointer specified using the + // properties 'alignment' and 'awidth'. + // a properties object is deduced from the list the of property values + auto arg_b = annotated_ptr(ptr, alignment<4>, awidth<32>); + + q.submit([=]{ + ... + *arg_a = (*arg_a) * 2; + }); + ... +} +---- + +The section below and the table that follows, describe the constructors and +member functions for `annotated_ptr`. + +NOTE: `annotated_ptr` is a trivially copyable type. + +The section below refers to an `annotated_ref` class which is described in the +section following this one. + +[source,c++] +---- +namespace sycl::ext::oneapi::experimental { +template > +class annotated_ptr { + public: + using reference = annotated_ref; + + annotated_ptr() noexcept; + annotated_ptr(T *Ptr, const PropertyListT &P = properties{}) noexcept; + template + annotated_ptr(T *Ptr, PropertyValueTs... props) noexcept; + + template explicit annotated_ptr( + const annotated_ptr&) noexcept; + template + explicit annotated_ptr(const annotated_ptr&, + properties) noexcept; + + annotated_ptr(const annotated_ptr&) noexcept = default; + + reference operator*() const noexcept; + + reference operator[](std::ptrdiff_t) const noexcept; + + annotated_ptr operator+(size_t) const noexcept; + std::ptrdiff_t operator-(annotated_ptr) const noexcept; + + explicit operator bool() const noexcept; + + // Implicit conversion is not supported + operator T*() noexcept = delete; + operator T*() const noexcept = delete; + + T* get() const noexcept; + + annotated_ptr& operator=(T*) noexcept; + annotated_ptr& operator=(const annotated_ptr&) noexcept = default; + + annotated_ptr& operator++() noexcept; + annotated_ptr operator++(int) noexcept; + + annotated_ptr& operator--() noexcept; + annotated_ptr operator--(int) noexcept; + + template + static constexpr bool has_property(); + + // The return type is an unspecified internal class used to represent + // instances of propertyT + template + static constexpr /*unspecified*/ get_property(); + }; + +//Deduction guides +template +annotated_ptr(T, PropertyValueTs... values) -> + annotated_ptr; + +template +annotated_ptr(annotated_ptr, PropertiesB>) -> + annotated_ptr; +} // namespace sycl::ext::oneapi::experimental +---- + +[frame="topbot",options="header"] +|=== +|Functions |Description + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ptr() noexcept; +---- +| +Constructs an `annotated_ptr` object. Does not allocate new storage. The +underlying pointer is initialized to `nullptr`. + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ptr(T *Ptr, const PropertyListT &P = properties{}) noexcept; +---- +| +Constructs an `annotated_ptr` object. Does not allocate new storage. The +underlying pointer is initialized with `Ptr`. + +The new property set `PropertyListT` contains all properties in `P`. + +// --- ROW BREAK --- +a| +[source,c++] +---- +template +annotated_ptr(T *Ptr, const PropertyValueTs... props) noexcept; +---- +| +Constructs an `annotated_ptr` object. Does not allocate new storage. The +underlying pointer is initialized with `Ptr`. + +The new property set `PropertyListT` contains all properties in `P`. + +// --- ROW BREAK --- +a| +[source,c++] +---- +template explicit annotated_ptr( + const annotated_ptr &ConvertFrom); +---- +| +Constructs the `annotated_ptr` object from the `ConvertFrom` object if +the list of properties in `PropertyListT` is a superset of the list of +properties in `P`. + +`T2*` must be implicitly convertible to `T*`. + +// --- ROW BREAK --- +a| +[source,c++] +---- +template +explicit annotated_ptr(const annotated_ptr &Ptr, + properties P) noexcept; +---- +| +Constructs an `annotated_ptr` object. Does not allocate new storage. The +underlying pointer is initialized with `Ptr`. + +The new `PropertyListT` is the union of all properties contained within +`PropertyListU` and `PropertyListV`. If there are any common properties +in the two lists with different values, a compile-time error is triggered. +Common properties with the same value (or no value) are allowed. + +`T2*` must be implicitly convertible to `T*`. + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ptr(const annotated_ptr &) noexcept = default; +---- +| +Constructs an `annotated_ptr` object from another `annotated_ptr` with the same +type and properties. + +// --- ROW BREAK --- +a| +[source,c++] +---- +reference operator*() const; +---- +| +Returns an `annotated_ref` reference wrapper which can be used to read or write +to the underlying pointer. Reads/Writes using the reference will retain the +annotations. + +// --- ROW BREAK --- +a| +[source,c++] +---- +reference operator[](std::ptrdiff_t Index) const; +---- +| +Returns an `annotated_ref` reference wrapper to the object at offset `Index`. + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ptr operator+(size_t Offset) const; +---- +| +Returns an `annotated_ptr` that points to `this[Offset]`. + +// --- ROW BREAK --- +a| +[source,c++] +---- +std::ptrdiff_t operator-(annotated_ptr FromPtr) const; +---- +| +Returns the distance between the underlying pointer and the pointer encapsulated +by `FromPtr`. + +// --- ROW BREAK --- +a| +[source,c++] +---- +explicit operator bool() const noexcept; +---- +| +Returns `false` if the underlying pointer is null, returns `true` otherwise. + +// --- ROW BREAK --- +a| +[source,c++] +---- +/*unspecified*/ operator T*() noexcept = delete; +/*unspecified*/ operator const T*() const noexcept = delete; +---- +| +Implicit conversion to a pointer to the underlying type is not supported. + +// --- ROW BREAK --- +a| +[source,c++] +---- +T* get() const noexcept; +---- +| +Returns the underlying raw pointer. The raw pointer will not retain the +annotations. + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ptr& operator=(T*) noexcept; +---- +| +Allows assignment from a pointer to type `T`. + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ptr& operator=(const annotated_ptr &) noexcept = default; +---- +| +Allows assignment from an `annotated_ptr` with the same parameterization. + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ptr& operator++() noexcept; +---- +| +Prefix increment operator. + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ptr operator++(int) noexcept; +---- +| +Postfix increment operator. + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ptr& operator--() noexcept; +---- +| +Prefix decrement operator. + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ptr operator--(int) noexcept; +---- +| +Postfix decrement operator. + +// --- ROW BREAK --- +a| +[source,c++] +---- +template +static constexpr bool has_property(); +---- +| +Returns true if `PropertyListT` contains the property with property key class +`propertyT`. Returns false if it does not. + +Available only when `propertyT` is a property key class. + +// --- ROW BREAK --- +a| +[source,c++] +---- +template +static constexpr auto get_property(); +---- +| +Returns a copy of the property value contained in `PropertyListT`. Must produce +a compile error if `PropertyListT` does not contain a property with the +`propertyT` key. + +Available only if `propertyT` is the property key class of a compile-time +constant property. +|=== + +=== Add new reference wrapper class `annotated_ref` to enable `annotated_ptr` + +The purpose of the `annotated_ref` class template is to provide reference +wrapper semantics. It enables properties to be preserved on loads from and +stores to annotated_ptrs. + +```c++ +namespace sycl::ext::oneapi::experimental { +template > +class annotated_ref { + public: + annotated_ref(const annotated_ref&) = default; + operator T() const; + annotated_ref& operator=(const T&) const; + annotated_ref& operator=(const annotated_ref&) const = default; + }; +} // namespace sycl::ext::oneapi::experimental +``` + + +Member Functions are described in the table below +[frame="topbot",options="header"] +|=== +|Functions |Description + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ref(const annotated_ref&) = default; +---- +| +Compiler generated copy constructor. + +// --- ROW BREAK --- +a| +[source,c++] +---- +operator T() const; +---- +| +Reads the object of type `T` that is referenced by this wrapper, applying the +annotations when the object is loaded from memory. + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ref& operator=(const T &) const; +---- +| +Writes an object of type `T` to the location referenced by this wrapper, +applying the annotations when the object is stored to memory. + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ref& operator=(const annotated_ref&) const = default; +---- +| +Assign from another `annotated_ref` object. + +|=== + +=== Properties + +Below is a list of compile-time constant properties supported with +`annotated_ptr`. Other extensions can define additional compile-time constant +properties that can be supported with `annotated_ptr`. Runtime properties +are not supported. + +```c++ +namespace sycl::ext::oneapi::experimental { +struct alignment_key { + template + using value_t = property_value>; +}; + +template +inline constexpr alignment_key::value_t alignment; + +template<> +struct is_property_key : std::true_type {}; + +template +struct is_property_key_of< + alignment_key, annotated_ptr> : std::true_type {}; +template +} // namespace sycl::ext::oneapi::experimental +``` +-- +[options="header"] +|==== +| Property | Description +|`alignment` +| +This property is an assertion by the application that the `annotated_ptr` has +the given alignment, specified in bytes. The behavior is undefined if the +pointer value does not have the indicated alignment. +|==== +-- + +== Issues related to `annotated_ptr` + +1) [RESOLVED] Should we allow implicit conversion to base class by default? + +No, implicit conversion will not be allowed. + +2) [RESOLVED] How do we support `operator->`? + +We will not support `operator->` with the initial release, since we do +not have meaningful usecases that require this support. Building the support +is complicated + +3) [RESOLVED] Can `sycl::atomic_ref` be used with `annotated_ref`? + +`atomic_ref` will not work with `annotated_ref` as is since `atomic_ref` +restricts the types it can take. If we want, we can create a sycl extension for +`atomic_ref` to support `annotated_ref`. The implementation complexity will +depend on how we chose to implement `annotated_ref`: via builtins or via pointer +annotations. + +4) [RESOLVED] Should we provide conversion functions to convert to/from +multi_ptr? + +No we do not want to support multi_ptr conversion. 'multi_ptr's provide +a way to annotate address spaces. That can be built with annotated_ptr. + +5) We need a property to capture local, global or private address-spaces. Within +global space we may want to distinguish between general, host, and device memory +spaces. + +6) [RESOLVED] Ctor should clarify whether when constructing the object from two +different property lists, duplicates will exist or not. + +Updated ctor description with the resolution. + +7) Add `runtime_aligned` property back to this core spec once a way to query +the alignment is set up. + +== Revision History + +[cols="5,15,15,70"] +[grid="rows"] +[options="header"] +|======================================== +|Rev|Date|Author|Changes +|3|2022-04-05|Abhishek Tiwari|*Addressed review comments* +|2|2022-03-07|Abhishek Tiwari|*Corrected API and updated description* +|1|2021-11-01|Abhishek Tiwari|*Initial internal review version* +|======================================== diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_properties.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_properties.asciidoc new file mode 100644 index 0000000000000..b33fe8b559b21 --- /dev/null +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_properties.asciidoc @@ -0,0 +1,189 @@ += sycl_ext_oneapi_kernel_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 + +[%hardbreaks] +Copyright (C) 2022 Intel Corporation. All rights reserved. + +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 5. + +It depends on the following extensions: + + - link:../experimental/sycl_ext_oneapi_properties.asciidoc[sycl_ext_oneapi_properties] + - link:sycl_ext_oneapi_annotated_ptr.asciidoc[sycl_ext_oneapi_annotated_ptr] + - 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 that are applied to kernel arguments by +using the `sycl::ext::oneapi::experimental::annotated_ptr` and +`sycl::ext::oneapi::experimental::annotated_arg` classes. + +== 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_ONEAPI_KERNEL_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. +|=== + +=== `restrict` property + +The `restrict` property defined here is only meaningful on the kernel arguments +when the kernel argument is a pointer type. It is ignored for other types. + +This property is not meaningful within the kernel body. + + +```c++ +namespace sycl::ext::oneapi::experimental { +struct restrict_key { + using value_t = property_value; +}; + +inline constexpr restrict_key::value_t restrict; + +template<> +struct is_property_key : std::true_type {}; + +template +struct is_property_key_of< + restrict_key, annotated_ptr> : std::true_type {}; + +template +struct is_property_key_of< + restrict_key, annotated_arg> : std::true_type {}; +} // namespace sycl::ext::oneapi::experimental +``` +-- + +[frame="topbot",options="header"] +|=== +|Property |Description + +a| +[source,c++] +---- +restrict +---- +a| +This is an assertion by the application that the pointer kernel arguments marked +with this property do not alias with one another with the same semantics as the +C99 `restrict` keyword. The behavior is undefined if these pointer values do +alias. + +|=== +-- + +=== Usage Examples + +The example below shows a simple kernel with one +`sycl::ext::oneapi::experimental::annotated_ptr` kernel argument and one +`sycl::ext::oneapi::experimental::annotated_arg` kernel argument. + +.Usage Example +```c++ +using sycl::ext::oneapi::experimental; +{ + sycl::queue q{...}; + + // Allocate memory + auto ptr_a = ... + int* ptr_b = ...; + + // Add properties + auto arg_a = annotated_ptr(ptr_a, properties{restrict}); + auto arg_n = annotated_arg(ptr_b, properties{restrict}); + ... + + q.single_task([=] { + for (int i=0; i< kSize; i++) + arg_a[i] = arg_a[i] + arg_n[i]; + }).wait(); + + ... +} +``` + +== Issues + +None + +== Revision History + +[cols="5,15,15,70"] +[grid="rows"] +[options="header"] +|======================================== +|Rev|Date |Author |Changes +|1 |2022-07-1 |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. +//************************************************************************