diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 2e7b1bc966255..92a878522d25e 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2980,18 +2980,32 @@ argument to **clEnqueueNDRangeKernel** (in OpenCL) or to **parallel_for** in SYCL. This allows the compiler to optimize the generated code appropriately for the kernel to which attribute is applied. -While semantic of this attribute is the same between OpenCL and SYCL, -spelling is a bit different: - -SYCL 2020 describes the ``[[sycl::reqd_work_group_size(X, Y, Z)]]`` spelling -in detail. This attribute indicates that the kernel must be launched with the -specified work-group size. The order of the arguments matches the constructor -of the group class. 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. This spelling -allows the Y and Z arguments to be optional. If not provided by the user, the -value of Y and Z defaults to 1. See section 5.8.1 Kernel Attributes for more -details. +The arguments to ``reqd_work_group_size`` are ordered based on which index +increments the fastest. In OpenCL, the first argument is the index that +increments the fastest. In SYCL, the last argument is the index that increments +the fastest. + +In OpenCL C, this attribute is available with the GNU spelling +(``__attribute__((reqd_work_group_size(X, Y, Z)))``) and all three arguments +are required. See section 6.7.2 Optional Attribute Qualifiers of OpenCL 1.2 +specification for details. + +.. code-block:: c++ + + __kernel __attribute__((reqd_work_group_size(8, 16, 32))) void test() {} + +In SYCL, the attribute accepts either one, two, or three arguments; in each +form, the last (or only) argument is the index that increments fastest. The +number of arguments passed to the attribute must match the dimensionality of +the kernel the attribute is applied to. + +SYCL 2020 describes the ``[[sycl::reqd_work_group_size(dim0, dim1, dim2)]]`` +spelling in detail. This attribute indicates that the kernel must be launched +with the specified work-group size. The order of the arguments matches the +constructor of the ``range`` class. 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. See +section 5.8.1 Kernel Attributes for more details. In SYCL 1.2.1 mode, the ``cl::reqd_work_group_size`` and ``sycl::reqd_work_group_size`` attributes are propagated from the function they @@ -3016,18 +3030,9 @@ attributes are not propagated to the kernel. template [[sycl::reqd_work_group_size(N, N1, N2)]] void func() {} -The ``[[cl::reqd_work_group_size(X, Y, Z)]]`` and -``__attribute__((reqd_work_group_size(X, Y, Z)))`` spellings are both +The ``[[cl::reqd_work_group_size(dim0, dim1, dim2)]]`` and +``__attribute__((reqd_work_group_size(dim0, dim1, dim2)))`` spellings are both deprecated in SYCL 2020. - -In OpenCL C, this attribute is available with the GNU spelling -(``__attribute__((reqd_work_group_size(X, Y, Z)))``), see section -6.7.2 Optional Attribute Qualifiers of OpenCL 1.2 specification for details. - -.. code-block:: c++ - - __kernel __attribute__((reqd_work_group_size(8, 16, 32))) void test() {} - }]; } @@ -3041,6 +3046,15 @@ unsigned. 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 arguments to ``work_group_size_hint`` are ordered based on which index +increments the fastest. In OpenCL, the first argument is the index that +increments the fastest. In SYCL, the last argument is the index that increments +the fastest. + +In OpenCL C, this attribute is available with the GNU spelling +(``__attribute__((work_group_size_hint(X, Y, Z)))``) and all three arguments +are required. + The GNU spelling is deprecated in SYCL mode. .. code-block:: c++ @@ -3052,15 +3066,6 @@ The GNU spelling is deprecated in SYCL mode. [[sycl::work_group_size_hint(2, 2, 2)]] void operator()() const {} }; -The arguments to ``work_group_size_hint`` are ordered based on which index -increments the fastest. In OpenCL, the first argument is the index that -increments the fastest, and in SYCL, the last argument is the index that -increments the fastest. - -In OpenCL C, this attribute is available with the GNU spelling -(``__attribute__((work_group_size_hint(X, Y, Z)))``) and all -three arguments are required. - In SYCL, the attribute accepts either one, two, or three arguments; in each form, the last (or only) argument is the index that increments fastest. The number of arguments passed to the attribute must match the dimensionality of @@ -3077,9 +3082,11 @@ def SYCLIntelMaxWorkGroupSizeAttrDocs : Documentation { let Heading = "intel::max_work_group_size"; let Content = [{ Applies to a device function/lambda function. Indicates the maximum dimensions -of a work group. Values must be positive integers. This is similar to -reqd_work_group_size, but allows work groups that are smaller or equal to the -specified sizes. +of a work group. Values must be positive integers. This attribute behaves +similarly to ``reqd_work_group_size``, but allows work groups that are smaller +or equal to the specified sizes. The dimensionality behaves the same as with +the SYCL ``reqd_work_group_size`` attribute, but *all* dimensions must be +provided. In SYCL 1.2.1 mode, the ``intel::max_work_group_size`` attribute is propagated from the function it is applied to onto the kernel which calls the function.