diff --git a/sycl/doc/extensions/accessor_properties/SYCL_INTEL_accessor_properties.asciidoc b/sycl/doc/extensions/accessor_properties/SYCL_INTEL_accessor_properties.asciidoc new file mode 100644 index 0000000000000..b03d313ff6727 --- /dev/null +++ b/sycl/doc/extensions/accessor_properties/SYCL_INTEL_accessor_properties.asciidoc @@ -0,0 +1,412 @@ += SYCL_INTEL_accessor_properties + +== Introduction +This extension introduces a mechanism to add compile-time-constant properties to property_list and modifies accessor to be templated on its property list. +The goal of these two changes is to enable information about properties to propagate to the device compiler and thereby enable additional optimization of kernel code. +This extension introduces two new compile-time-constant accessor properties that make use of this mechanism: no_alias and no_offset. + +== Contributors +Joe Garvey, Intel + +Roland Schulz, Intel + +Ilya Burylov, Intel + +Michael Kinsner, Intel + +John Pennycook, Intel + +== Notice +Copyright (c) 2019-2020 Intel Corporation. All rights reserved. + +== Status + +Working Draft - *DO NOT SHIP* + +== Version + +Built On: {docdate} + +Revision: 1 + +== Dependencies + +This extension is written against the SYCL 2020 pre-provisional specification. + +== Overview + +The accessor class is a key vehicle through which compile-time-constant information about a kernel argument can be passed to a SYCL device compiler. +This class already has many template parameters and, if left unchecked, that number will continue to grow as vendors add template parameters through their own extensions. +The accessor constructors already take a property_list argument through which a user can pass additional optimization hints or directives that seems intended to resolve this scalability problem, +but the properties passed do not affect the type of the property_list or of the accessor and thus they can't be reliably determined at compile time. +This extension resolves this difficiency by introducing a mechanism for annotating additional compile-time-constant information on accessors through the existing property_list. +The goals of this extension are: + +. Enable compile-time-constant properties to be applied to an accessor +. Preserve the brevity of accessor construction that was introduced in SYCL 2020 +. Achieve both of these in a scalable way that allows vendors to easily add their own new properties + +To these ends, this extension templates property_list on the types of its properties and accessor on the type of its property_list. +It introduces a new property-querying struct to determine if a particular property is compile-time constant. +And it requires that accessors and property_lists are implicitly convertible to other accessors and property_lists with the same compile-time-constant properties. + +Adding a new compile-time-constant property consists of adding a class for it (as one has to do for any property) and ensuring that the appropriate property functions, including the new is_compile_time_property struct, are specialized for the new property. + +This extension introduces two compile-time-constant properties: no_offset and no_alias. +no_offset indicates that the corresponding accessor will never contain an offset. +no_alias indicates that the corresponding accessor will not alias with any other accessor used in the kernel. + +=== Examples +Due to deduction guides, users should rarely need to specify the template parameters of property_list or accessor directly. +For example, declaring a read/write accessor for buffer B and handler h will remain: + +```c++ +accessor A1(B, h); // A default empty property_list is used +``` + +An accessor that has the the no_alias property can be created as follows: + +```c++ +accessor A2(B, h, property_list{no_alias}); +``` + +Here are some examples of legal and illegal conversions: + +```c++ +accessor A2(B, h, property_list{no_alias}); +accessor A3(B, h, property_list{no_alias, noinit}); +accessor A4(B, h, property_list{no_alias, no_offset}); +accessor A5(B, h, proeprty_list{no_offset, no_alias}); +A2 = A3; // Legal because noinit is a runtime property and A2 and A3 otherwise have the same properties +A3 = A2; // Legal for the same reasons as the previous line +A2 = A4; // Illegal as A2 doesn't have the compile-time-constant property no_offset +A5 = A4; // Legal as the order of properties doesn't matter +``` + +Writing a function that takes an arbitrary property_list can be done as follows: + +```c++ +template +void my_func(property_list props); +``` + +And writing one that only accepts specific compile-time properties can be accomplished by explicitly listing those properties: + +```c++ +void my_func2(property_list props); + +... +my_func2(property_list{noinit}); // Illegal because no_offset was not specified +my_func2(property_list{no_offset, no_alias}); // Illegal because no_alias is a compile-time property that wasn't specified in the function declaration +my_func2(property_list{no_offset, noinit}); // Legal because no_offset was specified and noinit is not a compile-time property +``` + +Similarly, a function that takes an accessor does not need to specify the details of the property_list parameter unless it should only accept property_lists with specific compile-time properties: + +```c++ +template +void my_func3(accessor a); + +template +void my_func4(accessor> a); +... +my_func3(accessor{b, h, sycl::property_list{no_alias}}); // Legal. my_func3 accepts any kind of compile-time property +my_func4(accessor{b, h, sycl::property_list{no_alias}}); // Legal. my_func4 requires no_alias +my_func4(accessor{b, h, sycl::property_list{no_alias, no_offset}}); // Illegal. no_offset is a compile-time property +my_func4(accessor{b, h, sycl::property_list{no_alias, noinit}}); // Legal. noinit is a run-time property +``` + +== Modifications to the SYCL 2020 Pre-Provisional Specification + +=== Section 4.5.5.1 Properties interface + +In the text, replace uses of the term "property_list class" with "property_list class template". + +At the end of the text, add the sentence: + +A property_list type must be implicitly convertible to any other property_list type with the same compile-time-constant properties. + +Rewrite the code listing as follows: + +```c++ +namespace sycl { + +template +struct is_property; + +template +struct is_property_of; + +// New struct for querying whether a property is compile-time-constant +template +struct is_compile_time_property; + +class T { + ... + + template + bool has_property() const; + + template + propertyT get_property() const; + + // New member function for the runtime classes + // Only available for accessor as the other classes can't have compile-time properties + template + static constexpr bool has_compile_time_property(); + + // There is no get_property variant for compile-time-constant properties because all information about them is already embedded in the type + // For templated compile-time constant properties, implementations may want to add functions/structs to query if any specialization of a particular class template is present and if so to return it + ... +}; + +// Now templated +template +class property_list { + public: + // No longer templated (properties come from the class template) + property_list(properties... props); + + // Conversion operator + // Enabled only if newProperties... and properties... contain the same compile-time-constant properties (in any order and any number of times) + // Can also be disabled selectively for compile-time properties that are incompatible as defined by the definitions of those properties + template + operator property_list () const; +}; +} // namespace sycl +``` + +Add a row to Table 4.5: Traits for properties as follows: + +-- +[options="header"] +|==== +| Member function | Description +a| +```c++ +template +struct is_compile_time_property; +``` | An explicit specialization of is_compile_time_property that inherits from std::true_type must be provided for each compile-time constant property, where propertyT is the class defining the property. +All other specializations of is_compile_time_property must inherit from std::false_type. +|==== +-- + +Add a row to Table 4.6: Common member functions of the SYCL property interface as follows: + +-- +[options="header"] +|==== +| Member function | Description +a| +```c++ +template +static constexpr bool has_compile_time_property(); +``` | Returns true if T was constructed with the compile-time-constant property specified by propertyT. Return false if it was not. +|=== +-- + +NOTE: I don't think we can automatically apply the noOffset property to accessors created from the get_access functions that don't take an offset because if this buffer is a sub-buffer it might already have an offset even if one isn't supplied to its get_access call. Can we distinguish buffers from sub-buffers at compile time so that we can resolve this? + +=== Section 4.7.6 Accessors + +Introduce a new template parameter to the accessor class template at the end of the list: + +* A property_list class template to encode the compile-time-constant properties of the accessor. +If two accessor specializations differ only in their property_list template parameters and those property_list types are implicitly convertible then the accessors specializations must be implicitly convertible. + +=== Section 4.7.6.6 Accessor declaration + +Modify the accessor declaration to add an additional template parameter as follows: + +```c++ +namespace sycl { +template ? access_mode::read + : access_mode::read_write), + target accessTarget = target::global_buffer, + access::placeholder isPlaceholder = access::placeholder::false_t, // Deprecated in SYCL 2020 + typename property_listT = property_list<> +> +class accessor; +... +``` + +=== Section 4.7.6.8 Implicit accessor conversions + +At the end of this section, add the following: + +Any accessor types that are identical except for their non-compile-time-constant properties can be implicitly converted to one another. + +=== Section 4.7.6.9.1 Device buffer accessor interface +Modify the code listing to introduce an additional template parameter: + +```c++ +namespace sycl { +template > +class accessor { +... +``` + +Modify the code listing to make all the accessor constructors that take a property_list take a templated property_list that matches the type the accessor is templated with: + +```c++ + /* All constructors are only available when: (std::is_same>::value == true) + + /* Available only when: (dimensions == 0) */ + template + accessor(buffer &bufferRef, + const property_list &propList = {}); + + /* Available only when: (dimensions == 0) */ + template + accessor(buffer &bufferRef, + handler &commandGroupHandlerRef, const property_list &propList = {}); + + /* Available only when: (dimensions > 0) */ + template + accessor(buffer &bufferRef, + const property_list &propList = {}); + + /* Available only when: (dimensions > 0) */ + template + accessor(buffer &bufferRef, TagT tag, + const property_list &propList = {}); + + /* Available only when: (dimensions > 0) */ + template + accessor(buffer &bufferRef, + handler &commandGroupHandlerRef, const property_list &propList = {}); + + /* Available only when: (dimensions > 0) */ + template + accessor(buffer &bufferRef, + handler &commandGroupHandlerRef, TagT tag, + const property_list &propList = {}); + + /* Available only when: (dimensions > 0) */ + template + accessor(buffer &bufferRef, + range accessRange, const property_list &propList = {}); + + /* Available only when: (dimensions > 0) */ + template + accessor(buffer &bufferRef, + range accessRange, TagT tag, + const property_list &propList = {}); + + /* Available only when: (dimensions > 0) */ + template + accessor(buffer &bufferRef, + range accessRange, id accessOffset, + const property_list &propList = {}); + + /* Available only when: (dimensions > 0) */ + template + accessor(buffer &bufferRef, + range accessRange, id accessOffset, + TagT tag, const property_list &propList = {}); + + /* Available only when: (dimensions > 0) */ + template + accessor(buffer &bufferRef, + handler &commandGroupHandlerRef, range accessRange, + const property_list &propList = {}); + + /* Available only when: (dimensions > 0) */ + template + accessor(buffer &bufferRef, + handler &commandGroupHandlerRef, range accessRange, + TagT tag, const property_list &propList = {}); + + /* Available only when: (dimensions > 0) */ + template + accessor(buffer &bufferRef, + handler &commandGroupHandlerRef, range accessRange, + id accessOffset, const property_list &propList = {}); + + /* Available only when: (dimensions > 0) */ + template + accessor(buffer &bufferRef, + handler &commandGroupHandlerRef, range accessRange, + id accessOffset, TagT tag, + const property_list &propList = {}); +``` + +Apply the same changes to the accessor constructors in Table 4.48: Constructors of the accessor class template buffer specialization. + +NOTE: Oddly enough, due to the rules in section 4.7.6.3 about deduction guides I don't think I need to explicitly list the new deduction guides here. +Readers may find that confusing given that deduction guides are explicitly listed for other classes, but that's how the spec is written. +The deduction guides must ensure that property_listT is inferred to be property_list. + +Also add to the listing a conversion function: + +```c++ +/* Available only when new_property_listT is convertible to property_listT */ +template +operator accessor () const; +``` + +And add a new row to Table 4.49: Member functions of the accessor class template buffer specialization for this new function: + +-- +[options="header"] +|==== +| Member function | Description +a| +```c++ +template +operator accessor () const; +``` | Available only when property_listT is convertible to new_property_listT. Converts this accessor to an accessor with a different property_list. +|==== +-- + +=== Section 4.7.6.9.2 Device buffer accessor properties + +Add two new properties to the listing: + +```c++ +namespace sycl { +namespace property { + struct noinit {}; + struct no_offset {}; + struct no_alias {}; +} // namespace property + +inline constexpr property::noinit noinit; +inline constexpr property::no_offset no_offset; +inline constexpr property::no_alias no_alias; +``` + +Rewrite Table 4.50: Properties supported by the SYCL accessor class as follows, introducing two new rows and a new column to indicate if a property is compile-time constant: + +-- +[options="header"] +|==== +| Property | Description | Compile-time Constant +| property::noinit | The noinit property notifies the SYCL runtime that previous contents of a buffer can be discarded. Replaces deprecated discard_write and discard_read_write access modes. | No +| property::no_offset | The no_offset property notifies the SYCL device compiler that the accessor will never contain an offset. This may enable the compiler to make assumptions about the alignment of the accessor that it couldn't make otherwise. | Yes +| property::no_alias | The no_alias property notifies the SYCL device compiler that the accessor will not alias with any other accessors or USM pointers accessed in the same kernel. This is an unchecked assertion by the programmer and results in undefined behaviour if it is violated. | Yes +|==== +-- + +Add two rows to Table 4.51: Constructors of the accessor property classes: + +-- +[options="header"] +|=== +| Constructor | Description +| property::no_offset::no_offset() | Constructs a SYCL no_offset property instance. +| property::no_alias::no_alias() | Constructs a SYCL no_alias property instance. +-- + +== Revision History + +[cols="5,15,15,70"] +[grid="rows"] +[options="header"] +|======================================== +|Rev|Date|Author|Changes +|A|2020-06-18|Joe Garvey|Initial public draft +|========================================