From 64217f6246729de89ef4ef46150929392691902f Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Tue, 8 Mar 2022 07:40:30 -0800 Subject: [PATCH 01/55] post draft for review --- .../sycl_ext_oneapi_annotated_ptr.asciidoc | 510 ++++++++++++++++++ ...t_oneapi_annotated_ptr_properties.asciidoc | 158 ++++++ 2 files changed, 668 insertions(+) create mode 100755 sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc create mode 100755 sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr_properties.asciidoc 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 100755 index 0000000000000..eaccb930c2ab9 --- /dev/null +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -0,0 +1,510 @@ += 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++ + +== Introduction +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. + +== 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. + +== Status + +Working Draft + +This is a preview extension specification, intended to provide early access to +a feature for review and community feedback. When the feature matures, this +specification may be released as a formal extension. + +Because the interfaces defined by this specification are not final and are +subject to change they are not intended to be used by shipping software +products. + +== Version + +Built On: {docdate} + +Revision: 2 + +== Contact + +Abhishek Tiwari, Intel (abhishek2 'dot' tiwari 'at' intel 'dot' com) + +== Contributors + +Abhishek Tiwari, Intel + +Greg 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 specification, revision 4. + +It also depends on the `sycl_ext_oneapi_properties` extension. + +== Overview + +[NOTE] +==== +In this document, we use the shortened form `annotated_ptr` to refer to the +proposed `sycl::ext::oneapi::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. + +=== 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 requirements (1) and (2) listed above because: + + * Users have to replace all their pointer read/write code with special function calls. + * If the set or type of information represented by `I` evolves over time, it requires creating and/or updating these special functions and the underlying compiler implementation. + +Another mechanism could be that the compiler provide 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 +{ + [[special_annotations::use_case_x_attribute(I)]] SomeType* a; // 'I' is the compiler time constant information being annotated on 'a' + ... + *a = some_val; + ... + read_val = *a +} +``` +However, this mechanism fails to meet requirements (2) and (3) listed above since: + + * Adding new attributes for 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, the compiler may not be able to observe a particular pointer load/store call and successfully trace back to the declaration to deduce which annotation was applied to it. + +The `annotated_ptr` class described in this document is a templated type 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. + +=== 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 +|=== + +=== 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 { +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. + +Example uses of a property are: + +[source,c++] +---- +using namespace sycl::ext::oneapi; +{ + sycl::queue q; + annotated_ptr>> kernel_arg; // alignment of the pointer in bytes specified using the property 'align' + q.submit([=]{ + *kernel_arg = *kernel_arg * 2; + }); + ... +} +---- + +The section below and the table that follows, describe the constructors, member functions and factory methods for `annotated_ptr`. +The section below refers to an `annotated_ref` class which is described in the section following this one. + +[source,c++] +---- +namespace sycl::ext::oneapi { +template > +class annotated_ptr { + public: + using reference = annotated_ref; + + annotated_ptr() noexcept = default; + explicit annotated_ptr(T *Ptr) noexcept = default; + annotated_ptr(annotated_ptr const &) noexcept = default; + + // Conversion function + template explicit annotated_ptr(annotated_ptr const &) noexcept; + + reference operator*() const noexcept; + reference operator[](std::ptrdiff_t) const noexcept; + annotated_ptr operator+(unsigned long long) const noexcept; + std::ptrdiff_t operator-(annotated_ptr) const noexcept; + + operator bool() const noexcept; + + // Implicit conversion is not supported + operator T*() noexcept = delete; + operator const T*() const noexcept = delete; + + T* get() noexcept; + const T* get() const noexcept; + + annotated_ptr& operator=(const T*) noexcept; + annotated_ptr& operator=(annotated_ptr const&) noexcept; + + 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(); + + private: + T *Ptr; + }; +} // namespace sycl::ext::oneapi + +---- + +[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++] +---- +explicit annotated_ptr(T *Ptr); +---- +| +Constructs an `annotated_ptr` object. Does not allocate new storage. The underlying pointer is initialized with `Ptr`. + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ptr(annotated_ptr const &) noexcept = default; +---- +| +Constructs an `annotated_ptr` object from another `annotated_ptr` with the same template parameterization object. + +// --- ROW BREAK --- +a| +[source,c++] +---- +template explicit annotated_ptr(annotated_ptr const & ConvertFrom); +---- +| +Constructs the `annotated_ptr` object from the `ConvertFrom` object if the template parameter types are compatible. For example, if the properties on one of the operands conflict with the others then the compiler should issue a compile time error. + +// --- ROW BREAK --- +a| +[source,c++] +---- +reference operator*() const; +---- +| +Returns a 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+(unsigned long long Offset) const; +---- +| +Returns an `annotated_ptr` that points to the location `Offset` distance away from the underlying pointer. + +// --- 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++] +---- +operator bool() const noexcept; +---- +| +Returns `false` if the underlying pointer is null, returns `true` otherwise. + +// --- ROW BREAK --- +a| +[source,c++] +---- +operator T*() noexcept = delete; +operator const T*() const noexcept = delete; +---- +| +Implicit conversion to a pointer to the underlying type `T` is not supported. + +// --- ROW BREAK --- +a| +[source,c++] +---- +T* get() noexcept; +const 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=(const T*) noexcept; +---- +| +Allows assignment from an pointer to type `T`. + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ptr& operator=(annotated_ptr const&) noexcept; +---- +| +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++() noexcept; +---- +| +Postfix increment operator. + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ptr& operator--() noexcept; +---- +| +Prefix decrement operator. + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ptr operator--() noexcept; +---- +| +Postfix decrement operator. + +// --- ROW BREAK --- +a| +[source,c++] +---- +template +static constexpr bool has_property(); +---- +| Returns true if the `PropertyListT` contains the property specified by `propertyT`. Returns false if it does not. +Available only if `sycl::is_property_of_v` is true. + +// --- ROW BREAK --- +a| +[source,c++] +---- +template +static constexpr auto get_property(); +---- +| Returns an object of the class used to represent the value of property `propertyT`. +Must produce a compiler diagnostic if `PropertyListT` does not contain a `propertyT` property. +Available only if `sycl::is_property_of_v` is true. +|=== + +=== 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 the implementation to preserve the properties on loads from and stores to the pointers. + +```c++ +namespace sycl::ext::oneapi { +template > +class annotated_ref { + public: + annotated_ref(T *); + operator T() noexcept; + operator const T() const noexcept; + void operator=(const T &); + private: + T *Ptr; + }; +} // namespace sycl::ext::oneapi +``` + + +Member Functions are described in the table below +[frame="topbot",options="header"] +|=== +|Functions |Description + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ref(T * InputPtr); +---- +| +Constructs an `annotated_ref` object. Does not allocate new storage. The underlying pointer is initialized to `InputPtr`. + +// --- ROW BREAK --- +a| +[source,c++] +---- +operator T() noexcept; +operator const T() const noexcept; +---- +| +Implicit conversion to underlying type. + +// --- ROW BREAK --- +a| +[source,c++] +---- +void operator=(const T &); +---- +| +Enables assignment to the underlying pointer. + +|=== + +=== Properties for `annotated_ptr` variables + +See the extension `sycl_ext_oneapi_annotated_ptr_properties` for the full list of supported properties. + +== Issues + +1) [RESOLVED] Should we allow implicit conversion to base class by default? +*Ans: No.* + +2) [RESOLVED] How do we support `operator->`? +*Ans: Not with the initial release.* + +3) [RESOLVED] Can `sycl::atomic_ref` be used with `annotated_ref`? +*Ans: Yes. This discussion is an implementation detail discussion and does not impact the spec.* + +4) [RESOLVED] Should we provide conversion functions to convert to/from multi_ptr? +*Ans: No.* + +== Revision History + +[cols="5,15,15,70"] +[grid="rows"] +[options="header"] +|======================================== +|Rev|Date|Author|Changes +|2|2022-03-07|Abhishek Tiwari|*Corrected API and updated description* +|1|2021-11-01|Abhishek Tiwari|*Initial internal review version* +|======================================== \ No newline at end of file diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr_properties.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr_properties.asciidoc new file mode 100755 index 0000000000000..ec8108b34e163 --- /dev/null +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr_properties.asciidoc @@ -0,0 +1,158 @@ += sycl_ext_oneapi_annotated_ptr_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++ + +== Introduction +This extension defines the properties supported with the class `annotated_ptr` +which is defined by the extension `sycl_ext_oneapi_annotated_ptr`. These +properties are carried by the `annotated_ptr` class and implemented by the +toolchain. + +== 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. + +== Status + +Final Draft + +This is a preview extension specification, intended to provide early access to +a feature for review and community feedback. When the feature matures, this +specification may be released as a formal extension. + +Because the interfaces defined by this specification are not final and are +subject to change they are not intended to be used by shipping software +products. + +== Version + +Built On: {docdate} + +Revision: 1 + +== Contact + +Abhishek Tiwari, Intel (abhishek2 'dot' tiwari 'at' intel 'dot' com) + +== 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: + - sycl_ext_oneapi_annotated_ptr + - sycl_ext_intel_buffer_location + +== Overview + +This extension introduces properties that establish the annotations that will be +carried by the class `sycl::ext::oneapi::annotated_ptr`. + +An example of the syntax can be seen below. + +[source,c++] +---- +annotated_ptr>> aptr; +---- + +== 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_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 APIs the implementation supports. + +[%header,cols="1,5"] +|=== +|Value |Description +|1 |Initial extension version. Base features are supported. +|=== + + +=== New Section 4.7.9.1 annotated_ptr properties + +Below is a list of compile-time-constant properties supported with `annotated_ptr`. + +```c++ +namespace sycl::ext::oneapi::experimental { +struct align_key { + template + using value_t = property_value>; +}; + +template +inline constexpr align::value_t align; + +template<> +struct is_property_key : std::true_type {}; + +template +struct is_property_key_of : std::true_type {}; + +} // namespace experimental::oneapi::ext::sycl +``` +-- +[options="header"] +|==== +| Property | Description +|`align` +| The alignment of the pointer address in bytes. + +| Property | Description +|`buffer_location` +| The alignment of the pointer address in bytes. +See the extension `sycl_ext_intel_buffer_location` for more details. +|==== +-- + +== Issues + +== Revision History + +[cols="5,15,15,70"] +[grid="rows"] +[options="header"] +|======================================== +|Rev|Date|Author|Changes +|1|2022-02-17|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. +//************************************************************************ From c3499e5b527af8304e458421f4def57573a8cdad Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Tue, 8 Mar 2022 07:57:53 -0800 Subject: [PATCH 02/55] edit author name --- .../extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index eaccb930c2ab9..cf15e9e92a154 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -58,7 +58,7 @@ Abhishek Tiwari, Intel (abhishek2 'dot' tiwari 'at' intel 'dot' com) == Contributors Abhishek Tiwari, Intel + -Greg Lueck, Intel + +Gregory Lueck, Intel + Jason Sewall, Intel + Jessica Davies, Intel + Joe Garvey, Intel + From 22b2eb1752c419fd9243aade4b1df7de1abeabce Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Wed, 13 Apr 2022 00:15:02 -0700 Subject: [PATCH 03/55] address review comments, add fpga specific properties --- ...tel_fpga_annotated_ptr_properties.asciidoc | 616 ++++++++++++++++++ .../sycl_ext_oneapi_annotated_ptr.asciidoc | 356 +++++++--- ...t_oneapi_annotated_ptr_properties.asciidoc | 158 ----- 3 files changed, 868 insertions(+), 262 deletions(-) create mode 100755 sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_ptr_properties.asciidoc delete mode 100755 sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr_properties.asciidoc diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_ptr_properties.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_ptr_properties.asciidoc new file mode 100755 index 0000000000000..52aa8eb808e86 --- /dev/null +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_ptr_properties.asciidoc @@ -0,0 +1,616 @@ += sycl_ext_intel_fpga_annotated_ptr_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_ptr.asciidoc[sycl_ext_oneapi_annotated_ptr] + +== 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_ptr`. 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_ptr, kernel_arg_awidth<32>, kernel_arg_dwidth<>64>> ptr_a; +---- + +.Example 2 +[source,c++] +---- +auto ptr_p = malloc_shared(1, q, properties{mem_id<2>}); +auto ptr_q = annotated_ptr(ptr_a, properties{ + kernel_arg_awidth<32>, kernel_arg_dwidth<64>}); +---- + +== 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_PTR_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. +|=== + +=== Kernel Argument Interface Properties + +Below is a list of compile-time constant properties supported with +`annotated_ptr`. These properties control the kernel argument interface on FPGA +devices. + +```c++ +namespace sycl::ext::oneapi::experimental { +struct kernel_arg_mm_host_key { + using value_t = property_value; +}; + +struct kernel_arg_agent_register_key { + using value_t = property_value; +}; + +struct kernel_arg_stable_key { + using value_t = property_value; +}; + +struct mem_id_key { + template + using value_t = property_value>; +}; + +struct kernel_arg_awidth_key { + template + using value_t = property_value< + kernel_arg_awidth_key, std::integral_constant>; +}; + +struct kernel_arg_dwidth_key { + template + using value_t = property_value< + kernel_arg_dwidth_key, std::integral_constant>; +}; + +enum class kernel_arg_read_write_mode_enum { + read_write, + read, + write +}; + +struct kernel_arg_read_write_mode_key { + template + using value_t = property_value>; +}; + +struct kernel_arg_latency_key { + template + using value_t = property_value>; +}; + +struct kernel_arg_maxburst_key { + template + using value_t = property_value< + kernel_arg_maxburst_key, std::integral_constant>; +}; + +struct kernel_arg_wait_request_key { + template + using value_t = property_value>; +}; + +inline constexpr kernel_arg_mm_host_key::value_t kernel_arg_mm_host; +inline constexpr kernel_arg_agent_register_key::value_t + kernel_arg_agent_register; +inline constexpr kernel_arg_stable_key::value_t kernel_arg_stable; +template inline constexpr mem_id_key::value_t mem_id; +template inline constexpr kernel_arg_awidth_key::value_t + kernel_arg_awidth; +template inline constexpr kernel_arg_dwidth_key::value_t + kernel_arg_dwidth; +template +inline constexpr kernel_arg_read_write_mode_key::value_t + kernel_arg_read_write_mode; +inline constexpr kernel_arg_read_write_mode_key::value_t< + kernel_arg_read_write_mode_enum::read> kernel_arg_read_write_mode_read; +inline constexpr kernel_arg_read_write_mode_key::value_t< + kernel_arg_read_write_mode_enum::write> kernel_arg_read_write_mode_write; +inline constexpr kernel_arg_read_write_mode_key::value_t< + kernel_arg_read_write_mode_enum::read_write> + kernel_arg_read_write_mode_readwrite; +template inline constexpr kernel_arg_latency_key::value_t + kernel_arg_latency; +template inline constexpr kernel_arg_maxburst_key::value_t + kernel_arg_maxburst; +template inline constexpr kernel_arg_wait_request_key::value_t + kernel_arg_wait_request; +inline constexpr kernel_arg_wait_request_key::value_t + kernel_arg_wait_request_requested; +inline constexpr kernel_arg_wait_request_key::value_t + kernel_arg_wait_request_not_requested; + +template<> struct is_property_key : std::true_type {}; +template<> struct is_property_key< + kernel_arg_agent_register_key> : std::true_type {}; +template<> struct is_property_key : std::true_type {}; +template<> struct is_property_key : std::true_type {}; +template<> struct is_property_key : std::true_type {}; +template<> struct is_property_key : std::true_type {}; +template<> struct is_property_key< + kernel_arg_read_write_mode_key> : std::true_type {}; +template<> struct is_property_key< + kernel_arg_latency_key> : std::true_type {}; +template<> struct is_property_key : std::true_type {}; +template<> struct is_property_key< + kernel_arg_wait_request_key> : std::true_type {}; + +template +struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> : std::true_type {}; +} // namespace experimental::oneapi::ext::sycl +``` +-- + +NOTE: All of the properties defined in this extension except `mem_id` are +meaningful only the kernel argument and are not meaningful within the device +code. + +Hence, all the properties except `mem_id` allow implicit conversion of +`annotated_ptr` to `T`, effectively losing the annotation +within the device and only retaining them on the kernel argument. + +[frame="topbot",options="header"] +|=== +|Property |Description + +a| +[source,c++] +---- +kernel_arg_mm_host +---- +a| +Directs the compiler to generate a pointer argument interface with a dedicated +input port on the kernel to input the pointer address into the kernel. + +The interface can be further customized with the other properties listed in this +extension. However, specifying this property applies the following properties +with their default values: + +- mem_id +- kernel_arg_awidth +- kernel_arg_dwidth +- kernel_arg_read_write_mode +- kernel_arg_wait_request +- kernel_arg_maxburst +- kernel_arg_latency + +If any of these need to have a value different from the default, they must be +specified as separate property on the type. + +a| +[source,c++] +---- +kernel_arg_agent_register +---- +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++] +---- +kernel_arg_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. + +a| +[source,c++] +---- +mem_id +---- +a| +Specifies a global memory identifier for the pointer interface. A default +value is set by the implementation. + +a| +[source,c++] +---- +kernel_arg_awidth +---- +a| +This property can only be used if the property `kernel_arg_mm_host` has been +specified. + +Specifies the width of the memory-mapped address bus in bits. The default is +set to 64. Valid values: Integer value in the range 1 – 64. + +a| +[source,c++] +---- +kernel_arg_dwidth +---- +a| +This property can only be used if the property `kernel_arg_mm_host` has been +specified. + +Specifies the width of the memory-mapped data bus in bits. The default is set +to 64. Valid values: 8, 16, 32, 64, 128, 256, 512, or 1024. + +a| +[source,c++] +---- +kernel_arg_read_write_mode +---- +a| +This property can only be used if the property `kernel_arg_mm_host` has been +specified. + +Specifies the port direction of the interface. `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: + + - kernel_arg_read_write_mode_read + - kernel_arg_read_write_mode_write + - kernel_arg_read_write_mode_readwrite + +a| +[source,c++] +---- +kernel_arg_latency +---- +a| +This property can only be used if the property `kernel_arg_mm_host` has been +specified. + +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. + +a| +[source,c++] +---- +kernel_arg_maxburst +---- +a| +This property can only be used if the property `kernel_arg_mm_host` has been +specified. + +Specifies the maximum number of data transfers that can be associated with a +read or write transaction. The default is set to 1. + +Legal values: Integer value in the range 1 – 1024. + +a| +[source,c++] +---- +kernel_arg_wait_request +---- +a| +This property can only be used if the property `kernel_arg_mm_host` has been +specified. + +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: + + - kernel_arg_wait_request_requested + - kernel_arg_wait_request_not_requested +|=== +-- + +=== Aliases provided for convenience + +[source,c++] +---- +namespace sycl::ext::oneapi::experimental{ + template + using mm_host = annotated_ptr; + + template + using agent_register = annotated_ptr; +}; // namespace sycl::ext::oneapi::experimental +---- + +=== Usage Examples + +The example below shows a simple kernel with one customized pointer argument +interface `a` and a scalar kernel argument `n`. The interface has default values +for the customizable settings as the alias `mm_host` specifies the property +`kernel_arg_mm_host`. + +.Usage Example 1 +```c++ +using sycl::ext::oneapi::experimental; +struct MyKernel { + mm_host a; + int n; + MyKernel(mm_host a_, int n_) : a(a_), n(n_) {} + void operator()() const { + for (int i=0; i + auto array_a = malloc_shared(kN, q); + // ... + // MyKernel member 'a' has type mm_host, which can be constructed from an + // annotated_ptr or even an integer USM pointer since an + // annotated_ptr can be constructed from T* or from an + // annotated_ptr as long as 'PropertyListB' is a valid + // subset of 'PropertyListA' + q.single_task(MyKernel{array_a, kN}).wait(); + + // ... + sycl::free(array_a, q); +} +``` + +In the example below, the kernel arguments are two customized pointer interfaces +and a scalar argument. The two pointers point to separate memories as specified +by the property `mem_id`. + +.Usage Example 2 +```c++ +using sycl::ext::oneapi::experimental; +struct MyKernel { + using MyMMHostA = mm_host, kernel_arg_awidth<32>, + kernel_arg_dwidth<32>>>; + using MyMMHostB = mm_host, kernel_arg_awidth<128>, + kernel_arg_dwidth<128>>>; + // struct members become kernel arguments + MyMMHostA a; + MyMMHostB b; + int n; + MyKernel(MyMMHostA a_, MyMMHostB b_, int n_) : a(a_), b(b_), n(n_) {} + void operator()() const { + for (int i=0; i>> + auto array_a = malloc_shared( + kN, q, MyKernel::MyMMHostA::get_property()); + // Constructs an object of type annotated_ptr>> + auto array_b = malloc_shared( + kN, q, MyKernel::MyMMHostB::get_property()); + // ... + + // 'array_a', 'array_b' and MyKernel members 'a', and 'b' are all + // annotated_ptr objects which wrap pointers to integers + // 'a' can be constructed from 'array_a' because the properties on the type of + // the object 'a' are a legal super set of the properties on the type of + // object 'array_a'. Same applies for 'b' and 'array_b' + q.single_task(MyKernel{array_a, array_b, kN}).wait(); + + // ... + sycl::free(array_a, q); + sycl::free(array_b, q); +} +``` + +In the example below, the kernel arguments are two pointer interfaces where the +pointers' base addresses are stored in agent registers. The two pointers point +to separate memories as specified by the property `mem_id`. + +.Usage Example 3 +```c++ +using sycl::ext::oneapi::experimental; +struct MyKernel { + using MyMMHostA = mm_host< + int, properties>>; + using MyMMHostB = mm_host< + int, properties>>; + MyMMHostA a; + MyMMHostB b; + MyKernel(MyMMHostA a_, MyMMHostB b_) : a(a_), b(b_) {} + void operator()() const { + *a = (*a) + 2; + *b = (*b) * 2; + } +}; + +int main () { + sycl::queue q; + auto ptr1 = malloc_shared(1, q, + MyKernel::MyMMHostA::get_property()); + auto ptr2 = malloc_shared(1, q, + MyKernel::MyMMHostB::get_property()); + *ptr1 = 5; + *ptr2 = 10; + // Again, struct members can be constructed from ptr1 and ptr2 since the + // type of the struct members are annotated_ptrs with a valid superset of the + // properties of ptr1 and ptr2 and all of the types involved are wrappers over + // integer pointers + q.single_task(MyKernel{ptr1, ptr2}).wait(); + + // ... + sycl::free(ptr1, q); + sycl::free(ptr2, q); +} +``` + +== Issues + +1. Should property defaults and ranges be specified in the spec document or + should we just say that the defaults and ranges are defined by the + implementation? + +2. I am not sure if the property `kernel_arg_mm_host` is needed since +all annotated_ptrs will result in a pointer interface for the FPGA. Should I +remove it? If yes, should we change the `mm_host` alias to +`using mm_host = annotated_ptr` instead? Or remove +that from the extension all together? + +3. `kernel_arg` prefix is too long. Is a `mem_` prefix acceptable? Or maybe +just an `arg_` prefix? Is that coupled with the description/doc enough to +emphasize that these only make sense on the kernel arguments? + +4. Should we add a new property argument to `kernel_arg_latency` to separate +specifying fixed latency and variable latency. +Yes, this is a TODO. + +5. Examples here also depend on USM `malloc*` API returning `annotated_ptr`. +Should I link to that extension (create one if one doesn't exist), or not show +the `malloc` calls in the examples? I think I should do the latter to keep these +extensions separate. An implementation could support these properties and +not the `malloc` changes and still have value. + +6. Does the `annotated_ptr` spec need more clarification for the ctor: +`template explicit annotated_ptr( + annotated_ptr const & ConvertFrom)` ? What is a valid superset? How to +convey that constructing an +`annotated_ptr>>` with +`annotated_ptr>>` is illegal if `N` is not +32? + +7. Should the presence of `mem_id` cause deletion of the implicit conversion +operator `annotated_ptr::operator T*()`? I am unsure. + +== 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_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index cf15e9e92a154..b7bf125115c1a 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -21,36 +21,14 @@ :cpp: C++ :dpcpp: DPC++ -== Introduction -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. - == Notice -Copyright (c) 2021 Intel Corporation. All rights reserved. +Copyright (c) 2021-2022 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. -== Status - -Working Draft - -This is a preview extension specification, intended to provide early access to -a feature for review and community feedback. When the feature matures, this -specification may be released as a formal extension. - -Because the interfaces defined by this specification are not final and are -subject to change they are not intended to be used by shipping software -products. - -== Version - -Built On: {docdate} + -Revision: 2 - == Contact Abhishek Tiwari, Intel (abhishek2 'dot' tiwari 'at' intel 'dot' com) @@ -70,10 +48,24 @@ Roland Schulz, Intel This extension is written against the SYCL 2020 specification, revision 4. -It also depends on the `sycl_ext_oneapi_properties` extension. +It also depends on the +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 @@ -84,18 +76,45 @@ 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: +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. + 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. +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++ @@ -113,54 +132,55 @@ To code that looks like: { SomeType* a; ... - __special_store_with_annotation(a, some_val, I /*some compile time constant information*/); + __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*/); + __special_load_with_annotation(a, &read_val, + I /*some compile time constant information*/); } ``` -This mechanism does not meet requirements (1) and (2) listed above because: +This mechanism does not meet requirement (1) listed above because: - * Users have to replace all their pointer read/write code with special function calls. - * If the set or type of information represented by `I` evolves over time, it requires creating and/or updating these special functions and the underlying compiler implementation. + * Users have to replace all their pointer read/write code with special + function calls. -Another mechanism could be that the compiler provide 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: +Another mechanism could be that the compiler provide 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 { - [[special_annotations::use_case_x_attribute(I)]] SomeType* a; // 'I' is the compiler time constant information being annotated on 'a' + // '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 for 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, the compiler may not be able to observe a particular pointer load/store call and successfully trace back to the declaration to deduce which annotation was applied to it. +However, this mechanism fails to meet requirements (2) and (3) listed above +since: -The `annotated_ptr` class described in this document is a templated type 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. + * 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 observe a + particular pointer load/store call and successfully trace back to the + declaration to deduce which annotation was applied to it. -=== Feature test macro +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. -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 -|=== === 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`. +`annotated_ptr` is a class template, parameterized by the type of the underlying +allocation `T`, and a list of properties `PropertyListT`. [source,c++] ---- @@ -174,14 +194,16 @@ class annotated_ptr { Properties may be specified for an `annotated_ptr` to provide semantic modification or optimization hint information. -Example uses of a property are: +Here's an example of how a property could be used: [source,c++] ---- using namespace sycl::ext::oneapi; { sycl::queue q; - annotated_ptr>> kernel_arg; // alignment of the pointer in bytes specified using the property 'align' + // alignment of the pointer in bytes specified using the property 'alignment' + annotated_ptr>> kernel_arg = ...; + q.submit([=]{ *kernel_arg = *kernel_arg * 2; }); @@ -189,34 +211,40 @@ using namespace sycl::ext::oneapi; } ---- -The section below and the table that follows, describe the constructors, member functions and factory methods for `annotated_ptr`. -The section below refers to an `annotated_ref` class which is described in the section following this one. +The section below and the table that follows, describe the constructors, member +functions and factory methods for `annotated_ptr`. +The section below refers to an `annotated_ref` class which is described in the +section following this one. [source,c++] ---- namespace sycl::ext::oneapi { -template > +template > class annotated_ptr { public: using reference = annotated_ref; - annotated_ptr() noexcept = default; - explicit annotated_ptr(T *Ptr) noexcept = default; - annotated_ptr(annotated_ptr const &) noexcept = default; - - // Conversion function - template explicit annotated_ptr(annotated_ptr const &) noexcept; + annotated_ptr() noexcept; + explicit annotated_ptr(T *Ptr, const properties &) noexcept; + annotated_ptr(annotated_ptr const &) noexcept; + template explicit annotated_ptr( + annotated_ptr const &) noexcept; + template + explicit annotated_ptr(annotated_ptr const &, + properties) noexcept; reference operator*() const noexcept; reference operator[](std::ptrdiff_t) const noexcept; - annotated_ptr operator+(unsigned long long) const noexcept; + annotated_ptr operator+(size_t) const noexcept; std::ptrdiff_t operator-(annotated_ptr) const noexcept; operator bool() const noexcept; - // Implicit conversion is not supported - operator T*() noexcept = delete; - operator const T*() const noexcept = delete; + // Implicit conversion is not supported in all cases, the "unspecified" + // SFINAE logic will ensure these are generated only when it is legal to do + // so + /*unspecified*/ operator T*() noexcept; + /*unspecified*/ operator const T*() const noexcept; T* get() noexcept; const T* get() const noexcept; @@ -255,16 +283,33 @@ a| annotated_ptr() noexcept; ---- | -Constructs an `annotated_ptr` object. Does not allocate new storage. The underlying pointer is initialized to `nullptr`. +Constructs an `annotated_ptr` object. Does not allocate new storage. The +underlying pointer is initialized to `nullptr`. + +// --- ROW BREAK --- +a| +[source,c++] +---- +explicit annotated_ptr(T *Ptr, const properties &P); +---- +| +Constructs an `annotated_ptr` object. Does not allocate new storage. The +underlying pointer is initialized with `Ptr`. `P` is used to specify the +`PropertyListT` type on the class. // --- ROW BREAK --- a| [source,c++] ---- -explicit annotated_ptr(T *Ptr); +template +explicit annotated_ptr( + annotated_ptr const &Ptr, + properties P) noexcept; ---- | -Constructs an `annotated_ptr` object. Does not allocate new storage. The underlying pointer is initialized with `Ptr`. +Constructs an `annotated_ptr` object. Does not allocate new storage. The +underlying pointer is initialized with `Ptr`. `PropertyListU` and +`PropertyListV` will be combined to construct `PropertyListT`. // --- ROW BREAK --- a| @@ -273,16 +318,20 @@ a| annotated_ptr(annotated_ptr const &) noexcept = default; ---- | -Constructs an `annotated_ptr` object from another `annotated_ptr` with the same template parameterization object. +Constructs an `annotated_ptr` object from another `annotated_ptr` with the same +template parameterization object. // --- ROW BREAK --- a| [source,c++] ---- -template explicit annotated_ptr(annotated_ptr const & ConvertFrom); +template explicit annotated_ptr( + annotated_ptr const & ConvertFrom); ---- | -Constructs the `annotated_ptr` object from the `ConvertFrom` object if the template parameter types are compatible. For example, if the properties on one of the operands conflict with the others then the compiler should issue a compile time error. +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`. // --- ROW BREAK --- a| @@ -291,7 +340,9 @@ a| reference operator*() const; ---- | -Returns a reference wrapper which can be used to read or write to the underlying pointer. Reads/Writes using the reference will retain the annotations. +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| @@ -306,10 +357,10 @@ Returns an `annotated_ref` reference wrapper to the object at offset `Index`. a| [source,c++] ---- -annotated_ptr operator+(unsigned long long Offset) const; +annotated_ptr operator+(size_t Offset) const; ---- | -Returns an `annotated_ptr` that points to the location `Offset` distance away from the underlying pointer. +Returns an `annotated_ptr` that points to `this[Offset]`. // --- ROW BREAK --- a| @@ -318,7 +369,8 @@ a| std::ptrdiff_t operator-(annotated_ptr FromPtr) const; ---- | -Returns the distance between the underlying pointer and the pointer encapsulated by `FromPtr`. +Returns the distance between the underlying pointer and the pointer encapsulated +by `FromPtr`. // --- ROW BREAK --- a| @@ -333,11 +385,12 @@ Returns `false` if the underlying pointer is null, returns `true` otherwise. a| [source,c++] ---- -operator T*() noexcept = delete; -operator const T*() const noexcept = delete; +/*unspecified*/ operator T*() noexcept; +/*unspecified*/ operator const T*() const noexcept; ---- | -Implicit conversion to a pointer to the underlying type `T` is not supported. +Implicit conversion to a pointer to the underlying type is available only when +all the properties within `PropertyListT` allow implicit conversion. // --- ROW BREAK --- a| @@ -347,7 +400,8 @@ T* get() noexcept; const T* get() const noexcept; ---- | -Returns the underlying raw pointer. The raw pointer will not retain the annotations. +Returns the underlying raw pointer. The raw pointer will not retain the +annotations. // --- ROW BREAK --- a| @@ -356,7 +410,7 @@ a| annotated_ptr& operator=(const T*) noexcept; ---- | -Allows assignment from an pointer to type `T`. +Allows assignment from a pointer to type `T`. // --- ROW BREAK --- a| @@ -410,8 +464,11 @@ a| template static constexpr bool has_property(); ---- -| Returns true if the `PropertyListT` contains the property specified by `propertyT`. Returns false if it does not. -Available only if `sycl::is_property_of_v` is true. +| +Returns true if the property list 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| @@ -420,18 +477,24 @@ a| template static constexpr auto get_property(); ---- -| Returns an object of the class used to represent the value of property `propertyT`. -Must produce a compiler diagnostic if `PropertyListT` does not contain a `propertyT` property. -Available only if `sycl::is_property_of_v` is true. +| +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. + +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 the implementation to preserve the properties on loads from and stores to the pointers. +The purpose of the `annotated_ref` class template is to provide reference +wrapper semantics. It enables the implementation to preserve the properties on +loads from and stores to the pointers. ```c++ namespace sycl::ext::oneapi { -template > +template > class annotated_ref { public: annotated_ref(T *); @@ -457,7 +520,8 @@ a| annotated_ref(T * InputPtr); ---- | -Constructs an `annotated_ref` object. Does not allocate new storage. The underlying pointer is initialized to `InputPtr`. +Constructs an `annotated_ref` object. Does not allocate new storage. The +underlying pointer is initialized to `InputPtr`. // --- ROW BREAK --- a| @@ -467,7 +531,8 @@ operator T() noexcept; operator const T() const noexcept; ---- | -Implicit conversion to underlying type. +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| @@ -476,27 +541,109 @@ a| void operator=(const T &); ---- | -Enables assignment to the underlying pointer. +Writes an object of type `T` to the location referenced by this wrapper, +applying the annotations when the object is stored to memory. |=== -=== Properties for `annotated_ptr` variables +=== Properties + +Below is a list of compile-time constant properties supported with +`annotated_ptr`. + +When the implicit conversion operator which converts +`annotated_ptr` to `T*` is used, the annotations will not +be retained in device code when the `T*` pointer is used. This is acceptable for +some properties and not for others. + +This implicit conversion operator is deleted if the property specifies that +it needs to be retained within the device code. -See the extension `sycl_ext_oneapi_annotated_ptr_properties` for the full list of supported properties. +```c++ +namespace sycl::ext::oneapi::experimental { +struct alignment_key { + template + using value_t = property_value>; +}; + +struct kernel_arg_restrict_key { + using value_t = property_value; +}; + +struct runtime_aligned_key { + using value_t = property_value; +}; + +template +inline constexpr alignment_key::value_t alignment; +inline constexpr kernel_arg_restrict_key::value_t kernel_arg_restrict; +inline constexpr runtime_aligned_key::value_t runtime_aligned; + +template<> +struct is_property_key : std::true_type {}; +template<> +struct is_property_key : std::true_type {}; +template<> +struct is_property_key : std::true_type {}; + +template +struct is_property_key_of< + alignment_key, annotated_ptr> : std::true_type {}; +template +struct is_property_key_of< + kernel_arg_restrict_key, annotated_ptr> : std::true_type {}; +template +struct is_property_key_of< + runtime_aligned_key, annotated_ptr> : std::true_type {}; +} // namespace experimental::oneapi::ext::sycl +``` +-- +[options="header"] +|==== +| Property | Description +|`alignment` +| +Sets the alignment of the pointer address in bytes. + +|`kernel_arg_restrict` +| +Informs the compiler that the pointer kernel argument cannot alias with other +pointer kernel arguments. + +This property does not need to be retained in the device code but it must be +applied to the kernel argument. + +|`runtime_aligned` +| +Informs the compiler that the pointer has the alignment as determined by the +runtime specification. +|==== +-- -== Issues +== Issues related to `annotated_ptr` 1) [RESOLVED] Should we allow implicit conversion to base class by default? -*Ans: No.* +Ans: Enabling conversion to underlying pointer will result in loss of the +annotations when that underlying pointer is used. Some use-cases will benefit +from this if they only need to retain the annotations on the kernel function +interface (and not on the load/store sites within the device code). Hence, +conversion will be allowed except when the property list contains properties +whose definitions disallow implicit conversion. +This can be implemented via SFINAE. 2) [RESOLVED] How do we support `operator->`? -*Ans: Not with the initial release.* +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`? -*Ans: Yes. This discussion is an implementation detail discussion and does not impact the spec.* +Yes. This discussion is an implementation detail discussion and does not +impact the annotated_ptr spec. -4) [RESOLVED] Should we provide conversion functions to convert to/from multi_ptr? -*Ans: No.* +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. == Revision History @@ -505,6 +652,7 @@ See the extension `sycl_ext_oneapi_annotated_ptr_properties` for the full list o [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* |======================================== \ No newline at end of file diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr_properties.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr_properties.asciidoc deleted file mode 100755 index ec8108b34e163..0000000000000 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr_properties.asciidoc +++ /dev/null @@ -1,158 +0,0 @@ -= sycl_ext_oneapi_annotated_ptr_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++ - -== Introduction -This extension defines the properties supported with the class `annotated_ptr` -which is defined by the extension `sycl_ext_oneapi_annotated_ptr`. These -properties are carried by the `annotated_ptr` class and implemented by the -toolchain. - -== 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. - -== Status - -Final Draft - -This is a preview extension specification, intended to provide early access to -a feature for review and community feedback. When the feature matures, this -specification may be released as a formal extension. - -Because the interfaces defined by this specification are not final and are -subject to change they are not intended to be used by shipping software -products. - -== Version - -Built On: {docdate} + -Revision: 1 - -== Contact - -Abhishek Tiwari, Intel (abhishek2 'dot' tiwari 'at' intel 'dot' com) - -== 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: - - sycl_ext_oneapi_annotated_ptr - - sycl_ext_intel_buffer_location - -== Overview - -This extension introduces properties that establish the annotations that will be -carried by the class `sycl::ext::oneapi::annotated_ptr`. - -An example of the syntax can be seen below. - -[source,c++] ----- -annotated_ptr>> aptr; ----- - -== 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_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 APIs the implementation supports. - -[%header,cols="1,5"] -|=== -|Value |Description -|1 |Initial extension version. Base features are supported. -|=== - - -=== New Section 4.7.9.1 annotated_ptr properties - -Below is a list of compile-time-constant properties supported with `annotated_ptr`. - -```c++ -namespace sycl::ext::oneapi::experimental { -struct align_key { - template - using value_t = property_value>; -}; - -template -inline constexpr align::value_t align; - -template<> -struct is_property_key : std::true_type {}; - -template -struct is_property_key_of : std::true_type {}; - -} // namespace experimental::oneapi::ext::sycl -``` --- -[options="header"] -|==== -| Property | Description -|`align` -| The alignment of the pointer address in bytes. - -| Property | Description -|`buffer_location` -| The alignment of the pointer address in bytes. -See the extension `sycl_ext_intel_buffer_location` for more details. -|==== --- - -== Issues - -== Revision History - -[cols="5,15,15,70"] -[grid="rows"] -[options="header"] -|======================================== -|Rev|Date|Author|Changes -|1|2022-02-17|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. -//************************************************************************ From f4ba6d51c094d5236339c55350f39bf56e8c3d10 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Wed, 13 Apr 2022 00:22:57 -0700 Subject: [PATCH 04/55] update ctor row in table to match synopsis --- .../extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index b7bf125115c1a..b2757b2b1d1de 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -290,7 +290,7 @@ underlying pointer is initialized to `nullptr`. a| [source,c++] ---- -explicit annotated_ptr(T *Ptr, const properties &P); +explicit annotated_ptr(T *Ptr, const properties &P) noexcept; ---- | Constructs an `annotated_ptr` object. Does not allocate new storage. The From a3cf4c76ece56f7762c9f1e26b20d8a954f871a9 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Wed, 13 Apr 2022 00:36:37 -0700 Subject: [PATCH 05/55] fix description sentence --- .../proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index b2757b2b1d1de..b9dccd75de9c8 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -166,9 +166,9 @@ since: 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 observe a - particular pointer load/store call and successfully trace back to the - declaration to deduce which annotation was applied to it. + 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 From 52ae99299e1f3bbc4c3ea18a1ddeed20118fb9b6 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Mon, 18 Apr 2022 20:36:47 -0700 Subject: [PATCH 06/55] rename to buffer_location, +conduit property, address comments, rm mmhost prop --- ...tel_fpga_annotated_ptr_properties.asciidoc | 233 +++++------------- .../sycl_ext_oneapi_annotated_ptr.asciidoc | 10 +- 2 files changed, 66 insertions(+), 177 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_ptr_properties.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_ptr_properties.asciidoc index 52aa8eb808e86..9a7f7561fae27 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_ptr_properties.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_ptr_properties.asciidoc @@ -70,13 +70,13 @@ Some examples of the syntax are shown below. [source,c++] ---- annotated_ptr, kernel_arg_awidth<32>, kernel_arg_dwidth<>64>> ptr_a; + buffer_location<1>, kernel_arg_awidth<32>, kernel_arg_dwidth<>64>> ptr_a; ---- .Example 2 [source,c++] ---- -auto ptr_p = malloc_shared(1, q, properties{mem_id<2>}); +auto ptr_p = malloc_shared(1, q, properties{buffer_location<2>}); auto ptr_q = annotated_ptr(ptr_a, properties{ kernel_arg_awidth<32>, kernel_arg_dwidth<64>}); ---- @@ -102,29 +102,29 @@ implementation supports. |Initial version of this extension. |=== -=== Kernel Argument Interface Properties +=== `annotated_ptr` Properties Below is a list of compile-time constant properties supported with `annotated_ptr`. These properties control the kernel argument interface on FPGA devices. ```c++ -namespace sycl::ext::oneapi::experimental { -struct kernel_arg_mm_host_key { - using value_t = property_value; +struct kernel_arg_conduit_key { + using value_t = property_value; }; -struct kernel_arg_agent_register_key { - using value_t = property_value; +struct kernel_arg_register_map_key { + using value_t = property_value; }; struct kernel_arg_stable_key { using value_t = property_value; }; -struct mem_id_key { +struct buffer_location_key { template - using value_t = property_value>; + using value_t = property_value< + buffer_location_key, std::integral_constant>; }; struct kernel_arg_awidth_key { @@ -169,11 +169,13 @@ struct kernel_arg_wait_request_key { std::integral_constant>; }; -inline constexpr kernel_arg_mm_host_key::value_t kernel_arg_mm_host; -inline constexpr kernel_arg_agent_register_key::value_t - kernel_arg_agent_register; +inline constexpr kernel_arg_conduit_key::value_t + kernel_arg_conduit; +inline constexpr kernel_arg_register_map_key::value_t + kernel_arg_register_map; inline constexpr kernel_arg_stable_key::value_t kernel_arg_stable; -template inline constexpr mem_id_key::value_t mem_id; +template inline constexpr buffer_location_key::value_t + buffer_location; template inline constexpr kernel_arg_awidth_key::value_t kernel_arg_awidth; template inline constexpr kernel_arg_dwidth_key::value_t @@ -199,11 +201,11 @@ inline constexpr kernel_arg_wait_request_key::value_t inline constexpr kernel_arg_wait_request_key::value_t kernel_arg_wait_request_not_requested; -template<> struct is_property_key : std::true_type {}; +template<> struct is_property_key : std::true_type {}; template<> struct is_property_key< - kernel_arg_agent_register_key> : std::true_type {}; + kernel_arg_register_map_key> : std::true_type {}; template<> struct is_property_key : std::true_type {}; -template<> struct is_property_key : std::true_type {}; +template<> struct is_property_key : std::true_type {}; template<> struct is_property_key : std::true_type {}; template<> struct is_property_key : std::true_type {}; template<> struct is_property_key< @@ -215,16 +217,16 @@ template<> struct is_property_key< kernel_arg_wait_request_key> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template struct is_property_key_of` to `T`, effectively losing the annotation -within the device and only retaining them on the kernel argument. +within the device code when this conversion operator is used. The properties +are only preserved on the kernel argument. [frame="topbot",options="header"] |=== @@ -263,31 +265,16 @@ within the device and only retaining them on the kernel argument. a| [source,c++] ---- -kernel_arg_mm_host +kernel_arg_conduit ---- a| -Directs the compiler to generate a pointer argument interface with a dedicated -input port on the kernel to input the pointer address into the kernel. - -The interface can be further customized with the other properties listed in this -extension. However, specifying this property applies the following properties -with their default values: - -- mem_id -- kernel_arg_awidth -- kernel_arg_dwidth -- kernel_arg_read_write_mode -- kernel_arg_wait_request -- kernel_arg_maxburst -- kernel_arg_latency - -If any of these need to have a value different from the default, they must be -specified as separate property on the type. +Directs the compiler to create a dedicated input port on the kernel for the +input data. a| [source,c++] ---- -kernel_arg_agent_register +kernel_arg_register_map ---- a| Directs the compiler to create a register to store the base address of the @@ -310,11 +297,10 @@ undefined. a| [source,c++] ---- -mem_id +buffer_location ---- a| -Specifies a global memory identifier for the pointer interface. A default -value is set by the implementation. +Specifies a global memory identifier for the pointer interface. a| [source,c++] @@ -322,11 +308,8 @@ a| kernel_arg_awidth ---- a| -This property can only be used if the property `kernel_arg_mm_host` has been -specified. - Specifies the width of the memory-mapped address bus in bits. The default is -set to 64. Valid values: Integer value in the range 1 – 64. +set to 64. a| [source,c++] @@ -334,11 +317,8 @@ a| kernel_arg_dwidth ---- a| -This property can only be used if the property `kernel_arg_mm_host` has been -specified. - Specifies the width of the memory-mapped data bus in bits. The default is set -to 64. Valid values: 8, 16, 32, 64, 128, 256, 512, or 1024. +to 64. a| [source,c++] @@ -346,9 +326,6 @@ a| kernel_arg_read_write_mode ---- a| -This property can only be used if the property `kernel_arg_mm_host` has been -specified. - Specifies the port direction of the interface. `mode` can be one of: `read_write` - Interface can be used for read and write operations. @@ -371,9 +348,6 @@ a| kernel_arg_latency ---- a| -This property can only be used if the property `kernel_arg_mm_host` has been -specified. - 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. @@ -387,23 +361,15 @@ a| kernel_arg_maxburst ---- a| -This property can only be used if the property `kernel_arg_mm_host` has been -specified. - Specifies the maximum number of data transfers that can be associated with a read or write transaction. The default is set to 1. -Legal values: Integer value in the range 1 – 1024. - a| [source,c++] ---- kernel_arg_wait_request ---- a| -This property can only be used if the property `kernel_arg_mm_host` has been -specified. - 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`. @@ -415,35 +381,19 @@ For convenience, the following are provided: |=== -- -=== Aliases provided for convenience - -[source,c++] ----- -namespace sycl::ext::oneapi::experimental{ - template - using mm_host = annotated_ptr; - - template - using agent_register = annotated_ptr; -}; // namespace sycl::ext::oneapi::experimental ----- - === Usage Examples The example below shows a simple kernel with one customized pointer argument -interface `a` and a scalar kernel argument `n`. The interface has default values -for the customizable settings as the alias `mm_host` specifies the property -`kernel_arg_mm_host`. +interface `a` and a scalar kernel argument `n`. The pointer properties such +as `kernel_arg_awidth`, `kernel_arg_awidth`, etc will take the default values. .Usage Example 1 ```c++ using sycl::ext::oneapi::experimental; struct MyKernel { - mm_host a; + annotated_ptr a; int n; - MyKernel(mm_host a_, int n_) : a(a_), n(n_) {} + MyKernel(annotated_ptr a_, int n_) : a(a_), n(n_) {} void operator()() const { for (int i=0; i auto array_a = malloc_shared(kN, q); // ... - // MyKernel member 'a' has type mm_host, which can be constructed from an - // annotated_ptr or even an integer USM pointer since an - // annotated_ptr can be constructed from T* or from an - // annotated_ptr as long as 'PropertyListB' is a valid - // subset of 'PropertyListA' q.single_task(MyKernel{array_a, kN}).wait(); // ... @@ -471,21 +416,21 @@ int main () { In the example below, the kernel arguments are two customized pointer interfaces and a scalar argument. The two pointers point to separate memories as specified -by the property `mem_id`. +by the property `buffer_location`. .Usage Example 2 ```c++ using sycl::ext::oneapi::experimental; struct MyKernel { - using MyMMHostA = mm_host, kernel_arg_awidth<32>, - kernel_arg_dwidth<32>>>; - using MyMMHostB = mm_host, kernel_arg_awidth<128>, - kernel_arg_dwidth<128>>>; + using MyPtrA = annotated_ptr, kernel_arg_awidth<32>, kernel_arg_dwidth<32>>>; + using MyPtrB = annotated_ptr, kernel_arg_awidth<128>, kernel_arg_dwidth<128>>>; // struct members become kernel arguments - MyMMHostA a; - MyMMHostB b; + MyPtrA a; + MyPtrB b; int n; - MyKernel(MyMMHostA a_, MyMMHostB b_, int n_) : a(a_), b(b_), n(n_) {} + MyKernel(MyPtrA a_, MyPtrB b_, int n_) : a(a_), b(b_), n(n_) {} void operator()() const { for (int i=0; i>> + // Constructs an object of type annotated_ptr>> auto array_a = malloc_shared( - kN, q, MyKernel::MyMMHostA::get_property()); - // Constructs an object of type annotated_ptr>> + kN, q, MyKernel::MyPtrA::get_property()); + // Constructs an object of type annotated_ptr>> auto array_b = malloc_shared( - kN, q, MyKernel::MyMMHostB::get_property()); + kN, q, MyKernel::MyPtrB::get_property()); // ... // 'array_a', 'array_b' and MyKernel members 'a', and 'b' are all @@ -518,83 +463,27 @@ int main () { } ``` -In the example below, the kernel arguments are two pointer interfaces where the -pointers' base addresses are stored in agent registers. The two pointers point -to separate memories as specified by the property `mem_id`. - -.Usage Example 3 -```c++ -using sycl::ext::oneapi::experimental; -struct MyKernel { - using MyMMHostA = mm_host< - int, properties>>; - using MyMMHostB = mm_host< - int, properties>>; - MyMMHostA a; - MyMMHostB b; - MyKernel(MyMMHostA a_, MyMMHostB b_) : a(a_), b(b_) {} - void operator()() const { - *a = (*a) + 2; - *b = (*b) * 2; - } -}; - -int main () { - sycl::queue q; - auto ptr1 = malloc_shared(1, q, - MyKernel::MyMMHostA::get_property()); - auto ptr2 = malloc_shared(1, q, - MyKernel::MyMMHostB::get_property()); - *ptr1 = 5; - *ptr2 = 10; - // Again, struct members can be constructed from ptr1 and ptr2 since the - // type of the struct members are annotated_ptrs with a valid superset of the - // properties of ptr1 and ptr2 and all of the types involved are wrappers over - // integer pointers - q.single_task(MyKernel{ptr1, ptr2}).wait(); - - // ... - sycl::free(ptr1, q); - sycl::free(ptr2, q); -} -``` - == Issues -1. Should property defaults and ranges be specified in the spec document or - should we just say that the defaults and ranges are defined by the - implementation? - -2. I am not sure if the property `kernel_arg_mm_host` is needed since -all annotated_ptrs will result in a pointer interface for the FPGA. Should I -remove it? If yes, should we change the `mm_host` alias to -`using mm_host = annotated_ptr` instead? Or remove -that from the extension all together? - -3. `kernel_arg` prefix is too long. Is a `mem_` prefix acceptable? Or maybe +1. `kernel_arg` prefix is too long. Is a `mem_` prefix acceptable? Or maybe just an `arg_` prefix? Is that coupled with the description/doc enough to -emphasize that these only make sense on the kernel arguments? +emphasize that these properties only influence kernel arguments? -4. Should we add a new property argument to `kernel_arg_latency` to separate +2. Should we add a new property argument to `kernel_arg_latency` to separate specifying fixed latency and variable latency. -Yes, this is a TODO. +Yes, in a future extension we can introduce a separate property. -5. Examples here also depend on USM `malloc*` API returning `annotated_ptr`. +3. [RESOLVED] Examples here also depend on USM `malloc*` API returning +`annotated_ptr`. Should I link to that extension (create one if one doesn't exist), or not show the `malloc` calls in the examples? I think I should do the latter to keep these extensions separate. An implementation could support these properties and not the `malloc` changes and still have value. +And. Can remove malloc calls from the examples here for keeping the spec +simple and to keep it focussed on `annotated_ptr`. -6. Does the `annotated_ptr` spec need more clarification for the ctor: -`template explicit annotated_ptr( - annotated_ptr const & ConvertFrom)` ? What is a valid superset? How to -convey that constructing an -`annotated_ptr>>` with -`annotated_ptr>>` is illegal if `N` is not -32? - -7. Should the presence of `mem_id` cause deletion of the implicit conversion -operator `annotated_ptr::operator T*()`? I am unsure. +4. TODO: Link the fpga_kernel_properties spec to this one. Specifying certain +fpga kernel properties should result in changes to kernel arguments. == Revision History diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index b9dccd75de9c8..9228d5f117f03 100755 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -185,7 +185,7 @@ allocation `T`, and a list of properties `PropertyListT`. [source,c++] ---- namespace sycl::ext::oneapi { -template > +template > class annotated_ptr { ... ---- @@ -201,11 +201,11 @@ Here's an example of how a property could be used: using namespace sycl::ext::oneapi; { sycl::queue q; - // alignment of the pointer in bytes specified using the property 'alignment' + // alignment of the pointer in bytes specified using the property 'alignment' annotated_ptr>> kernel_arg = ...; q.submit([=]{ - *kernel_arg = *kernel_arg * 2; + *kernel_arg = (*kernel_arg) * 2; }); ... } @@ -607,8 +607,8 @@ Sets the alignment of the pointer address in bytes. |`kernel_arg_restrict` | -Informs the compiler that the pointer kernel argument cannot alias with other -pointer kernel arguments. +Informs the compiler that writes to the address pointed to by this pointer +are only done by this pointer or pointers derived from it. This property does not need to be retained in the device code but it must be applied to the kernel argument. From 36936bc0a754bb525af21ecb01c7863fe59fc0aa Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Fri, 27 May 2022 12:12:32 -0700 Subject: [PATCH 07/55] address meeting comments --- ...tel_fpga_annotated_ptr_properties.asciidoc | 326 +++++++----------- .../sycl_ext_oneapi_annotated_ptr.asciidoc | 42 +-- 2 files changed, 148 insertions(+), 220 deletions(-) mode change 100755 => 100644 sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_ptr_properties.asciidoc mode change 100755 => 100644 sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_ptr_properties.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_ptr_properties.asciidoc old mode 100755 new mode 100644 index 9a7f7561fae27..a76933bd5bccc --- a/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_ptr_properties.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_ptr_properties.asciidoc @@ -69,16 +69,16 @@ Some examples of the syntax are shown below. .Example 1 [source,c++] ---- -annotated_ptr, kernel_arg_awidth<32>, kernel_arg_dwidth<>64>> ptr_a; +annotated_ptr, awidth<32>, dwidth<64>> arg_a; ---- .Example 2 [source,c++] ---- -auto ptr_p = malloc_shared(1, q, properties{buffer_location<2>}); -auto ptr_q = annotated_ptr(ptr_a, properties{ - kernel_arg_awidth<32>, kernel_arg_dwidth<64>}); +// Allocate memory +auto ptr = ... +auto arg_a = annotated_ptr(ptr, properties{buffer_location<1>, awidth<32>, dwidth<64>}); ---- == Specification @@ -108,17 +108,20 @@ Below is a list of compile-time constant properties supported with `annotated_ptr`. These properties control the kernel argument interface on FPGA devices. +All of the properties defined in this extension are meaningful only on the +kernel argument and are not meaningful within the kernel body. + ```c++ -struct kernel_arg_conduit_key { - using value_t = property_value; +struct conduit_key { + using value_t = property_value; }; -struct kernel_arg_register_map_key { - using value_t = property_value; +struct register_map_key { + using value_t = property_value; }; -struct kernel_arg_stable_key { - using value_t = property_value; +struct stable_key { + using value_t = property_value; }; struct buffer_location_key { @@ -127,137 +130,129 @@ struct buffer_location_key { buffer_location_key, std::integral_constant>; }; -struct kernel_arg_awidth_key { +struct awidth_key { template using value_t = property_value< - kernel_arg_awidth_key, std::integral_constant>; + awidth_key, std::integral_constant>; }; -struct kernel_arg_dwidth_key { +struct dwidth_key { template using value_t = property_value< - kernel_arg_dwidth_key, std::integral_constant>; + dwidth_key, std::integral_constant>; }; -enum class kernel_arg_read_write_mode_enum { +enum class read_write_mode_enum { read_write, read, write }; -struct kernel_arg_read_write_mode_key { - template - using value_t = property_value>; +struct read_write_mode_key { + template + using value_t = property_value>; }; -struct kernel_arg_latency_key { +struct latency_key { template - using value_t = property_value>; }; -struct kernel_arg_maxburst_key { +struct maxburst_key { template using value_t = property_value< - kernel_arg_maxburst_key, std::integral_constant>; + maxburst_key, std::integral_constant>; }; -struct kernel_arg_wait_request_key { +struct wait_request_key { template - using value_t = property_value>; }; -inline constexpr kernel_arg_conduit_key::value_t - kernel_arg_conduit; -inline constexpr kernel_arg_register_map_key::value_t - kernel_arg_register_map; -inline constexpr kernel_arg_stable_key::value_t kernel_arg_stable; +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 kernel_arg_awidth_key::value_t - kernel_arg_awidth; -template inline constexpr kernel_arg_dwidth_key::value_t - kernel_arg_dwidth; -template -inline constexpr kernel_arg_read_write_mode_key::value_t - kernel_arg_read_write_mode; -inline constexpr kernel_arg_read_write_mode_key::value_t< - kernel_arg_read_write_mode_enum::read> kernel_arg_read_write_mode_read; -inline constexpr kernel_arg_read_write_mode_key::value_t< - kernel_arg_read_write_mode_enum::write> kernel_arg_read_write_mode_write; -inline constexpr kernel_arg_read_write_mode_key::value_t< - kernel_arg_read_write_mode_enum::read_write> - kernel_arg_read_write_mode_readwrite; -template inline constexpr kernel_arg_latency_key::value_t - kernel_arg_latency; -template inline constexpr kernel_arg_maxburst_key::value_t - kernel_arg_maxburst; -template inline constexpr kernel_arg_wait_request_key::value_t - kernel_arg_wait_request; -inline constexpr kernel_arg_wait_request_key::value_t - kernel_arg_wait_request_requested; -inline constexpr kernel_arg_wait_request_key::value_t - kernel_arg_wait_request_not_requested; - -template<> struct is_property_key : std::true_type {}; +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; + +template<> struct is_property_key : std::true_type {}; template<> struct is_property_key< - kernel_arg_register_map_key> : std::true_type {}; -template<> struct is_property_key : std::true_type {}; + register_map_key> : std::true_type {}; +template<> struct is_property_key : std::true_type {}; template<> struct is_property_key : std::true_type {}; -template<> struct is_property_key : std::true_type {}; -template<> struct is_property_key : std::true_type {}; +template<> struct is_property_key : std::true_type {}; +template<> struct is_property_key : std::true_type {}; template<> struct is_property_key< - kernel_arg_read_write_mode_key> : std::true_type {}; + read_write_mode_key> : std::true_type {}; template<> struct is_property_key< - kernel_arg_latency_key> : std::true_type {}; -template<> struct is_property_key : std::true_type {}; + latency_key> : std::true_type {}; +template<> struct is_property_key : std::true_type {}; template<> struct is_property_key< - kernel_arg_wait_request_key> : std::true_type {}; + wait_request_key> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; } // namespace experimental::oneapi::ext::sycl ``` -- -NOTE: All of the properties defined in this extension are meaningful only on the -kernel argument and are not meaningful within the device code. - -All the properties defined in this specification allow implicit conversion of -`annotated_ptr` to `T`, effectively losing the annotation -within the device code when this conversion operator is used. The properties -are only preserved on the kernel argument. - [frame="topbot",options="header"] |=== |Property |Description @@ -265,31 +260,30 @@ are only preserved on the kernel argument. a| [source,c++] ---- -kernel_arg_conduit +conduit ---- a| Directs the compiler to create a dedicated input port on the kernel for the -input data. +input. a| [source,c++] ---- -kernel_arg_register_map +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. +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++] ---- -kernel_arg_stable +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. +Specifies that the input 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. @@ -305,7 +299,7 @@ Specifies a global memory identifier for the pointer interface. a| [source,c++] ---- -kernel_arg_awidth +awidth ---- a| Specifies the width of the memory-mapped address bus in bits. The default is @@ -314,7 +308,7 @@ set to 64. a| [source,c++] ---- -kernel_arg_dwidth +dwidth ---- a| Specifies the width of the memory-mapped data bus in bits. The default is set @@ -323,10 +317,11 @@ to 64. a| [source,c++] ---- -kernel_arg_read_write_mode +read_write_mode ---- a| -Specifies the port direction of the interface. `mode` can be one of: +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. @@ -338,14 +333,14 @@ The default is set to `read_write`. For convenience, the following are provided: - - kernel_arg_read_write_mode_read - - kernel_arg_read_write_mode_write - - kernel_arg_read_write_mode_readwrite + - read_write_mode_read + - read_write_mode_write + - read_write_mode_readwrite a| [source,c++] ---- -kernel_arg_latency +latency ---- a| Specifies the guaranteed latency in cycles, from when a read command exits @@ -358,7 +353,7 @@ fixed latency. a| [source,c++] ---- -kernel_arg_maxburst +maxburst ---- a| Specifies the maximum number of data transfers that can be associated with a @@ -367,7 +362,7 @@ read or write transaction. The default is set to 1. a| [source,c++] ---- -kernel_arg_wait_request +wait_request ---- a| Specifies whether the 'wait request' signal is generated or not. This signal is @@ -376,114 +371,51 @@ request. The default is set to `false`. For convenience, the following are provided: - - kernel_arg_wait_request_requested - - kernel_arg_wait_request_not_requested + - wait_request_requested + - wait_request_not_requested |=== -- === Usage Examples -The example below shows a simple kernel with one customized pointer argument -interface `a` and a scalar kernel argument `n`. The pointer properties such -as `kernel_arg_awidth`, `kernel_arg_awidth`, etc will take the default values. - -.Usage Example 1 -```c++ -using sycl::ext::oneapi::experimental; -struct MyKernel { - annotated_ptr a; - int n; - MyKernel(annotated_ptr a_, int n_) : a(a_), n(n_) {} - void operator()() const { - for (int i=0; i - auto array_a = malloc_shared(kN, q); - // ... - q.single_task(MyKernel{array_a, kN}).wait(); +The example below shows a simple kernel with two annotated pointer kernel +arguments 'arg_a' and 'arg_b'. - // ... - sycl::free(array_a, q); -} -``` - -In the example below, the kernel arguments are two customized pointer interfaces -and a scalar argument. The two pointers point to separate memories as specified -by the property `buffer_location`. - -.Usage Example 2 +.Usage Example ```c++ using sycl::ext::oneapi::experimental; -struct MyKernel { - using MyPtrA = annotated_ptr, kernel_arg_awidth<32>, kernel_arg_dwidth<32>>>; - using MyPtrB = annotated_ptr, kernel_arg_awidth<128>, kernel_arg_dwidth<128>>>; - // struct members become kernel arguments - MyPtrA a; - MyPtrB b; - int n; - MyKernel(MyPtrA a_, MyPtrB b_, int n_) : a(a_), b(b_), n(n_) {} - void operator()() const { - for (int i=0; i>> - auto array_a = malloc_shared( - kN, q, MyKernel::MyPtrA::get_property()); - // Constructs an object of type annotated_ptr>> - auto array_b = malloc_shared( - kN, q, MyKernel::MyPtrB::get_property()); - // ... - - // 'array_a', 'array_b' and MyKernel members 'a', and 'b' are all - // annotated_ptr objects which wrap pointers to integers - // 'a' can be constructed from 'array_a' because the properties on the type of - // the object 'a' are a legal super set of the properties on the type of - // object 'array_a'. Same applies for 'b' and 'array_b' - q.single_task(MyKernel{array_a, array_b, kN}).wait(); - - // ... - sycl::free(array_a, q); - sycl::free(array_b, q); +{ + sycl::queue q{...}; + + // Allocate memory + auto ptr_a = ... + auto ptr_b = ... + + // Add properties + auto arg_a = annotated_ptr( + ptr_a, properties{buffer_location<1>, awidth<32>, dwidth<32>}); + auto arg_b = annotated_ptr( + ptr_b, properties{buffer_location<2>, awidth<64>, dwidth<128>}); + + q.single_task([=] { + ... + arg_a[index] *= 2; + arg_b[index] *= 4; + ... + }).wait(); + + ... } ``` == Issues -1. `kernel_arg` prefix is too long. Is a `mem_` prefix acceptable? Or maybe -just an `arg_` prefix? Is that coupled with the description/doc enough to -emphasize that these properties only influence kernel arguments? - -2. Should we add a new property argument to `kernel_arg_latency` to separate -specifying fixed latency and variable latency. +1. Should we add a new property argument to `latency` to separate specifying +fixed latency and variable latency. Yes, in a future extension we can introduce a separate property. -3. [RESOLVED] Examples here also depend on USM `malloc*` API returning -`annotated_ptr`. -Should I link to that extension (create one if one doesn't exist), or not show -the `malloc` calls in the examples? I think I should do the latter to keep these -extensions separate. An implementation could support these properties and -not the `malloc` changes and still have value. -And. Can remove malloc calls from the examples here for keeping the spec -simple and to keep it focussed on `annotated_ptr`. - -4. TODO: Link the fpga_kernel_properties spec to this one. Specifying certain -fpga kernel properties should result in changes to kernel arguments. +2. How do I link the fpga_kernel_properties spec to this one, to specify that +certain fpga kernel properties should result in changes to kernel arguments. == Revision History @@ -491,8 +423,8 @@ fpga kernel properties should result in changes to kernel arguments. [grid="rows"] [options="header"] |======================================== -|Rev|Date|Author|Changes -|1|2022-04-13|Abhishek Tiwari|*Initial draft* +|Rev|Date |Author |Changes +|1 |2022-04-13 |Abhishek Tiwari |*Initial draft* |======================================== //************************************************************************ diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc old mode 100755 new mode 100644 index 9228d5f117f03..a4a32c08925f0 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -240,11 +240,9 @@ class annotated_ptr { operator bool() const noexcept; - // Implicit conversion is not supported in all cases, the "unspecified" - // SFINAE logic will ensure these are generated only when it is legal to do - // so - /*unspecified*/ operator T*() noexcept; - /*unspecified*/ operator const T*() const noexcept; + // Implicit conversion is not supported + operator T*() noexcept = delete; + operator const T*() const noexcept = delete; T* get() noexcept; const T* get() const noexcept; @@ -297,6 +295,9 @@ Constructs an `annotated_ptr` object. Does not allocate new storage. The underlying pointer is initialized with `Ptr`. `P` is used to specify the `PropertyListT` type on the class. +The new property set `PropertyListT` must contain all properties from `P`, +and if any common property takes a value, the value must be the same. + // --- ROW BREAK --- a| [source,c++] @@ -385,12 +386,11 @@ Returns `false` if the underlying pointer is null, returns `true` otherwise. a| [source,c++] ---- -/*unspecified*/ operator T*() noexcept; -/*unspecified*/ operator const T*() const noexcept; +/*unspecified*/ operator T*() noexcept = delete; +/*unspecified*/ operator const T*() const noexcept = delete; ---- | -Implicit conversion to a pointer to the underlying type is available only when -all the properties within `PropertyListT` allow implicit conversion. +Implicit conversion to a pointer to the underlying type is not supported. // --- ROW BREAK --- a| @@ -551,14 +551,6 @@ applying the annotations when the object is stored to memory. Below is a list of compile-time constant properties supported with `annotated_ptr`. -When the implicit conversion operator which converts -`annotated_ptr` to `T*` is used, the annotations will not -be retained in device code when the `T*` pointer is used. This is acceptable for -some properties and not for others. - -This implicit conversion operator is deleted if the property specifies that -it needs to be retained within the device code. - ```c++ namespace sycl::ext::oneapi::experimental { struct alignment_key { @@ -605,14 +597,11 @@ struct is_property_key_of< | Sets the alignment of the pointer address in bytes. -|`kernel_arg_restrict` +|`restrict` | Informs the compiler that writes to the address pointed to by this pointer are only done by this pointer or pointers derived from it. -This property does not need to be retained in the device code but it must be -applied to the kernel argument. - |`runtime_aligned` | Informs the compiler that the pointer has the alignment as determined by the @@ -637,14 +626,21 @@ not have meaningful usecases that require this support. Building the support is complicated 3) [RESOLVED] Can `sycl::atomic_ref` be used with `annotated_ref`? -Yes. This discussion is an implementation detail discussion and does not -impact the annotated_ptr spec. +`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. + == Revision History [cols="5,15,15,70"] From 0c8e179bb579ed72a6658b17df21e920065be88d Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Fri, 27 May 2022 14:44:17 -0700 Subject: [PATCH 08/55] polish examples, condolidate properties into single extension --- ...tel_fpga_annotated_arg_properties.asciidoc | 280 ------------------ ...intel_fpga_kernel_arg_properties.asciidoc} | 90 ++++-- .../sycl_ext_oneapi_annotated_ptr.asciidoc | 31 +- 3 files changed, 77 insertions(+), 324 deletions(-) delete mode 100644 sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_arg_properties.asciidoc rename sycl/doc/extensions/proposed/{sycl_ext_intel_fpga_annotated_ptr_properties.asciidoc => sycl_ext_intel_fpga_kernel_arg_properties.asciidoc} (83%) 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_annotated_ptr_properties.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_arg_properties.asciidoc similarity index 83% rename from sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_ptr_properties.asciidoc rename to sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_arg_properties.asciidoc index a76933bd5bccc..4808fa33b2d5a 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_annotated_ptr_properties.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_arg_properties.asciidoc @@ -1,4 +1,4 @@ -= sycl_ext_intel_fpga_annotated_ptr_properties += sycl_ext_intel_fpga_kernel_arg_properties :source-highlighter: coderay :coderay-linenums-mode: table @@ -49,6 +49,7 @@ 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_ptr] == Status @@ -60,34 +61,51 @@ not rely on APIs defined in this specification.* == Overview -This extension introduces properties for the class -`sycl::ext::oneapi::annotated_ptr`. The properties will influence the kernel -argument interfaces for FPGA kernels and can be ignored for other devices. +This extension introduces properties for the classes +`sycl::ext::oneapi::annotated_ptr` and `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_ptr, awidth<32>, dwidth<64>> arg_a; ----- - -.Example 2 -[source,c++] ----- -// Allocate memory auto ptr = ... -auto arg_a = annotated_ptr(ptr, properties{buffer_location<1>, awidth<32>, dwidth<64>}); +auto arg_a = annotated_ptr( + ptr, properties{buffer_location<1>, awidth<32>, dwidth<64>}); + +... + +auto ptr2 = ... +auto arg_b = annotated_arg( + ptr2, properties{buffer_location<2>, awidth<32>, dwidth<64>}); + +... + +int val = 5; +auto arg_c = annotated_arg(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_ANNOTATED_PTR_PROPERTIES` to one of the values +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 @@ -102,15 +120,18 @@ implementation supports. |Initial version of this extension. |=== -=== `annotated_ptr` Properties +=== `annotated_ptr` and `annotated_arg` Properties Below is a list of compile-time constant properties supported with -`annotated_ptr`. These properties control the kernel argument interface on FPGA -devices. +the `annotated_ptr` and `annotated_arg` classes. These properties control the +kernel argument interface on FPGA devices. All of the properties defined in this extension are meaningful only on the kernel argument and are not meaningful within the kernel body. +NOTE: The properties that are defined for pointers will be meaningful for +annotated_arg only when T is a pointer type + ```c++ struct conduit_key { using value_t = property_value; @@ -296,6 +317,8 @@ buffer_location a| Specifies a global memory identifier for the pointer interface. +This property is only meaningful on pointer kernel arguments. + a| [source,c++] ---- @@ -305,6 +328,8 @@ a| Specifies the width of the memory-mapped address bus in bits. The default is set to 64. +This property is only meaningful on pointer kernel arguments. + a| [source,c++] ---- @@ -314,6 +339,8 @@ a| Specifies the width of the memory-mapped data bus in bits. The default is set to 64. +This property is only meaningful on pointer kernel arguments. + a| [source,c++] ---- @@ -337,6 +364,8 @@ For convenience, the following are provided: - read_write_mode_write - read_write_mode_readwrite +This property is only meaningful on pointer kernel arguments. + a| [source,c++] ---- @@ -350,6 +379,8 @@ 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 on pointer kernel arguments. + a| [source,c++] ---- @@ -359,6 +390,8 @@ 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 on pointer kernel arguments. + a| [source,c++] ---- @@ -373,13 +406,15 @@ For convenience, the following are provided: - wait_request_requested - wait_request_not_requested + +This property is only meaningful on pointer kernel arguments. |=== -- === Usage Examples -The example below shows a simple kernel with two annotated pointer kernel -arguments 'arg_a' and 'arg_b'. +The example below shows a simple kernel with one `annotated_ptr` kernel +argument and one `annotated_arg` kernel argument. .Usage Example ```c++ @@ -389,19 +424,16 @@ using sycl::ext::oneapi::experimental; // Allocate memory auto ptr_a = ... - auto ptr_b = ... + constexpr int kN = 10; // Add properties - auto arg_a = annotated_ptr( - ptr_a, properties{buffer_location<1>, awidth<32>, dwidth<32>}); - auto arg_b = annotated_ptr( - ptr_b, properties{buffer_location<2>, awidth<64>, dwidth<128>}); + auto arg_a = annotated_ptr(ptr_a, properties{ + register_map, buffer_location<1>, awidth<18>, dwidth<64>}); + auto arg_n = annotated_ptr(kN, properties{register_map, stable}); q.single_task([=] { - ... - arg_a[index] *= 2; - arg_b[index] *= 4; - ... + for (int i=0; i>> kernel_arg = ...; + auto arg_a = annotated_ptr(ptr, properties{alignment<4>}); q.submit([=]{ - *kernel_arg = (*kernel_arg) * 2; + ... + *arg_a = (*arg_a) * 2; }); ... } @@ -558,8 +561,8 @@ struct alignment_key { using value_t = property_value>; }; -struct kernel_arg_restrict_key { - using value_t = property_value; +struct restrict_key { + using value_t = property_value; }; struct runtime_aligned_key { @@ -568,13 +571,13 @@ struct runtime_aligned_key { template inline constexpr alignment_key::value_t alignment; -inline constexpr kernel_arg_restrict_key::value_t kernel_arg_restrict; +inline constexpr restrict_key::value_t restrict; inline constexpr runtime_aligned_key::value_t runtime_aligned; template<> struct is_property_key : std::true_type {}; template<> -struct is_property_key : std::true_type {}; +struct is_property_key : std::true_type {}; template<> struct is_property_key : std::true_type {}; @@ -583,7 +586,7 @@ struct is_property_key_of< alignment_key, annotated_ptr> : std::true_type {}; template struct is_property_key_of< - kernel_arg_restrict_key, annotated_ptr> : std::true_type {}; + restrict_key, annotated_ptr> : std::true_type {}; template struct is_property_key_of< runtime_aligned_key, annotated_ptr> : std::true_type {}; @@ -612,20 +615,17 @@ runtime specification. == Issues related to `annotated_ptr` 1) [RESOLVED] Should we allow implicit conversion to base class by default? -Ans: Enabling conversion to underlying pointer will result in loss of the -annotations when that underlying pointer is used. Some use-cases will benefit -from this if they only need to retain the annotations on the kernel function -interface (and not on the load/store sites within the device code). Hence, -conversion will be allowed except when the property list contains properties -whose definitions disallow implicit conversion. -This can be implemented via SFINAE. + +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 @@ -634,6 +634,7 @@ 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. From 0361b1ce3372a406443ea25c4cff30b4f0a3eaf1 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Wed, 22 Jun 2022 08:33:15 -0700 Subject: [PATCH 09/55] address review comments --- ..._intel_fpga_kernel_arg_properties.asciidoc | 59 ++++++++++--- .../sycl_ext_oneapi_annotated_ptr.asciidoc | 84 ++++++++++--------- 2 files changed, 94 insertions(+), 49 deletions(-) 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 index 4808fa33b2d5a..96dcd823ca1a3 100644 --- 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 @@ -49,7 +49,7 @@ 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_ptr] + - link:sycl_ext_oneapi_annotated_arg.asciidoc[sycl_ext_oneapi_annotated_arg] == Status @@ -122,7 +122,7 @@ implementation supports. === `annotated_ptr` and `annotated_arg` Properties -Below is a list of compile-time constant properties supported with +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. @@ -130,9 +130,10 @@ All of the properties defined in this extension are meaningful only on the kernel argument and are not meaningful within the kernel body. NOTE: The properties that are defined for pointers will be meaningful for -annotated_arg only when T is a pointer type +annotated_arg only when T is a pointer type. ```c++ +namespace sycl::ext::oneapi::experimental { struct conduit_key { using value_t = property_value; }; @@ -270,7 +271,39 @@ struct is_property_key_of struct is_property_key_of> : std::true_type {}; -} // namespace experimental::oneapi::ext::sycl + +template +struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> : std::true_type {}; +template +struct is_property_key_of> : std::true_type {}; + +} // namespace sycl::ext::oneapi::experimental ``` -- @@ -328,7 +361,8 @@ a| Specifies the width of the memory-mapped address bus in bits. The default is set to 64. -This property is only meaningful on pointer kernel arguments. +This property is only meaningful for pointer kernel arguments and only +when the `buffer_location` property is specified. a| [source,c++] @@ -339,7 +373,8 @@ a| Specifies the width of the memory-mapped data bus in bits. The default is set to 64. -This property is only meaningful on pointer kernel arguments. +This property is only meaningful for pointer kernel arguments and only +when the `buffer_location` property is specified. a| [source,c++] @@ -364,7 +399,8 @@ For convenience, the following are provided: - read_write_mode_write - read_write_mode_readwrite -This property is only meaningful on pointer kernel arguments. +This property is only meaningful for pointer kernel arguments and only +when the `buffer_location` property is specified. a| [source,c++] @@ -379,7 +415,8 @@ 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 on pointer kernel arguments. +This property is only meaningful for pointer kernel arguments and only +when the `buffer_location` property is specified. a| [source,c++] @@ -390,7 +427,8 @@ 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 on pointer kernel arguments. +This property is only meaningful for pointer kernel arguments and only +when the `buffer_location` property is specified. a| [source,c++] @@ -407,7 +445,8 @@ For convenience, the following are provided: - wait_request_requested - wait_request_not_requested -This property is only meaningful on pointer kernel arguments. +This property is only meaningful for pointer kernel arguments and only +when the `buffer_location` property is specified. |=== -- diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 2fbc65e4abe7d..4c577ed57f5e5 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -184,7 +184,7 @@ allocation `T`, and a list of properties `PropertyListT`. [source,c++] ---- -namespace sycl::ext::oneapi { +namespace sycl::ext::oneapi::experimental { template > class annotated_ptr { ... @@ -214,30 +214,34 @@ using namespace sycl::ext::oneapi::experimental; } ---- -The section below and the table that follows, describe the constructors, member -functions and factory methods for `annotated_ptr`. +The section below and the table that follows, describe the constructors and +member functions for `annotated_ptr`. The section below refers to an `annotated_ref` class which is described in the section following this one. [source,c++] ---- -namespace sycl::ext::oneapi { +namespace sycl::ext::oneapi::experimental { template > class annotated_ptr { public: using reference = annotated_ref; annotated_ptr() noexcept; - explicit annotated_ptr(T *Ptr, const properties &) noexcept; - annotated_ptr(annotated_ptr const &) noexcept; - template explicit annotated_ptr( - annotated_ptr const &) noexcept; - template - explicit annotated_ptr(annotated_ptr const &, + annotated_ptr(T *Ptr, const properties &P = properties{}) noexcept; + + template explicit annotated_ptr( + annotated_ptr const &) noexcept; + template + explicit annotated_ptr(annotated_ptr const &, properties) noexcept; + annotated_ptr(annotated_ptr const &) noexcept; + 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; @@ -255,6 +259,7 @@ class annotated_ptr { annotated_ptr& operator++() noexcept; annotated_ptr operator++(int) noexcept; + annotated_ptr& operator--() noexcept; annotated_ptr operator--(int) noexcept; @@ -267,9 +272,9 @@ class annotated_ptr { static constexpr /*unspecified*/ get_property(); private: - T *Ptr; + /* unspecified */ }; -} // namespace sycl::ext::oneapi +} // namespace sycl::ext::oneapi::experimental ---- @@ -291,7 +296,7 @@ underlying pointer is initialized to `nullptr`. a| [source,c++] ---- -explicit annotated_ptr(T *Ptr, const properties &P) noexcept; +annotated_ptr(T *Ptr, const properties &P = properties{}) noexcept; ---- | Constructs an `annotated_ptr` object. Does not allocate new storage. The @@ -305,37 +310,36 @@ and if any common property takes a value, the value must be the same. a| [source,c++] ---- -template -explicit annotated_ptr( - annotated_ptr const &Ptr, - properties P) noexcept; +template explicit annotated_ptr( + annotated_ptr const & ConvertFrom); ---- | -Constructs an `annotated_ptr` object. Does not allocate new storage. The -underlying pointer is initialized with `Ptr`. `PropertyListU` and -`PropertyListV` will be combined to construct `PropertyListT`. +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`. // --- ROW BREAK --- a| [source,c++] ---- -annotated_ptr(annotated_ptr const &) noexcept = default; +template +explicit annotated_ptr(annotated_ptr const &Ptr, + properties P) noexcept; ---- | -Constructs an `annotated_ptr` object from another `annotated_ptr` with the same -template parameterization object. +Constructs an `annotated_ptr` object. Does not allocate new storage. The +underlying pointer is initialized with `Ptr`. `PropertyListU` and +`PropertyListV` will be combined to construct `PropertyListT`. // --- ROW BREAK --- a| [source,c++] ---- -template explicit annotated_ptr( - annotated_ptr const & ConvertFrom); +annotated_ptr(annotated_ptr const &) noexcept = default; ---- | -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`. +Constructs an `annotated_ptr` object from another `annotated_ptr` with the same +template parameterization object. // --- ROW BREAK --- a| @@ -437,7 +441,7 @@ Prefix increment operator. a| [source,c++] ---- -annotated_ptr operator++() noexcept; +annotated_ptr operator++(int) noexcept; ---- | Postfix increment operator. @@ -455,7 +459,7 @@ Prefix decrement operator. a| [source,c++] ---- -annotated_ptr operator--() noexcept; +annotated_ptr operator--(int) noexcept; ---- | Postfix decrement operator. @@ -492,11 +496,11 @@ 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 the implementation to preserve the properties on -loads from and stores to the pointers. +wrapper semantics. It enables properties to be preserved on loads from and +stores to annotated_ptrs. ```c++ -namespace sycl::ext::oneapi { +namespace sycl::ext::oneapi::experimental { template > class annotated_ref { public: @@ -507,7 +511,7 @@ class annotated_ref { private: T *Ptr; }; -} // namespace sycl::ext::oneapi +} // namespace sycl::ext::oneapi::experimental ``` @@ -552,7 +556,8 @@ applying the annotations when the object is stored to memory. === Properties Below is a list of compile-time constant properties supported with -`annotated_ptr`. +`annotated_ptr`. Other extensions can define additional compile-time constant or +runtime properties that can be supported with `annotated_ptr`. ```c++ namespace sycl::ext::oneapi::experimental { @@ -590,7 +595,7 @@ struct is_property_key_of< template struct is_property_key_of< runtime_aligned_key, annotated_ptr> : std::true_type {}; -} // namespace experimental::oneapi::ext::sycl +} // namespace sycl::ext::oneapi::experimental ``` -- [options="header"] @@ -598,7 +603,7 @@ struct is_property_key_of< | Property | Description |`alignment` | -Sets the alignment of the pointer address in bytes. +Indicates the alignment of the pointer in bytes. |`restrict` | @@ -607,8 +612,9 @@ are only done by this pointer or pointers derived from it. |`runtime_aligned` | -Informs the compiler that the pointer has the alignment as determined by the -runtime specification. +Informs the compiler that the pointer has at least the default alignment for a +pointer allocated through the SYCL runtime. This is always safe to apply to any +pointer returned by a SYCL memory allocation function. |==== -- From 1cf8b984e4414cdd45b5833063ede21fbefa76d9 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Thu, 23 Jun 2022 12:28:34 -0700 Subject: [PATCH 10/55] make corrections to the annotated_arg spec --- .../sycl_ext_oneapi_annotated_arg.asciidoc | 573 ++---------------- 1 file changed, 37 insertions(+), 536 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index 7252b053c3f68..bbbd1e318f85f 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -55,15 +55,9 @@ incompatible ways before it is finalized. *Shipping software products should not rely on APIs defined in this specification.* == Overview -This extension introduces a wrapper class that provides a mechanism to attach -compile-time constant information to a variable in a manner that allows the -compiler to reliably maintain and analyze the information when such variables are used as arguments to functions and kernels. - -== Overview - -The purpose of this document is to clearly describe and specify -`annotated_arg` and related concepts, types, and mechanisms, and to give -examples and context for their usage. +This extension introduces a wrapper class `sycl::ext::oneapi::annotated_arg` +that provides a mechanism to attach compile-time constant information to +kernel arguments in a reliable manner. [NOTE] ==== @@ -71,11 +65,12 @@ In this document, we use the shortened form `annotated_arg` to refer to the proposed `sycl::ext::oneapi::annotated_arg` class. ==== -The purpose of this document is to clearly describe and specify `annotated_arg` -and related concepts, types, and mechanisms, and to give examples and context -for their usage. +The purpose of this document is to clearly describe and specify +`annotated_arg` 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 @@ -96,9 +91,9 @@ supports. === Introduction The `annotated_arg` class enables users to attach compile-time constant -information to the kernel arguments with properties. +information to kernel arguments with properties. -The example below shows a use case with some properties `PropA` and `PropB`. +The example below shows a use case with some properties. .Toy Example [source,c++] @@ -124,12 +119,12 @@ int main () { ---- `PropA`, `PropB` and `PropC` are expected to be preserved on the kernel -arguments in a reliable manner. For example, the properties may be converted to -kernel function level metadata. +arguments in a reliable manner. -NOTE: `PropC` will just influence the kernel argument and not the device code -where the pointer is dereferenced. Use the `sycl::ext::oneapi::annotated_ptr` -class template to apply properties that must be preserved in the device code. +NOTE: `PropC` will only influence the kernel argument and not any pointer +dereference sites. Use the `sycl::ext::oneapi::annotated_ptr` +class template to apply properties that must be preserved at pointer dereference +sites. The example below shows a type of use-case which is not meant to be supported by `annotated_arg`: @@ -164,7 +159,7 @@ processed. [source,c++] ---- -namespace sycl::ext::oneapi { +namespace sycl::ext::oneapi::experimental { template < typename T, typename PropertyListT = properties<>> class annotated_arg { ... @@ -181,10 +176,11 @@ The section below describes the constructors and member functions for [source,c++] ---- -namespace sycl::ext::oneapi { - template > - class annotated_arg { - T data; +namespace sycl::ext::oneapi::experimental { +template > +class annotated_arg { + private: + /* unspecified */ public: annotated_arg(); @@ -194,142 +190,15 @@ namespace sycl::ext::oneapi { 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; - - // 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; - 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, return + // type will match the return type of T::operator[](std::ptrdiff_t) + /* ... */ operator[](std::ptrdiff_t idx) 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; - - // 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 + // template T::operator()(Args... args) + template /* unspecified */ operator()(Args... args) noexcept; + template /* unspecified */ operator()(Args... args) const noexcept; template static constexpr bool has_property(); @@ -338,8 +207,8 @@ namespace sycl::ext::oneapi { // instances of propertyT template static constexpr /*unspecified*/ get_property(); - }; -}; // namespace sycl::ext::oneapi +}; +} // namespace sycl::ext::oneapi::experimental ---- [frame="topbot",options="header"] @@ -377,381 +246,22 @@ operator const T&() const noexcept; 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` - -// --- 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!() 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; +/* unspecified */ operator[](std::ptrdiff_t idx) const noexcept; ---- | -Available if the `operator==(const T&)` 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| [source,c++] ---- -bool operator!=(const T&) const noexcept; +template /* unspecified */ operator()(Args... args) noexcept; +template /* unspecified */ operator()(Args... args) 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++] ----- -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++] ----- -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++() noexcept; ----- -| -Available if the `operator++` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T operator++(int) noexcept; ----- -| -Available if the `operator++(int)` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T& operator--() noexcept; ----- -| -Available if the `operator--` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -T operator--(int) noexcept; ----- -| -Available if the `operator--(int)` 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` - -Provides member access through `T` that is a pointer or a class which defines -`operator->`. - -// --- ROW BREAK --- -a| -[source,c++] ----- -T& operator[](std::ptrdiff_t idx) noexcept; -const T& operator[](std::ptrdiff_t idx) const noexcept; ----- -| -Available if the `operator[]` is valid for objects of type `T` - -// --- ROW BREAK --- -a| -[source,c++] ----- -template auto operator()(Args... args) noexcept; -template auto operator()(Args... args) 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 'call operator' defined for `T`. // --- ROW BREAK --- a| @@ -805,12 +315,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. -//************************************************************************ From 65f9ea4af34f3b67f10bb5ae3a17b1a43fc3ec3b Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Thu, 23 Jun 2022 12:42:09 -0700 Subject: [PATCH 11/55] namespace corrections --- .../proposed/sycl_ext_oneapi_annotated_arg.asciidoc | 7 ++++--- .../proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 2 +- 2 files changed, 5 insertions(+), 4 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index bbbd1e318f85f..f22ce9d3eb929 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -55,14 +55,15 @@ incompatible ways before it is finalized. *Shipping software products should not rely on APIs defined in this specification.* == Overview -This extension introduces a wrapper class `sycl::ext::oneapi::annotated_arg` +This extension introduces a wrapper class +`sycl::ext::oneapi::experimental::annotated_arg` that provides a mechanism to attach compile-time constant information to kernel arguments in a reliable manner. [NOTE] ==== In this document, we use the shortened form `annotated_arg` to refer to the -proposed `sycl::ext::oneapi::annotated_arg` class. +proposed `sycl::ext::oneapi::experimental::annotated_arg` class. ==== The purpose of this document is to clearly describe and specify @@ -122,7 +123,7 @@ int main () { arguments in a reliable manner. NOTE: `PropC` will only influence the kernel argument and not any pointer -dereference sites. Use the `sycl::ext::oneapi::annotated_ptr` +dereference sites. Use the `sycl::ext::oneapi::experimental::annotated_ptr` class template to apply properties that must be preserved at pointer dereference sites. diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 4c577ed57f5e5..c8619f07aa643 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -69,7 +69,7 @@ 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::annotated_ptr` class. +proposed `sycl::ext::oneapi::experimental::annotated_ptr` class. ==== The purpose of this document is to clearly describe and specify `annotated_ptr` From 9e92ae476194ec1aa0a7df3a4ddc2289003348d0 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Thu, 23 Jun 2022 12:50:34 -0700 Subject: [PATCH 12/55] add note about the new classes being legal kernel parameters --- .../proposed/sycl_ext_oneapi_annotated_arg.asciidoc | 4 ++++ .../proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 4 ++++ 2 files changed, 8 insertions(+) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index f22ce9d3eb929..05bb1943517f9 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -175,6 +175,10 @@ separate extensions. The section below describes the constructors and member functions for `annotated_arg`. +NOTE: The template parameter `T` in the definition of `annotated_arg` template +below must be a device copyable type or a legal parameter type as defined by the +SYCL specification. + [source,c++] ---- namespace sycl::ext::oneapi::experimental { diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index c8619f07aa643..b249a568c57a4 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -216,6 +216,10 @@ using namespace sycl::ext::oneapi::experimental; The section below and the table that follows, describe the constructors and member functions for `annotated_ptr`. + +NOTE: `annotated_ptr` is a device copyable type since it is a wrapper over +a pointer and pointers are trivially copyable. + The section below refers to an `annotated_ref` class which is described in the section following this one. From 25f55466e517fb6f5da57487a0d66bacceea594f Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Mon, 4 Jul 2022 07:10:47 -0700 Subject: [PATCH 13/55] address review comments, move restrict to a separate ext --- ..._intel_fpga_kernel_arg_properties.asciidoc | 4 +- .../sycl_ext_oneapi_annotated_arg.asciidoc | 18 +- .../sycl_ext_oneapi_annotated_ptr.asciidoc | 38 ++-- ...eapi_kernel_arg_restrict_property.asciidoc | 187 ++++++++++++++++++ 4 files changed, 209 insertions(+), 38 deletions(-) create mode 100644 sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc 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 index 96dcd823ca1a3..ddb6b5e782e78 100644 --- 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 @@ -359,7 +359,7 @@ awidth ---- a| Specifies the width of the memory-mapped address bus in bits. The default is -set to 64. +determined by the implementation. This property is only meaningful for pointer kernel arguments and only when the `buffer_location` property is specified. @@ -468,7 +468,7 @@ using sycl::ext::oneapi::experimental; // Add properties auto arg_a = annotated_ptr(ptr_a, properties{ register_map, buffer_location<1>, awidth<18>, dwidth<64>}); - auto arg_n = annotated_ptr(kN, properties{register_map, stable}); + auto arg_n = annotated_arg(kN, properties{register_map, stable}); q.single_task([=] { for (int i=0; i T::operator()(Args... args) - template /* unspecified */ operator()(Args... args) noexcept; - template /* unspecified */ operator()(Args... args) const noexcept; + template /* ... */ operator()(Args... args) noexcept; + template /* ... */ operator()(Args... args) const noexcept; template static constexpr bool has_property(); @@ -251,7 +251,7 @@ operator const T&() const noexcept; a| [source,c++] ---- -/* unspecified */ operator[](std::ptrdiff_t idx) const noexcept; +/* ... */ operator[](std::ptrdiff_t idx) const noexcept; ---- | Available if the `operator[]` is valid for objects of type `T`. This function @@ -261,8 +261,8 @@ will call the subscript operator defined for `T`. a| [source,c++] ---- -template /* unspecified */ operator()(Args... args) noexcept; -template /* unspecified */ operator()(Args... args) const noexcept; +template /* ... */ operator()(Args... args) noexcept; +template /* ... */ operator()(Args... args) const noexcept; ---- | Available if the `operator()` is valid for objects of type `T`. This function @@ -276,7 +276,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. @@ -289,9 +289,9 @@ 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. diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index b249a568c57a4..049ec523f03e9 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -140,12 +140,10 @@ To code that looks like: } ``` -This mechanism does not meet requirement (1) listed above because: +This mechanism does not meet requirement (1) listed above because users have to +replace all their pointer read/write code with special function calls. - * Users have to replace all their pointer read/write code with special - function calls. - -Another mechanism could be that the compiler provide attributes which can be +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 @@ -343,7 +341,7 @@ annotated_ptr(annotated_ptr const &) noexcept = default; ---- | Constructs an `annotated_ptr` object from another `annotated_ptr` with the same -template parameterization object. +type and properties. // --- ROW BREAK --- a| @@ -476,7 +474,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. @@ -489,9 +487,9 @@ template static constexpr auto 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. @@ -560,8 +558,9 @@ applying the annotations when the object is stored to memory. === Properties Below is a list of compile-time constant properties supported with -`annotated_ptr`. Other extensions can define additional compile-time constant or -runtime properties that can be supported with `annotated_ptr`. +`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 { @@ -570,33 +569,23 @@ struct alignment_key { using value_t = property_value>; }; -struct restrict_key { - using value_t = property_value; -}; - struct runtime_aligned_key { using value_t = property_value; }; template inline constexpr alignment_key::value_t alignment; -inline constexpr restrict_key::value_t restrict; inline constexpr runtime_aligned_key::value_t runtime_aligned; template<> struct is_property_key : std::true_type {}; template<> -struct is_property_key : std::true_type {}; -template<> struct is_property_key : std::true_type {}; template struct is_property_key_of< alignment_key, annotated_ptr> : std::true_type {}; template -struct is_property_key_of< - restrict_key, annotated_ptr> : std::true_type {}; -template struct is_property_key_of< runtime_aligned_key, annotated_ptr> : std::true_type {}; } // namespace sycl::ext::oneapi::experimental @@ -609,11 +598,6 @@ struct is_property_key_of< | Indicates the alignment of the pointer in bytes. -|`restrict` -| -Informs the compiler that writes to the address pointed to by this pointer -are only done by this pointer or pointers derived from it. - |`runtime_aligned` | Informs the compiler that the pointer has at least the default alignment for a diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc new file mode 100644 index 0000000000000..b31d1ce0829f0 --- /dev/null +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc @@ -0,0 +1,187 @@ += sycl_ext_oneapi_kernel_arg_restrict_property + +: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_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 the `restrict` kernel argument property for the +classes `sycl::ext::oneapi::annotated_ptr` and +`sycl::ext::oneapi::annotated_arg`. + +== 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_RESTRICT_PROPERTY` 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 below is supported with the +`sycl::ext::oneapi::experimental::annotated_ptr` and +`sycl::ext::oneapi::experimental::annotated_arg` classes. They are only +meaningful on the kernel argument when the kernel argument is a pointer type or +a pointer wrapper type and are 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 a hint to the compiler that the pointer kernel arguments marked with +this property do not alias with one another. + +|=== +-- + +=== 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. +//************************************************************************ From 50f2ca6af05e84bdab28fe7875f1706396c439ff Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Thu, 11 Aug 2022 12:40:04 -0700 Subject: [PATCH 14/55] move properties to sycl::ext::intel namespace --- .../sycl_ext_intel_fpga_kernel_arg_properties.asciidoc | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) 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 index ddb6b5e782e78..58f45c557275d 100644 --- 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 @@ -133,7 +133,8 @@ NOTE: The properties that are defined for pointers will be meaningful for annotated_arg only when T is a pointer type. ```c++ -namespace sycl::ext::oneapi::experimental { +namespace sycl::ext::intel::experimental { +using sycl::ext::oneapi::experimental::properties; struct conduit_key { using value_t = property_value; }; @@ -225,7 +226,10 @@ 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 +namespace sycl::ext::oneapi::experimental { +using sycl::ext::intel::experimental; template<> struct is_property_key : std::true_type {}; template<> struct is_property_key< register_map_key> : std::true_type {}; @@ -302,7 +306,6 @@ struct is_property_key_of struct is_property_key_of> : std::true_type {}; - } // namespace sycl::ext::oneapi::experimental ``` -- @@ -457,7 +460,7 @@ argument and one `annotated_arg` kernel argument. .Usage Example ```c++ -using sycl::ext::oneapi::experimental; +using sycl::ext::intel::experimental; { sycl::queue q{...}; From 61c3a0bb67718a13d8052c88c015a9797fd8d5cc Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Wed, 7 Sep 2022 10:37:12 -0700 Subject: [PATCH 15/55] specify namespaces --- ..._intel_fpga_kernel_arg_properties.asciidoc | 118 +++++++++++------- 1 file changed, 73 insertions(+), 45 deletions(-) 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 index 58f45c557275d..e49362e871efa 100644 --- 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 @@ -134,34 +134,36 @@ annotated_arg only when T is a pointer type. ```c++ namespace sycl::ext::intel::experimental { -using sycl::ext::oneapi::experimental::properties; struct conduit_key { - using value_t = property_value; + using value_t = sycl::ext::oneapi::experimental::properties::property_value< + conduit_key>; }; struct register_map_key { - using value_t = property_value; + using value_t = sycl::ext::oneapi::experimental::properties::property_value< + register_map_key>; }; struct stable_key { - using value_t = property_value; + using value_t = sycl::ext::oneapi::experimental::properties::property_value< + stable_key>; }; struct buffer_location_key { template - using value_t = property_value< + using value_t = sycl::ext::oneapi::experimental::properties::property_value< buffer_location_key, std::integral_constant>; }; struct awidth_key { template - using value_t = property_value< + using value_t = sycl::ext::oneapi::experimental::properties::property_value< awidth_key, std::integral_constant>; }; struct dwidth_key { template - using value_t = property_value< + using value_t = sycl::ext::oneapi::experimental::properties::property_value< dwidth_key, std::integral_constant>; }; @@ -173,26 +175,26 @@ enum class read_write_mode_enum { struct read_write_mode_key { template - using value_t = property_value>; + using value_t = sycl::ext::oneapi::experimental::properties::property_value< + read_write_mode_key, std::integral_constant>; }; struct latency_key { template - using value_t = property_value>; + using value_t = sycl::ext::oneapi::experimental::properties::property_value< + latency_key, std::integral_constant>; }; struct maxburst_key { template - using value_t = property_value< + using value_t = sycl::ext::oneapi::experimental::properties::property_value< maxburst_key, std::integral_constant>; }; struct wait_request_key { template - using value_t = property_value>; + using value_t = sycl::ext::oneapi::experimental::properties::property_value< + wait_request_key, std::integral_constant>; }; inline constexpr conduit_key::value_t @@ -228,83 +230,109 @@ inline constexpr wait_request_key::value_t wait_request_not_requested; } // namespace sycl::ext::intel::experimental +// Type trait specializations namespace sycl::ext::oneapi::experimental { -using sycl::ext::intel::experimental; -template<> struct is_property_key : std::true_type {}; template<> struct is_property_key< - register_map_key> : std::true_type {}; -template<> struct is_property_key : std::true_type {}; -template<> struct is_property_key : std::true_type {}; -template<> struct is_property_key : std::true_type {}; -template<> struct is_property_key : std::true_type {}; + 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< - read_write_mode_key> : std::true_type {}; + sycl::ext::intel::experimental::latency_key> : std::true_type {}; template<> struct is_property_key< - latency_key> : std::true_type {}; -template<> struct is_property_key : std::true_type {}; + sycl::ext::intel::experimental::maxburst_key> : std::true_type {}; template<> struct is_property_key< - wait_request_key> : std::true_type {}; + sycl::ext::intel::experimental::wait_request_key> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; template -struct is_property_key_of> : std::true_type {}; } // namespace sycl::ext::oneapi::experimental ``` From 9ec66c4edc88bc03211d4ad6c611b5c2a4d44396 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Wed, 7 Sep 2022 15:18:33 -0700 Subject: [PATCH 16/55] add alignment property note --- .../extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 049ec523f03e9..6b0c688ef6c05 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -596,7 +596,8 @@ struct is_property_key_of< | Property | Description |`alignment` | -Indicates the alignment of the pointer in bytes. +Indicates the alignment of the pointer in bytes. Alignment cannot be smaller +than the size of type `T`. |`runtime_aligned` | From f51eed34d1619e2775fb17f63826be178360aaeb Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Mon, 26 Sep 2022 11:59:02 -0700 Subject: [PATCH 17/55] address review comments 1. Update tempalte to state 2022 copyright, reference revision 5 and use the updated Notice section template 2. Remove annotated_ref ctor from spec 3. Place 'const' keyword consistently before the type 4. Add a note about trivial copy-ablity 5. Correct bad use of auto 6. Correct namespace in the Overview section 7. Add note about properties being ignored on non FPGA devices --- ..._intel_fpga_kernel_arg_properties.asciidoc | 21 ++++--- .../sycl_ext_oneapi_annotated_ptr.asciidoc | 55 +++++++++---------- ...eapi_kernel_arg_restrict_property.asciidoc | 11 ++-- 3 files changed, 44 insertions(+), 43 deletions(-) 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 index e49362e871efa..46d90e55f4b92 100644 --- 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 @@ -23,11 +23,12 @@ == Notice -Copyright (c) 2021 Intel Corporation. All rights reserved. +[%hardbreaks] +Copyright (C) 2022-2022 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. +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 @@ -43,7 +44,7 @@ Joseph Garvey, Intel == Dependencies -This extension is written against the SYCL 2020 specification, revision 4. +This extension is written against the SYCL 2020 specification, revision 5. It depends on the following extensions: @@ -62,9 +63,13 @@ not rely on APIs defined in this specification.* == Overview This extension introduces properties for the classes -`sycl::ext::oneapi::annotated_ptr` and `sycl::ext::oneapi::annotated_arg`. The -properties will influence the kernel argument interfaces for FPGA kernels and -can be ignored for other devices. +`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. + +NOTE: These properties are only valid on FPGA kernel arguments and should be +ignored on other devices. Some examples of the syntax are shown below. diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 6b0c688ef6c05..5505667a2d3ed 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -23,15 +23,19 @@ == Notice -Copyright (c) 2021-2022 Intel Corporation. All rights reserved. +[%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. -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 -Abhishek Tiwari, Intel (abhishek2 'dot' tiwari 'at' intel 'dot' com) +To report problems with this extension, please open a new issue at: + +https://github.com/intel/llvm/issues == Contributors @@ -46,10 +50,13 @@ Roland Schulz, Intel == Dependencies -This extension is written against the SYCL 2020 specification, revision 4. +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: -It also depends on the -link:../experimental/sycl_ext_oneapi_properties.asciidoc[sycl_ext_oneapi_properties] +* link:../experimental/sycl_ext_oneapi_properties.asciidoc[sycl_ext_oneapi_properties] extension. == Status @@ -199,8 +206,7 @@ Here's an example of how a property could be used: using namespace sycl::ext::oneapi::experimental; { sycl::queue q; - // A pointer to an integer - auto ptr = ... + int* ptr = ... // alignment of the pointer in bytes specified using the property 'alignment' auto arg_a = annotated_ptr(ptr, properties{alignment<4>}); @@ -215,8 +221,7 @@ using namespace sycl::ext::oneapi::experimental; The section below and the table that follows, describe the constructors and member functions for `annotated_ptr`. -NOTE: `annotated_ptr` is a device copyable type since it is a wrapper over -a pointer and pointers are trivially copyable. +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. @@ -233,12 +238,12 @@ class annotated_ptr { annotated_ptr(T *Ptr, const properties &P = properties{}) noexcept; template explicit annotated_ptr( - annotated_ptr const &) noexcept; + const annotated_ptr&) noexcept; template - explicit annotated_ptr(annotated_ptr const &, + explicit annotated_ptr(const annotated_ptr&, properties) noexcept; - annotated_ptr(annotated_ptr const &) noexcept; + annotated_ptr(const annotated_ptr&) noexcept; reference operator*() const noexcept; @@ -257,7 +262,7 @@ class annotated_ptr { const T* get() const noexcept; annotated_ptr& operator=(const T*) noexcept; - annotated_ptr& operator=(annotated_ptr const&) noexcept; + annotated_ptr& operator=(const annotated_ptr&) noexcept; annotated_ptr& operator++() noexcept; annotated_ptr operator++(int) noexcept; @@ -313,7 +318,7 @@ a| [source,c++] ---- template explicit annotated_ptr( - annotated_ptr const & ConvertFrom); + const annotated_ptr &ConvertFrom); ---- | Constructs the `annotated_ptr` object from the `ConvertFrom` object if @@ -325,7 +330,7 @@ a| [source,c++] ---- template -explicit annotated_ptr(annotated_ptr const &Ptr, +explicit annotated_ptr(const annotated_ptr &Ptr, properties P) noexcept; ---- | @@ -337,7 +342,7 @@ underlying pointer is initialized with `Ptr`. `PropertyListU` and a| [source,c++] ---- -annotated_ptr(annotated_ptr const &) noexcept = default; +annotated_ptr(const annotated_ptr &) noexcept = default; ---- | Constructs an `annotated_ptr` object from another `annotated_ptr` with the same @@ -506,12 +511,12 @@ namespace sycl::ext::oneapi::experimental { template > class annotated_ref { public: - annotated_ref(T *); operator T() noexcept; operator const T() const noexcept; void operator=(const T &); private: T *Ptr; + ... }; } // namespace sycl::ext::oneapi::experimental ``` @@ -522,16 +527,6 @@ Member Functions are described in the table below |=== |Functions |Description -// --- ROW BREAK --- -a| -[source,c++] ----- -annotated_ref(T * InputPtr); ----- -| -Constructs an `annotated_ref` object. Does not allocate new storage. The -underlying pointer is initialized to `InputPtr`. - // --- ROW BREAK --- a| [source,c++] diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc index b31d1ce0829f0..5b444c9b0c356 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc @@ -23,11 +23,12 @@ == Notice -Copyright (c) 2021 Intel Corporation. All rights reserved. +[%hardbreaks] +Copyright (C) 2022-2022 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. +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 @@ -43,7 +44,7 @@ Joseph Garvey, Intel == Dependencies -This extension is written against the SYCL 2020 specification, revision 4. +This extension is written against the SYCL 2020 specification, revision 5. It depends on the following extensions: From 67617fc6cd3765c92efbb132849d08a6827c530a Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Mon, 26 Sep 2022 12:00:41 -0700 Subject: [PATCH 18/55] update copyright date and sycl spec revision number --- .../proposed/sycl_ext_oneapi_annotated_arg.asciidoc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index 2f3ec679bba17..ef8140315e89a 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -23,7 +23,7 @@ == Notice [%hardbreaks] -Copyright (c) 2021-2022 Intel Corporation. All rights reserved. +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 @@ -41,7 +41,7 @@ Abhishek Tiwari, Intel == Dependencies -This extension is written against the SYCL 2020 specification, Revision 4 and +This extension is written against the SYCL 2020 specification, Revision 5 and the following extensions: - link:../experimental/sycl_ext_oneapi_properties.asciidoc[sycl_ext_oneapi_properties] From 37ed507f071faf4818f43b2dfced2e7d331ac41b Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Mon, 26 Sep 2022 12:05:48 -0700 Subject: [PATCH 19/55] leave private members as unspecified --- .../extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 5505667a2d3ed..67872eac6d8e2 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -515,8 +515,7 @@ class annotated_ref { operator const T() const noexcept; void operator=(const T &); private: - T *Ptr; - ... + /* unspecified */ }; } // namespace sycl::ext::oneapi::experimental ``` From 9474fc54a9a167d40935fa7638850fd940a361c9 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Mon, 26 Sep 2022 12:06:57 -0700 Subject: [PATCH 20/55] remove references to private members all together --- .../proposed/sycl_ext_oneapi_annotated_arg.asciidoc | 3 --- .../proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 5 ----- 2 files changed, 8 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index ef8140315e89a..dc07cd6b1deba 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -184,9 +184,6 @@ SYCL specification. namespace sycl::ext::oneapi::experimental { template > class annotated_arg { - private: - /* unspecified */ - public: annotated_arg(); annotated_arg(const T& v_); diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 67872eac6d8e2..1818cf83c2b52 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -277,9 +277,6 @@ class annotated_ptr { // instances of propertyT template static constexpr /*unspecified*/ get_property(); - - private: - /* unspecified */ }; } // namespace sycl::ext::oneapi::experimental @@ -514,8 +511,6 @@ class annotated_ref { operator T() noexcept; operator const T() const noexcept; void operator=(const T &); - private: - /* unspecified */ }; } // namespace sycl::ext::oneapi::experimental ``` From 46a4fe345e70dd779a283a83d17818f8e826d07b Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Mon, 26 Sep 2022 12:20:52 -0700 Subject: [PATCH 21/55] update annotated_ptr ctor notes --- .../proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 1818cf83c2b52..e421147dfde2c 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -322,6 +322,8 @@ 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++] @@ -333,7 +335,10 @@ explicit annotated_ptr(const annotated_ptr &Ptr, | Constructs an `annotated_ptr` object. Does not allocate new storage. The underlying pointer is initialized with `Ptr`. `PropertyListU` and -`PropertyListV` will be combined to construct `PropertyListT`. +`PropertyListV` will be merged together to construct `PropertyListT`. If any +common properties have different values then an error will be issued. + +`T2*` must be implicitly convertible to `T*`. // --- ROW BREAK --- a| From 45e39029505f176ea5fe878da05a10db8634f91f Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Mon, 26 Sep 2022 12:28:25 -0700 Subject: [PATCH 22/55] make bool operator explicit, correct get() function --- .../proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 10 ++++------ 1 file changed, 4 insertions(+), 6 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index e421147dfde2c..7b3ef3339ad74 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -252,14 +252,13 @@ class annotated_ptr { annotated_ptr operator+(size_t) const noexcept; std::ptrdiff_t operator-(annotated_ptr) const noexcept; - operator bool() const noexcept; + explicit operator bool() const noexcept; // Implicit conversion is not supported operator T*() noexcept = delete; operator const T*() const noexcept = delete; - T* get() noexcept; - const T* get() const noexcept; + T* get() const noexcept; annotated_ptr& operator=(const T*) noexcept; annotated_ptr& operator=(const annotated_ptr&) noexcept; @@ -393,7 +392,7 @@ by `FromPtr`. a| [source,c++] ---- -operator bool() const noexcept; +explicit operator bool() const noexcept; ---- | Returns `false` if the underlying pointer is null, returns `true` otherwise. @@ -412,8 +411,7 @@ Implicit conversion to a pointer to the underlying type is not supported. a| [source,c++] ---- -T* get() noexcept; -const T* get() const noexcept; +T* get() const noexcept; ---- | Returns the underlying raw pointer. The raw pointer will not retain the From 5df917be74e6da07dabeebc97fb81b995421e559 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Tue, 27 Sep 2022 07:10:50 -0700 Subject: [PATCH 23/55] normative note for fpga properties on non fpga devices --- ...sycl_ext_intel_fpga_kernel_arg_properties.asciidoc | 11 +++++------ 1 file changed, 5 insertions(+), 6 deletions(-) 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 index 46d90e55f4b92..fe70ba9dedaab 100644 --- 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 @@ -68,9 +68,6 @@ This extension introduces properties for the classes the kernel argument interfaces for FPGA kernels and can be ignored for other devices. -NOTE: These properties are only valid on FPGA kernel arguments and should be -ignored on other devices. - Some examples of the syntax are shown below. [source,c++] @@ -129,13 +126,15 @@ implementation supports. 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. +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. -NOTE: The properties that are defined for pointers will be meaningful for -annotated_arg only when T is a pointer type. +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 { From e7219e19d67ec680eefb70258d95f3014b72c1a8 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Tue, 27 Sep 2022 09:05:58 -0700 Subject: [PATCH 24/55] add default copy ctor and assignment operator --- .../sycl_ext_oneapi_annotated_arg.asciidoc | 27 ++++++++++++++++--- 1 file changed, 23 insertions(+), 4 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index dc07cd6b1deba..053ea24362305 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -153,7 +153,7 @@ 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 +kernel arguments and the other properties on members of `MyType` will not be processed. === Representation of `annotated_arg` @@ -175,9 +175,9 @@ separate extensions. The section below describes the constructors and member functions for `annotated_arg`. -NOTE: The template parameter `T` in the definition of `annotated_arg` template -below must be a device copyable type or a legal parameter type as defined by the -SYCL specification. +The template parameter `T` in the definition of `annotated_arg` template below +must be a device copyable type or a legal parameter type as defined by the SYCL +specification. [source,c++] ---- @@ -188,6 +188,9 @@ class annotated_arg { annotated_arg(); annotated_arg(const T& v_); + annotated_arg(const annotated_arg&) = default; + annotated_arg& operator=(annotated_arg&) = default; + // Conversion operator to convert to the underlying type operator T&() noexcept; operator const T&() const noexcept; @@ -235,6 +238,22 @@ annotated_arg(const T& v_); | Not available in device code. Constructs an `annotated_arg` object from the input object `v_`. +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_arg(const annotated_arg&) = default; +---- +Compiler generated copy constructor. + +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_arg& operator=(annotated_arg&) = default; +---- +Compiler generated assignment operator. + // --- ROW BREAK --- a| [source,c++] From 29ef127c1ff8d2da69935a09c5dfc5f5fe4421c2 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Tue, 27 Sep 2022 09:08:00 -0700 Subject: [PATCH 25/55] remove destructor --- .../proposed/sycl_ext_oneapi_annotated_arg.asciidoc | 9 --------- 1 file changed, 9 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index 053ea24362305..a92e0c00ef6cc 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -312,15 +312,6 @@ a compile error if `PropertyListT` does not contain a property with the 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 From c099b7433abbea67ed6467f1b6b0badc62a15b67 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Tue, 27 Sep 2022 09:19:32 -0700 Subject: [PATCH 26/55] correct assignment op arg --- .../proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 7b3ef3339ad74..ef3183e32ff2c 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -260,7 +260,7 @@ class annotated_ptr { T* get() const noexcept; - annotated_ptr& operator=(const T*) noexcept; + annotated_ptr& operator=(T*) noexcept; annotated_ptr& operator=(const annotated_ptr&) noexcept; annotated_ptr& operator++() noexcept; @@ -421,7 +421,7 @@ annotations. a| [source,c++] ---- -annotated_ptr& operator=(const T*) noexcept; +annotated_ptr& operator=(T*) noexcept; ---- | Allows assignment from a pointer to type `T`. @@ -430,7 +430,7 @@ Allows assignment from a pointer to type `T`. a| [source,c++] ---- -annotated_ptr& operator=(annotated_ptr const&) noexcept; +annotated_ptr& operator=(const annotated_ptr &) noexcept; ---- | Allows assignment from an `annotated_ptr` with the same parameterization. From 988052c4902b2dda578620465ab9c52823f57d31 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Tue, 27 Sep 2022 09:40:43 -0700 Subject: [PATCH 27/55] add annotated_ref copy ctor --- .../proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index ef3183e32ff2c..f66d1a3cef313 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -511,6 +511,7 @@ namespace sycl::ext::oneapi::experimental { template > class annotated_ref { public: + annotated_ref(const annotated_ref&) = default; operator T() noexcept; operator const T() const noexcept; void operator=(const T &); @@ -524,6 +525,15 @@ Member Functions are described in the table below |=== |Functions |Description +// --- ROW BREAK --- +a| +[source,c++] +---- +annotated_ref(const annotated_ref&) = default; +---- +| +Compiler generated copy constructor. + // --- ROW BREAK --- a| [source,c++] From 11e7bd3ffc36d20c57b72fb1ffac72f52079948b Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Tue, 27 Sep 2022 09:47:52 -0700 Subject: [PATCH 28/55] correct const member func --- .../proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index f66d1a3cef313..d18475ab20e04 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -512,8 +512,7 @@ template > class annotated_ref { public: annotated_ref(const annotated_ref&) = default; - operator T() noexcept; - operator const T() const noexcept; + operator T() const noexcept; void operator=(const T &); }; } // namespace sycl::ext::oneapi::experimental @@ -538,8 +537,7 @@ Compiler generated copy constructor. a| [source,c++] ---- -operator T() noexcept; -operator const T() const noexcept; +operator T() const noexcept; ---- | Reads the object of type `T` that is referenced by this wrapper, applying the From 2a8c8355c593630b0f67974afe3d40d3051f080e Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Tue, 27 Sep 2022 09:53:15 -0700 Subject: [PATCH 29/55] correction to const member fn --- .../proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index d18475ab20e04..464c2f9151111 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -512,8 +512,8 @@ template > class annotated_ref { public: annotated_ref(const annotated_ref&) = default; - operator T() const noexcept; - void operator=(const T &); + operator T() const; + void operator=(const T &) const; }; } // namespace sycl::ext::oneapi::experimental ``` @@ -537,7 +537,7 @@ Compiler generated copy constructor. a| [source,c++] ---- -operator T() const noexcept; +operator T() const; ---- | Reads the object of type `T` that is referenced by this wrapper, applying the @@ -547,7 +547,7 @@ annotations when the object is loaded from memory. a| [source,c++] ---- -void operator=(const T &); +void operator=(const T &) const; ---- | Writes an object of type `T` to the location referenced by this wrapper, From 32564a4fed2ec8bdae182123b0d6b4cbce1590f2 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Wed, 28 Sep 2022 06:54:46 -0700 Subject: [PATCH 30/55] update alignment property description --- .../proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 464c2f9151111..5240b576ffaee 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -596,8 +596,9 @@ struct is_property_key_of< | Property | Description |`alignment` | -Indicates the alignment of the pointer in bytes. Alignment cannot be smaller -than the size of type `T`. +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. |`runtime_aligned` | From d3b5f3357626eef0c1e5877fb7bd0ed8fe26b9a5 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Thu, 29 Sep 2022 10:11:06 -0700 Subject: [PATCH 31/55] add issue about ctor clarification --- .../extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 3 +++ 1 file changed, 3 insertions(+) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 5240b576ffaee..657e08691319b 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -638,6 +638,9 @@ a way to annotate address spaces. That can be built with annotated_ptr. global space we may want to distinguish between general, host, and device memory spaces. +6) Ctor should clarify whether when constructing the object from two different +property lists, duplicates will exist or not. + == Revision History [cols="5,15,15,70"] From f6c72d14783af13661416a6581e643582e669b38 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Thu, 29 Sep 2022 13:01:23 -0700 Subject: [PATCH 32/55] default the copy ctor for annotated_ptr, add assignment op for annotated_ref --- .../sycl_ext_oneapi_annotated_ptr.asciidoc | 14 ++++++++++++-- 1 file changed, 12 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 657e08691319b..de6e3acd9e023 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -243,7 +243,7 @@ class annotated_ptr { explicit annotated_ptr(const annotated_ptr&, properties) noexcept; - annotated_ptr(const annotated_ptr&) noexcept; + annotated_ptr(const annotated_ptr&) noexcept = default; reference operator*() const noexcept; @@ -513,7 +513,8 @@ class annotated_ref { public: annotated_ref(const annotated_ref&) = default; operator T() const; - void operator=(const T &) const; + void operator=(const T&) const; + void operator=(const annotated_ref&) const; }; } // namespace sycl::ext::oneapi::experimental ``` @@ -553,6 +554,15 @@ void 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++] +---- +void operator=(const annotated_ref&) const; +---- +| +Assign from another `annotated_ref` object. + |=== === Properties From 7f06cbf5a0267a76085924f8dadf40fe1576266c Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Thu, 29 Sep 2022 13:05:25 -0700 Subject: [PATCH 33/55] default the copy assignment ops --- .../proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index de6e3acd9e023..c699b2d9f7ca4 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -261,7 +261,7 @@ class annotated_ptr { T* get() const noexcept; annotated_ptr& operator=(T*) noexcept; - annotated_ptr& operator=(const annotated_ptr&) noexcept; + annotated_ptr& operator=(const annotated_ptr&) noexcept = default; annotated_ptr& operator++() noexcept; annotated_ptr operator++(int) noexcept; @@ -430,7 +430,7 @@ Allows assignment from a pointer to type `T`. a| [source,c++] ---- -annotated_ptr& operator=(const annotated_ptr &) noexcept; +annotated_ptr& operator=(const annotated_ptr &) noexcept = default; ---- | Allows assignment from an `annotated_ptr` with the same parameterization. @@ -514,7 +514,7 @@ class annotated_ref { annotated_ref(const annotated_ref&) = default; operator T() const; void operator=(const T&) const; - void operator=(const annotated_ref&) const; + void operator=(const annotated_ref&) const = default; }; } // namespace sycl::ext::oneapi::experimental ``` @@ -558,7 +558,7 @@ applying the annotations when the object is stored to memory. a| [source,c++] ---- -void operator=(const annotated_ref&) const; +void operator=(const annotated_ref&) const = default; ---- | Assign from another `annotated_ref` object. From 3743ae471ed96e290efd925c1b96d7ccb3dd6c05 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Thu, 29 Sep 2022 13:47:37 -0700 Subject: [PATCH 34/55] clarify when restrict property is meaningful --- .../sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc index 5b444c9b0c356..6119015d6e711 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc @@ -91,9 +91,11 @@ implementation supports. The `restrict` property defined below is supported with the `sycl::ext::oneapi::experimental::annotated_ptr` and -`sycl::ext::oneapi::experimental::annotated_arg` classes. They are only -meaningful on the kernel argument when the kernel argument is a pointer type or -a pointer wrapper type and are not meaningful within the kernel body. +`sycl::ext::oneapi::experimental::annotated_arg` classes. It is only meaningful +on the kernel argument when the kernel argument is a pointer type or a pointer +wrapper type and is ignored for other types. + +This property is not meaningful within the kernel body. ```c++ From 7e5b2ae20b7fbba04248af099aa1e13bfe043947 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Thu, 29 Sep 2022 13:50:57 -0700 Subject: [PATCH 35/55] clarify restrict description --- .../sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc index 6119015d6e711..8e496e3c828c6 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc @@ -131,7 +131,8 @@ restrict ---- a| This is a hint to the compiler that the pointer kernel arguments marked with -this property do not alias with one another. +this property do not alias with one another with the same semantics as the C99 +`restrict` keyword. |=== -- From 10c2e9a92cc8fa33a609781ea6ac16ff739a5632 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Thu, 29 Sep 2022 14:37:04 -0700 Subject: [PATCH 36/55] clarify what types are not supported for annotations --- .../proposed/sycl_ext_oneapi_annotated_arg.asciidoc | 4 ++++ .../proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 4 ++++ 2 files changed, 8 insertions(+) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index a92e0c00ef6cc..d711459ba05a4 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -169,6 +169,10 @@ 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::sampler, and +sycl::half_type. + The properties supported with `annotated_arg` may be defined in separate extensions. diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index c699b2d9f7ca4..761fdf79683a9 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -199,6 +199,10 @@ class annotated_ptr { Properties may be specified for an `annotated_ptr` to provide semantic modification or optimization hint information. +`T` can be any type except the following types or a structure containing one of +the following types: sycl::accessor, sycl::stream, sycl::sampler, and +sycl::half_type. + Here's an example of how a property could be used: [source,c++] From 4eca4ccb050fe86b2ad4025adb693cbe8012cd3b Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Thu, 29 Sep 2022 14:50:52 -0700 Subject: [PATCH 37/55] turn restrict extension into kernel arg properties extension --- ...eapi_kernel_arg_restrict_property.asciidoc | 24 +++++++++---------- 1 file changed, 11 insertions(+), 13 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc index 8e496e3c828c6..56c97fe8eab4c 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc @@ -1,4 +1,4 @@ -= sycl_ext_oneapi_kernel_arg_restrict_property += sycl_ext_oneapi_kernel_arg_properties :source-highlighter: coderay :coderay-linenums-mode: table @@ -62,9 +62,9 @@ not rely on APIs defined in this specification.* == Overview -This extension introduces the `restrict` kernel argument property for the -classes `sycl::ext::oneapi::annotated_ptr` and -`sycl::ext::oneapi::annotated_arg`. +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 @@ -72,10 +72,10 @@ classes `sycl::ext::oneapi::annotated_ptr` and 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_RESTRICT_PROPERTY` 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 +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"] @@ -89,11 +89,9 @@ implementation supports. === `restrict` property -The `restrict` property defined below is supported with the -`sycl::ext::oneapi::experimental::annotated_ptr` and -`sycl::ext::oneapi::experimental::annotated_arg` classes. It is only meaningful -on the kernel argument when the kernel argument is a pointer type or a pointer -wrapper type and is ignored for other types. +The `restrict` property defined here is only meaningful on the kernel arguments +when the kernel argument is a pointer type or a pointer wrapper type and is +ignored for other types. This property is not meaningful within the kernel body. From eb2e830f6c40240cc963781049e447873df0906d Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Mon, 3 Oct 2022 09:07:19 -0700 Subject: [PATCH 38/55] make unsupported use case illegal, add note about stable property --- ...l_ext_intel_fpga_kernel_arg_properties.asciidoc | 14 +++++++++----- .../sycl_ext_oneapi_annotated_arg.asciidoc | 13 ++++--------- 2 files changed, 13 insertions(+), 14 deletions(-) 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 index fe70ba9dedaab..fae96a50ca151 100644 --- 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 @@ -370,12 +370,16 @@ a| stable ---- a| -Specifies that the input to the kernel will not change during the execution of -the kernel. The input can still change after all active kernel invocations have -finished. +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. -If the input is changed while the kernel is executing, the behavior is -undefined. +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++] diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index d711459ba05a4..40c8aa19f8088 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -140,21 +140,16 @@ struct MyType { }; struct MyKernel { - annotated_arg> 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 arguments 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` From e1ab16ea90daf57fef32bbae1cde8a9c65fab379 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Mon, 3 Oct 2022 09:14:56 -0700 Subject: [PATCH 39/55] clarify unsupported types with annotated_ptr --- .../proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 761fdf79683a9..9b7fbaef93ac3 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -200,8 +200,14 @@ Properties may be specified for an `annotated_ptr` to provide semantic modification or optimization hint information. `T` can be any type except the following types or a structure containing one of -the following types: sycl::accessor, sycl::stream, sycl::sampler, and -sycl::half_type. +the following types: + +* sycl::accessor +* sycl::stream +* sycl::local_accessor +* sycl::unsampled_image_accessor +* sycl::sampled_image_accessor +* sycl::half Here's an example of how a property could be used: From d7b4e4bd05b5110eaf6cbbedc59f2d48abc50487 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Mon, 3 Oct 2022 09:16:21 -0700 Subject: [PATCH 40/55] clarify unsupported types with annotated_arg --- .../proposed/sycl_ext_oneapi_annotated_arg.asciidoc | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index 40c8aa19f8088..001e39cd73c09 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -165,8 +165,14 @@ class annotated_arg { 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::sampler, and -sycl::half_type. +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. From 838fec991fb60755a43d71f531bcf7d2cb1702c4 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Mon, 3 Oct 2022 09:17:14 -0700 Subject: [PATCH 41/55] clarify annotated_arg statement --- .../extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index 001e39cd73c09..9e9b6a19b2902 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -181,8 +181,7 @@ 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 device copyable type or a legal parameter type as defined by the SYCL -specification. +must be a legal parameter type as defined by the SYCL specification. [source,c++] ---- From 5432e546100d7fd7dd01d3af0f8b7c1f361b008b Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Mon, 3 Oct 2022 09:25:30 -0700 Subject: [PATCH 42/55] remove type restrictions from annotated_ptr --- .../proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 10 ---------- 1 file changed, 10 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 9b7fbaef93ac3..c699b2d9f7ca4 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -199,16 +199,6 @@ class annotated_ptr { Properties may be specified for an `annotated_ptr` to provide semantic modification or optimization hint information. -`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 - Here's an example of how a property could be used: [source,c++] From b639fb59dd7ce85f5b28654df15b538aa424a11e Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Mon, 3 Oct 2022 13:58:58 -0700 Subject: [PATCH 43/55] rename restrict ext to kernel arg properties --- ...ty.asciidoc => sycl_ext_oneapi_kernel_arg_properties.asciidoc} | 0 1 file changed, 0 insertions(+), 0 deletions(-) rename sycl/doc/extensions/proposed/{sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc => sycl_ext_oneapi_kernel_arg_properties.asciidoc} (100%) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_properties.asciidoc similarity index 100% rename from sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_restrict_property.asciidoc rename to sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_properties.asciidoc From 63e8beda28d405ad2711e287e55c0ec3ca735dcc Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Mon, 3 Oct 2022 14:03:46 -0700 Subject: [PATCH 44/55] clarify what restrict applies to --- .../proposed/sycl_ext_oneapi_kernel_arg_properties.asciidoc | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) 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 index 56c97fe8eab4c..428770d69df6b 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_properties.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_properties.asciidoc @@ -90,8 +90,7 @@ implementation supports. === `restrict` property The `restrict` property defined here is only meaningful on the kernel arguments -when the kernel argument is a pointer type or a pointer wrapper type and is -ignored for other types. +when the kernel argument is a pointer type. It is ignored for other types. This property is not meaningful within the kernel body. From 98e53ad7005cdb517aeb45d03eed2497af4b48ae Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Tue, 4 Oct 2022 07:30:38 -0700 Subject: [PATCH 45/55] clarify ctor description, close issue --- .../sycl_ext_oneapi_annotated_ptr.asciidoc | 14 +++++++++----- 1 file changed, 9 insertions(+), 5 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index c699b2d9f7ca4..26f6fc987426c 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -333,9 +333,12 @@ explicit annotated_ptr(const annotated_ptr &Ptr, ---- | Constructs an `annotated_ptr` object. Does not allocate new storage. The -underlying pointer is initialized with `Ptr`. `PropertyListU` and -`PropertyListV` will be merged together to construct `PropertyListT`. If any -common properties have different values then an error will be issued. +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*`. @@ -648,8 +651,9 @@ a way to annotate address spaces. That can be built with annotated_ptr. global space we may want to distinguish between general, host, and device memory spaces. -6) Ctor should clarify whether when constructing the object from two different -property lists, duplicates will exist or not. +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. == Revision History From 0619a6626e0020072861a17cff0d287859a8c775 Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Tue, 4 Oct 2022 07:35:59 -0700 Subject: [PATCH 46/55] remove runtime_aligned property, open an issue --- .../sycl_ext_oneapi_annotated_ptr.asciidoc | 19 ++++--------------- 1 file changed, 4 insertions(+), 15 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 26f6fc987426c..597023aac9fe7 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -582,25 +582,16 @@ struct alignment_key { using value_t = property_value>; }; -struct runtime_aligned_key { - using value_t = property_value; -}; - template inline constexpr alignment_key::value_t alignment; -inline constexpr runtime_aligned_key::value_t runtime_aligned; template<> struct is_property_key : std::true_type {}; -template<> -struct is_property_key : std::true_type {}; template struct is_property_key_of< alignment_key, annotated_ptr> : std::true_type {}; template -struct is_property_key_of< - runtime_aligned_key, annotated_ptr> : std::true_type {}; } // namespace sycl::ext::oneapi::experimental ``` -- @@ -612,12 +603,6 @@ struct is_property_key_of< 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. - -|`runtime_aligned` -| -Informs the compiler that the pointer has at least the default alignment for a -pointer allocated through the SYCL runtime. This is always safe to apply to any -pointer returned by a SYCL memory allocation function. |==== -- @@ -653,8 +638,12 @@ 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"] From f14819047cbdd1062af520d6cbfc3da3f647203f Mon Sep 17 00:00:00 2001 From: Abhishek Tiwari Date: Tue, 4 Oct 2022 07:38:43 -0700 Subject: [PATCH 47/55] clarify restrict definition --- .../sycl_ext_oneapi_kernel_arg_properties.asciidoc | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) 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 index 428770d69df6b..a864bb8d4e22a 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_properties.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_properties.asciidoc @@ -127,9 +127,10 @@ a| restrict ---- a| -This is a hint to the compiler that the pointer kernel arguments marked with -this property do not alias with one another with the same semantics as the C99 -`restrict` keyword. +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. |=== -- From fe8f7439f8b877e2d36cbfc4bac7bb021ef1ca9e Mon Sep 17 00:00:00 2001 From: "Chen, Brox" Date: Tue, 1 Nov 2022 09:37:11 -0700 Subject: [PATCH 48/55] update the description of the annotated_arg/ptr spec --- .../sycl_ext_oneapi_annotated_arg.asciidoc | 65 +++++++++++++++++-- .../sycl_ext_oneapi_annotated_ptr.asciidoc | 31 +++++++-- 2 files changed, 86 insertions(+), 10 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index 9e9b6a19b2902..aa5a7de6d4827 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -189,8 +189,16 @@ namespace sycl::ext::oneapi::experimental { template > class annotated_arg { public: - annotated_arg(); - annotated_arg(const T& v_); + annotated_arg() noexcept; + annotated_arg(const T& v_, const properties &P = properties{}) noexcept; + template + annotated_arg(T *Ptr, PropertyValueTs... props) noexcept; + + template explicit annotated_arg( + const annotated_arg&) noexcept; + template + explicit annotated_arg(const annotated_arg&, + properties) noexcept; annotated_arg(const annotated_arg&) = default; annotated_arg& operator=(annotated_arg&) = default; @@ -228,7 +236,7 @@ class annotated_arg { a| [source,c++] ---- -annotated_arg(); +annotated_arg() noexcept; ---- | Not available in device code. Constructs an `annotated_arg` object which is default initialized. @@ -237,11 +245,60 @@ Constructs an `annotated_arg` object which is default initialized. a| [source,c++] ---- -annotated_arg(const T& v_); +annotated_arg(const T& v_, const properties &P = properties{}) noexcept; ---- | Not available in device code. Constructs an `annotated_arg` object from the input object `v_`. +The new property set `PropertyListT` contains all properties in `P`. +If there were duplicate property in `P`, the value of this property must be same. + +// --- ROW BREAK --- +a| +[source,c++] +---- +template +annotated_arg(const T& v_, const PropertyValueTs... props) noexcept; +---- +| +Constructs an `annotated_arg` object from the input object `v_`. + +The new property set `PropertyListT` contains all properties listed in `props`. +If there were duplicate property in the list of `props`, the value of this property must be same. + +// --- ROW BREAK --- +a| +[source,c++] +---- +template explicit annotated_arg( + const annotated_arg &ConvertFrom); +---- +| +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`. + +`T2` must be implicitly convertible to `T`. + +// --- ROW BREAK --- +a| +[source,c++] +---- +template +explicit annotated_arg(const annotated_arg& v_, + properties P) noexcept; +---- +| +Constructs the `annotated_arg` object from the input object `v_`. + +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++] diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 597023aac9fe7..b0bd4ec2b1c6e 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -210,6 +210,11 @@ using namespace sycl::ext::oneapi::experimental; // 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 + // property 'alignment' and 'awidth'. + // a properties object is deducted from the list the of property values + auto arg_b = annotated_ptr(ptr, alignment<4>, awidth<32>); + q.submit([=]{ ... *arg_a = (*arg_a) * 2; @@ -236,6 +241,8 @@ class annotated_ptr { annotated_ptr() noexcept; annotated_ptr(T *Ptr, const properties &P = properties{}) noexcept; + template + annotated_ptr(T *Ptr, PropertyValueTs... props) noexcept; template explicit annotated_ptr( const annotated_ptr&) noexcept; @@ -278,7 +285,6 @@ class annotated_ptr { static constexpr /*unspecified*/ get_property(); }; } // namespace sycl::ext::oneapi::experimental - ---- [frame="topbot",options="header"] @@ -303,11 +309,24 @@ annotated_ptr(T *Ptr, const properties &P = properties{}) noexcep ---- | Constructs an `annotated_ptr` object. Does not allocate new storage. The -underlying pointer is initialized with `Ptr`. `P` is used to specify the -`PropertyListT` type on the class. +underlying pointer is initialized with `Ptr`. + +The new property set `PropertyListT` contains all properties in `P`. +If there were duplicate property in `P`, the value of this property must be same. + +// --- 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` must contain all properties from `P`, -and if any common property takes a value, the value must be the same. +The new property set `PropertyListT` contains all properties listed in `props`. +If there were duplicate property in the list of `props`, the value of this property must be same. // --- ROW BREAK --- a| @@ -654,4 +673,4 @@ the alignment is set up. |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* -|======================================== \ No newline at end of file +|======================================== From 1ea381adb55bdd6cb6f44449e425346cdcf7c303 Mon Sep 17 00:00:00 2001 From: "Chen, Brox" Date: Wed, 9 Nov 2022 13:39:43 -0800 Subject: [PATCH 49/55] address comments and change wording --- .../proposed/sycl_ext_oneapi_annotated_arg.asciidoc | 8 ++++---- .../proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 8 ++++---- 2 files changed, 8 insertions(+), 8 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index aa5a7de6d4827..b8b6363e326da 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -192,7 +192,7 @@ class annotated_arg { annotated_arg() noexcept; annotated_arg(const T& v_, const properties &P = properties{}) noexcept; template - annotated_arg(T *Ptr, PropertyValueTs... props) noexcept; + annotated_arg(const T& v_, PropertyValueTs... props) noexcept; template explicit annotated_arg( const annotated_arg&) noexcept; @@ -251,7 +251,7 @@ annotated_arg(const T& v_, const properties &P = properties{}) no Constructs an `annotated_arg` object from the input object `v_`. The new property set `PropertyListT` contains all properties in `P`. -If there were duplicate property in `P`, the value of this property must be same. +If there are duplicate properties present in the property list of `P`, the values of the duplicate properties must be the same. // --- ROW BREAK --- a| @@ -263,8 +263,8 @@ annotated_arg(const T& v_, const PropertyValueTs... props) noexcept; | Constructs an `annotated_arg` object from the input object `v_`. -The new property set `PropertyListT` contains all properties listed in `props`. -If there were duplicate property in the list of `props`, the value of this property must be same. +The new property set `PropertyListT` contains all properties in `P`. +If there are duplicate properties present in the property list of `P`, the values of the duplicate properties must be the same. // --- ROW BREAK --- a| diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index b0bd4ec2b1c6e..7078e552bc2db 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -211,7 +211,7 @@ using namespace sycl::ext::oneapi::experimental; auto arg_a = annotated_ptr(ptr, properties{alignment<4>}); // alignment in bytes and address bus width of the pointer specified using the - // property 'alignment' and 'awidth'. + // properties 'alignment' and 'awidth'. // a properties object is deducted from the list the of property values auto arg_b = annotated_ptr(ptr, alignment<4>, awidth<32>); @@ -312,7 +312,7 @@ 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`. -If there were duplicate property in `P`, the value of this property must be same. +If there are duplicate properties present in the property list of `P`, the values of the duplicate properties must be the same. // --- ROW BREAK --- a| @@ -325,8 +325,8 @@ 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 listed in `props`. -If there were duplicate property in the list of `props`, the value of this property must be same. +The new property set `PropertyListT` contains all properties in `P`. +If there are duplicate properties present in the property list of `P`, the values of the duplicate properties must be the same. // --- ROW BREAK --- a| From 16f764d981023b8846ad64ebfe4f1b567bf4406e Mon Sep 17 00:00:00 2001 From: Brox Chen Date: Thu, 10 Nov 2022 09:33:36 -0500 Subject: [PATCH 50/55] Update the description of the annotated_arg/ptr spec (#1) * update the description of the annotated_arg/ptr spec * address comments and change wording --- .../sycl_ext_oneapi_annotated_arg.asciidoc | 65 +++++++++++++++++-- .../sycl_ext_oneapi_annotated_ptr.asciidoc | 31 +++++++-- 2 files changed, 86 insertions(+), 10 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index 9e9b6a19b2902..b8b6363e326da 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -189,8 +189,16 @@ namespace sycl::ext::oneapi::experimental { template > class annotated_arg { public: - annotated_arg(); - annotated_arg(const T& v_); + annotated_arg() noexcept; + annotated_arg(const T& v_, const properties &P = properties{}) noexcept; + template + annotated_arg(const T& v_, PropertyValueTs... props) noexcept; + + template explicit annotated_arg( + const annotated_arg&) noexcept; + template + explicit annotated_arg(const annotated_arg&, + properties) noexcept; annotated_arg(const annotated_arg&) = default; annotated_arg& operator=(annotated_arg&) = default; @@ -228,7 +236,7 @@ class annotated_arg { a| [source,c++] ---- -annotated_arg(); +annotated_arg() noexcept; ---- | Not available in device code. Constructs an `annotated_arg` object which is default initialized. @@ -237,11 +245,60 @@ Constructs an `annotated_arg` object which is default initialized. a| [source,c++] ---- -annotated_arg(const T& v_); +annotated_arg(const T& v_, const properties &P = properties{}) noexcept; ---- | Not available in device code. Constructs an `annotated_arg` object from the input object `v_`. +The new property set `PropertyListT` contains all properties in `P`. +If there are duplicate properties present in the property list of `P`, the values of the duplicate properties must be the same. + +// --- ROW BREAK --- +a| +[source,c++] +---- +template +annotated_arg(const T& v_, const PropertyValueTs... props) noexcept; +---- +| +Constructs an `annotated_arg` object from the input object `v_`. + +The new property set `PropertyListT` contains all properties in `P`. +If there are duplicate properties present in the property list of `P`, the values of the duplicate properties must be the same. + +// --- ROW BREAK --- +a| +[source,c++] +---- +template explicit annotated_arg( + const annotated_arg &ConvertFrom); +---- +| +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`. + +`T2` must be implicitly convertible to `T`. + +// --- ROW BREAK --- +a| +[source,c++] +---- +template +explicit annotated_arg(const annotated_arg& v_, + properties P) noexcept; +---- +| +Constructs the `annotated_arg` object from the input object `v_`. + +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++] diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 597023aac9fe7..7078e552bc2db 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -210,6 +210,11 @@ using namespace sycl::ext::oneapi::experimental; // 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 deducted from the list the of property values + auto arg_b = annotated_ptr(ptr, alignment<4>, awidth<32>); + q.submit([=]{ ... *arg_a = (*arg_a) * 2; @@ -236,6 +241,8 @@ class annotated_ptr { annotated_ptr() noexcept; annotated_ptr(T *Ptr, const properties &P = properties{}) noexcept; + template + annotated_ptr(T *Ptr, PropertyValueTs... props) noexcept; template explicit annotated_ptr( const annotated_ptr&) noexcept; @@ -278,7 +285,6 @@ class annotated_ptr { static constexpr /*unspecified*/ get_property(); }; } // namespace sycl::ext::oneapi::experimental - ---- [frame="topbot",options="header"] @@ -303,11 +309,24 @@ annotated_ptr(T *Ptr, const properties &P = properties{}) noexcep ---- | Constructs an `annotated_ptr` object. Does not allocate new storage. The -underlying pointer is initialized with `Ptr`. `P` is used to specify the -`PropertyListT` type on the class. +underlying pointer is initialized with `Ptr`. + +The new property set `PropertyListT` contains all properties in `P`. +If there are duplicate properties present in the property list of `P`, the values of the duplicate properties must be the same. + +// --- 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` must contain all properties from `P`, -and if any common property takes a value, the value must be the same. +The new property set `PropertyListT` contains all properties in `P`. +If there are duplicate properties present in the property list of `P`, the values of the duplicate properties must be the same. // --- ROW BREAK --- a| @@ -654,4 +673,4 @@ the alignment is set up. |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* -|======================================== \ No newline at end of file +|======================================== From ec6f9dda39b26c45864d74caf8b0426277726ab6 Mon Sep 17 00:00:00 2001 From: "Chen, Brox" Date: Mon, 21 Nov 2022 14:29:41 -0800 Subject: [PATCH 51/55] fix errors, added unsupported cases in annotated_ptr page --- .../sycl_ext_oneapi_annotated_arg.asciidoc | 6 ++-- .../sycl_ext_oneapi_annotated_ptr.asciidoc | 28 ++++++++++++++++--- 2 files changed, 27 insertions(+), 7 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index b8b6363e326da..c106fc70696a8 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -191,7 +191,7 @@ class annotated_arg { public: annotated_arg() noexcept; annotated_arg(const T& v_, const properties &P = properties{}) noexcept; - template + template annotated_arg(const T& v_, PropertyValueTs... props) noexcept; template explicit annotated_arg( @@ -245,7 +245,7 @@ Constructs an `annotated_arg` object which is default initialized. a| [source,c++] ---- -annotated_arg(const T& v_, const properties &P = properties{}) noexcept; +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_`. @@ -257,7 +257,7 @@ If there are duplicate properties present in the property list of `P`, the value a| [source,c++] ---- -template +template annotated_arg(const T& v_, const PropertyValueTs... props) noexcept; ---- | diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 7078e552bc2db..62f1f8924d776 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -181,6 +181,26 @@ 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` @@ -240,8 +260,8 @@ class annotated_ptr { using reference = annotated_ref; annotated_ptr() noexcept; - annotated_ptr(T *Ptr, const properties &P = properties{}) noexcept; - template + annotated_ptr(T *Ptr, const PropertyListT &P = properties{}) noexcept; + template annotated_ptr(T *Ptr, PropertyValueTs... props) noexcept; template explicit annotated_ptr( @@ -305,7 +325,7 @@ underlying pointer is initialized to `nullptr`. a| [source,c++] ---- -annotated_ptr(T *Ptr, const properties &P = properties{}) noexcept; +annotated_ptr(T *Ptr, const PropertyListT &P = properties{}) noexcept; ---- | Constructs an `annotated_ptr` object. Does not allocate new storage. The @@ -318,7 +338,7 @@ If there are duplicate properties present in the property list of `P`, the value a| [source,c++] ---- -template +template annotated_ptr(T *Ptr, const PropertyValueTs... props) noexcept; ---- | From 3fdeabd5b39851bb162ed5f78a2bc409f8e44722 Mon Sep 17 00:00:00 2001 From: "Chen, Brox" Date: Mon, 28 Nov 2022 13:35:05 -0800 Subject: [PATCH 52/55] added template deduction guide to annotated_arg/ptr spec --- .../proposed/sycl_ext_oneapi_annotated_arg.asciidoc | 9 +++++++++ .../proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 9 +++++++++ 2 files changed, 18 insertions(+) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index c106fc70696a8..a1281fa9e7d12 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -225,6 +225,15 @@ class annotated_arg { template static constexpr /*unspecified*/ get_property(); }; + +//Deduction guides +template +annotated_arg(T, Args... args) -> annotated_arg>; + +template +annotated_arg(annotated_arg, otherProp other) + -> annotated_arg>; + } // namespace sycl::ext::oneapi::experimental ---- diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 62f1f8924d776..43e888e7cc6a3 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -304,6 +304,15 @@ class annotated_ptr { template static constexpr /*unspecified*/ get_property(); }; + +//Deduction guides +template +annotated_ptr(T, Args... args) -> annotated_ptr>; + +template +annotated_ptr(annotated_arg, otherProp other) + -> annotated_ptr>; + } // namespace sycl::ext::oneapi::experimental ---- From 9b1d33e081f5dd8ba846ba5516ec7df580fb73b8 Mon Sep 17 00:00:00 2001 From: "Chen, Brox" Date: Fri, 2 Dec 2022 12:01:05 -0800 Subject: [PATCH 53/55] added template deductions, fixed minor errors --- .../sycl_ext_oneapi_annotated_arg.asciidoc | 55 +++++++------------ .../sycl_ext_oneapi_annotated_ptr.asciidoc | 31 ++++++----- 2 files changed, 38 insertions(+), 48 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index a1281fa9e7d12..41265c6fc2765 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -175,13 +175,15 @@ the following types: * 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. +must be a legal parameter type as defined by the SYCL specification. Given `annotated_arg`, +`T` must be a trivially copy-able type. [source,c++] ---- @@ -190,7 +192,7 @@ template > class annotated_arg { public: annotated_arg() noexcept; - annotated_arg(const T& v_, const properties &P = properties{}) noexcept; + annotated_arg(const T& v_, const PropertyListT &P = properties{}) noexcept; template annotated_arg(const T& v_, PropertyValueTs... props) noexcept; @@ -204,19 +206,13 @@ class annotated_arg { annotated_arg& operator=(annotated_arg&) = default; // Conversion operator to convert to the underlying type - operator T&() noexcept; - operator const T&() const noexcept; + operator T() noexcept; + operator const T() 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; - // Available if the operator() is valid for objects of type T, return - // type will match the return type of - // template T::operator()(Args... args) - template /* ... */ operator()(Args... args) noexcept; - template /* ... */ operator()(Args... args) const noexcept; - template static constexpr bool has_property(); @@ -228,12 +224,16 @@ class annotated_arg { //Deduction guides template -annotated_arg(T, Args... args) -> annotated_arg>; +annotated_arg(T, Args... args) -> + annotated_arg>; -template -annotated_arg(annotated_arg, otherProp other) - -> annotated_arg>; +template +annotated_arg(T, properties>) -> + annotated_arg>; +template +annotated_arg(annotated_arg, properties>) -> + annotated_arg>>; } // namespace sycl::ext::oneapi::experimental ---- @@ -260,7 +260,6 @@ annotated_arg(const T& v_, const PropertyListT &P = properties{}) noexcept; Constructs an `annotated_arg` object from the input object `v_`. The new property set `PropertyListT` contains all properties in `P`. -If there are duplicate properties present in the property list of `P`, the values of the duplicate properties must be the same. // --- ROW BREAK --- a| @@ -269,11 +268,10 @@ a| template annotated_arg(const T& v_, const PropertyValueTs... props) noexcept; ---- -| +| Not available in device code. Constructs an `annotated_arg` object from the input object `v_`. -The new property set `PropertyListT` contains all properties in `P`. -If there are duplicate properties present in the property list of `P`, the values of the duplicate properties must be the same. +The new property set `PropertyListT` contains all properties listed in `props`. // --- ROW BREAK --- a| @@ -282,7 +280,7 @@ a| template explicit annotated_arg( const annotated_arg &ConvertFrom); ---- -| +| 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`. @@ -297,7 +295,7 @@ template explicit annotated_arg(const annotated_arg& v_, properties P) noexcept; ---- -| +| Not available in device code. Constructs the `annotated_arg` object from the input object `v_`. The new `PropertyListT` is the union of all properties contained within @@ -307,13 +305,13 @@ properties with the same value (or no value) are allowed. `T2` must be implicitly convertible to `T`. - // --- ROW BREAK --- a| [source,c++] ---- annotated_arg(const annotated_arg&) = default; ---- +| Not available in device code. Compiler generated copy constructor. // --- ROW BREAK --- @@ -322,7 +320,7 @@ a| ---- annotated_arg& operator=(annotated_arg&) = default; ---- -Compiler generated assignment operator. +| Compiler generated assignment operator. // --- ROW BREAK --- a| @@ -343,17 +341,6 @@ a| Available if the `operator[]` is valid for objects of type `T`. This function will call the subscript operator defined for `T`. -// --- ROW BREAK --- -a| -[source,c++] ----- -template /* ... */ operator()(Args... args) noexcept; -template /* ... */ operator()(Args... args) const noexcept; ----- -| -Available if the `operator()` is valid for objects of type `T`. This function -will call the 'call operator' defined for `T`. - // --- ROW BREAK --- a| [source,c++] diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 43e888e7cc6a3..555c27baf51f9 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -217,7 +217,8 @@ 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. +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: @@ -307,12 +308,16 @@ class annotated_ptr { //Deduction guides template -annotated_ptr(T, Args... args) -> annotated_ptr>; +annotated_arg(T, Args... args) -> + annotated_arg>; -template -annotated_ptr(annotated_arg, otherProp other) - -> annotated_ptr>; +template +annotated_arg(T, properties>) -> + annotated_arg>; +template +annotated_arg(annotated_arg, properties>) -> + annotated_arg>>; } // namespace sycl::ext::oneapi::experimental ---- @@ -341,7 +346,6 @@ 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`. -If there are duplicate properties present in the property list of `P`, the values of the duplicate properties must be the same. // --- ROW BREAK --- a| @@ -355,7 +359,6 @@ 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`. -If there are duplicate properties present in the property list of `P`, the values of the duplicate properties must be the same. // --- ROW BREAK --- a| @@ -384,9 +387,9 @@ 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. +`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*`. @@ -564,8 +567,8 @@ class annotated_ref { public: annotated_ref(const annotated_ref&) = default; operator T() const; - void operator=(const T&) const; - void operator=(const annotated_ref&) const = default; + annotated_ref& operator=(const T&) const; + annotated_ref& operator=(const annotated_ref&) const = default; }; } // namespace sycl::ext::oneapi::experimental ``` @@ -599,7 +602,7 @@ annotations when the object is loaded from memory. a| [source,c++] ---- -void operator=(const T &) const; +annotated_ref& operator=(const T &) const; ---- | Writes an object of type `T` to the location referenced by this wrapper, @@ -609,7 +612,7 @@ applying the annotations when the object is stored to memory. a| [source,c++] ---- -void operator=(const annotated_ref&) const = default; +annotated_ref& operator=(const annotated_ref&) const = default; ---- | Assign from another `annotated_ref` object. From 1b2a3fbd7a263651cf4f6080d08dc195913a93f9 Mon Sep 17 00:00:00 2001 From: "Chen, Brox" Date: Fri, 2 Dec 2022 13:45:48 -0800 Subject: [PATCH 54/55] fixed template deduction regarding Joe's comment --- .../sycl_ext_oneapi_annotated_arg.asciidoc | 20 ++++++++----------- .../sycl_ext_oneapi_annotated_ptr.asciidoc | 16 ++++++--------- 2 files changed, 14 insertions(+), 22 deletions(-) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index 41265c6fc2765..61feb0648039c 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -223,17 +223,13 @@ class annotated_arg { }; //Deduction guides -template -annotated_arg(T, Args... args) -> - annotated_arg>; +template +annotated_arg(T, PropertyValueTs... values) -> + annotated_arg; -template -annotated_arg(T, properties>) -> - annotated_arg>; - -template -annotated_arg(annotated_arg, properties>) -> - annotated_arg>>; +template +annotated_arg(annotated_arg, PropertiesB>) -> + annotated_arg; } // namespace sycl::ext::oneapi::experimental ---- @@ -326,8 +322,8 @@ annotated_arg& operator=(annotated_arg&) = default; a| [source,c++] ---- -operator T&() noexcept; -operator const T&() const noexcept; +operator T() noexcept; +operator const T() const noexcept; ---- | Implicit conversion to a reference to the underlying type `T`. diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 555c27baf51f9..02830ec24d8c5 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -307,17 +307,13 @@ class annotated_ptr { }; //Deduction guides -template -annotated_arg(T, Args... args) -> - annotated_arg>; +template +annotated_ptr(T, PropertyValueTs... values) -> + annotated_ptr; -template -annotated_arg(T, properties>) -> - annotated_arg>; - -template -annotated_arg(annotated_arg, properties>) -> - annotated_arg>>; +template +annotated_ptr(annotated_ptr, PropertiesB>) -> + annotated_ptr; } // namespace sycl::ext::oneapi::experimental ---- From 36d521e0edef3fab4444e2964c24aa5f10879f63 Mon Sep 17 00:00:00 2001 From: "Chen, Brox" Date: Wed, 14 Dec 2022 08:49:52 -0800 Subject: [PATCH 55/55] fixed copyrigth, remove const from conversion fucntions, simplify example use case --- .../sycl_ext_intel_fpga_kernel_arg_properties.asciidoc | 10 +++++----- .../proposed/sycl_ext_oneapi_annotated_arg.asciidoc | 6 +++--- .../proposed/sycl_ext_oneapi_annotated_ptr.asciidoc | 6 +++--- .../sycl_ext_oneapi_kernel_arg_properties.asciidoc | 2 +- 4 files changed, 12 insertions(+), 12 deletions(-) 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 index fae96a50ca151..6852a3f4c4a26 100644 --- 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 @@ -73,19 +73,19 @@ Some examples of the syntax are shown below. [source,c++] ---- auto ptr = ... -auto arg_a = annotated_ptr( - ptr, properties{buffer_location<1>, awidth<32>, dwidth<64>}); +annotated_ptr arg_a{ + ptr, properties{buffer_location<1>, awidth<32>, dwidth<64>}}; ... auto ptr2 = ... -auto arg_b = annotated_arg( - ptr2, properties{buffer_location<2>, awidth<32>, dwidth<64>}); +annotated_arg arg_b{ + ptr2, properties{buffer_location<2>, awidth<32>, dwidth<64>}}; ... int val = 5; -auto arg_c = annotated_arg(val, properties{register_map, stable}); +annotated_arg arg_c{val, properties{register_map, stable}}; ... diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc index 61feb0648039c..eb8d8e045aa5c 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_arg.asciidoc @@ -23,7 +23,7 @@ == Notice [%hardbreaks] -Copyright (c) 2022-2022 Intel Corporation. All rights reserved. +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 @@ -183,7 +183,7 @@ The section below describes the constructors and member functions for 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 trivially copy-able type. +`T` must be a device copy-able type. [source,c++] ---- @@ -207,7 +207,7 @@ class annotated_arg { // Conversion operator to convert to the underlying type operator T() noexcept; - operator const T() const noexcept; + operator T() 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) diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc index 02830ec24d8c5..3fe61ef389234 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_annotated_ptr.asciidoc @@ -24,7 +24,7 @@ == Notice [%hardbreaks] -Copyright (C) 2022-2022 Intel Corporation. All rights reserved. +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 @@ -233,7 +233,7 @@ using namespace sycl::ext::oneapi::experimental; // alignment in bytes and address bus width of the pointer specified using the // properties 'alignment' and 'awidth'. - // a properties object is deducted from the list the of property values + // a properties object is deduced from the list the of property values auto arg_b = annotated_ptr(ptr, alignment<4>, awidth<32>); q.submit([=]{ @@ -284,7 +284,7 @@ class annotated_ptr { // Implicit conversion is not supported operator T*() noexcept = delete; - operator const T*() const noexcept = delete; + operator T*() const noexcept = delete; T* get() const noexcept; 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 index a864bb8d4e22a..b33fe8b559b21 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_properties.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_arg_properties.asciidoc @@ -24,7 +24,7 @@ == Notice [%hardbreaks] -Copyright (C) 2022-2022 Intel Corporation. All rights reserved. +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