[SYCL][Doc] Function-type kernel attribute extension

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.
+

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(<type>)`
+|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.
+//************************************************************************