diff --git a/sycl/doc/extensions/KernelRHSAttributes/README.md b/sycl/doc/extensions/KernelRHSAttributes/README.md new file mode 100644 index 0000000000000..f647a761f40d7 --- /dev/null +++ b/sycl/doc/extensions/KernelRHSAttributes/README.md @@ -0,0 +1,6 @@ +# SYCL_INTEL_attribute_style + +Extension that deprecates use of function attributes (left-sided) applied to +device functions for kernel attributes, and introduces instead function-type +attributes (right-sided) for kernel attributes that apply directly to kernel functions. + diff --git a/sycl/doc/extensions/KernelRHSAttributes/SYCL_INTEL_attribute_style.asciidoc b/sycl/doc/extensions/KernelRHSAttributes/SYCL_INTEL_attribute_style.asciidoc new file mode 100755 index 0000000000000..b9f7555a54332 --- /dev/null +++ b/sycl/doc/extensions/KernelRHSAttributes/SYCL_INTEL_attribute_style.asciidoc @@ -0,0 +1,219 @@ += SYCL_INTEL_attribute_style + +: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++ + +== Introduction +IMPORTANT: This specification is a draft. + +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. + +NOTE: This document is better viewed when rendered as html with asciidoctor. +GitHub does not render image icons. + +This document describes an extension that deprecates use of function attributes +(left-sided) for kernel attributes, and introduces use of function-type +attributes (right-sided) for kernel attributes. This allows SYCL kernel +attributes to be applied directly to kernels defined as lambdas, and no longer requires +propagation of attributes across call trees (which left-sided function attributes require). + +== Name Strings + ++SYCL_INTEL_attribute_style+ + +== Notice + +Copyright (c) 2020 Intel Corporation. All rights reserved. + +== 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: 1 + +== Contact +Michael Kinsner, Intel (michael 'dot' kinsner 'at' intel 'dot' com) + +== Dependencies + +This extension is written against the SYCL 1.2.1 specification, Revision 6. + +== Overview + +SYCL 1.2.1 defines kernel attributes as applying to device functions (functions called +by kernels), and describes a call tree-based propagation scheme in which the attributes would +propagate to calling kernels. This extension instead enables attributes to be applied +directly to kernel functions, avoiding complex and error prone call tree propagation, and +making it clear to which kernel an attribute applies. + +A kernel attribute applied to the function as required by SYCL 1.2.1 looks like: + +[source,c++] +---- +[[attrib]] void foo1() {}; + +class f { + [[attrib]] void foo2() {}; +}; +---- + +where `attrib` is a placeholder for any of the kernel attributes defined by the SYCL specification or extensions. + +This extension deprecates the SYCL 1.2.1 attribute style (attribute applied to +a device function) and instead defines kernel attributes as attributes that apply to the +function type. The location of the resulting attributes looks like: + +[source,c++] +---- +void bar1() [[attrib]] {}; + +class f { + void bar2() [[attrib]] {}; +}; + +class KernelFunctor { + public: + void operator()(sycl::item<1> item) [[attrib]] {}; +}; + +auto bar3 = []()[[attrib]]{}; // Works on lambdas. operator() type +---- + +The function type attributes have an effect when applied to a kernel function, +do not propagate up or down call trees unless specified by a specific attribute, +and the effect when applied to non-kernel functions or non-functions is implementation defined. + +== Modifications of SYCL 1.2.1 Specification + +=== Modify Section 6.7 (Attributes) + +==== Rename the section + +Rename Section 6.7 from "Attributes" to "Kernel attributes". + +==== Replace the entire section with: + +The SYCL programming interface defines attributes that augment the +information available while generating the device code for a particular platform. +The attributes in Table 1 are defined in the `sycl::` namespace +and are applied to the function-type of kernel function declarations using +{cpp}11 attribute specifier syntax. + +A given attribute-token shall appear at most once in each attribute-list. The +first declaration of a function shall specify an attribute if any declaration +of that function specifies the same attribute. If a function is declared with +an attribute in one translation unit and the same function is declared without +the same attribute in another translation unit, the program is ill-formed and +no diagnostic is required. + +If there are any conflicts between different kernel attributes, then the behavior +is undefined. The attributes have an effect when applied to a device function and no +effect otherwise (i.e. no effect on non-device functions and on anything other than a +function). If an attribute is applied to a device function that is not a kernel function +(but that is potentially called from a kernel function), then the effect is implementation defined. +It is implementation defined whether any diagnostic is produced when an attribute is applied +to anything other than the function-type of a kernel function declaration. + +.Attributes supported by the SYCL programming interface +[cols="2*"] +|=== +|`reqd_work_group_size(dim0)` +`reqd_work_group_size(dim0, dim1)` +`reqd_work_group_size(dim0, dim1, dim2)` +|Indicates that the kernel must be launched with the specified work-group size. The sizes +are written in row-major format. Each argument to the attribute must be an integral +constant expression. The dimensionality of the attribute variant used must match the +dimensionality of the work-group used to invoke the kernel. + +SYCL device compilers should give a compilation error if the required work-group size +is unsupported. If the kernel is submitted for execution using an incompatible +work-group size, the SYCL runtime must throw an `nd_range_error`. + +|`work_group_size_hint(dim0)` +`work_group_size_hint(dim0, dim1)` +`work_group_size_hint(dim0, dim1, dim2)` +|Hint to the compiler on the work-group size most likely to be used when launching the kernel +at runtime. Each argument must be an integral constant expression, and the number of dimensional +values defined provide additional information to the compiler on the dimensionality most likely +to be used when launching the kernel at runtime. The effect of this attribute, if any, is +implementation defined. + +|`vec_type_hint()` +|Hint to the compiler on the vector computational width of of the kernel. The argument must be +one of the vector types defined in section 4.10.2. This attribute is deprecated by this +extension (available for use, but will be removed in the future and is not recommended for +use in new code). +|=== + + + +==== Add new sub-section 6.7.1: Deprecated attribute syntax +The SYCL 1.2.1 specification defined two mechanisms for kernel attributes to be specified, +which are deprecated by this extension. Deprecation means that the syntaxes are supported, +but will be removed in the future, and are therefore not recommended for use. Specifically, +the following two attribute syntaxes defined by the SYCL 1.2.1 specification are deprecated: + +1. The `attribute` syntax defined by the OpenCL C specification within device +code (`__attribute__\((attrib))`). +2. {cpp}11 attribute specifier syntax (`\[[attrib]]`) applied to device functions +(not the function-type), including automatic propagation of the attribute to any +caller of those device functions. + + +== Issues + +None. + +//. asd +//+ +//-- +//*RESOLUTION*: Not resolved. +//-- + +== Revision History + +[cols="5,15,15,70"] +[grid="rows"] +[options="header"] +|======================================== +|Rev|Date|Author|Changes +|1|2020-04-08|Michael Kinsner|*Initial public 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. +//************************************************************************