From b486a5b8ed1e58945974b85536b792bf4d536c98 Mon Sep 17 00:00:00 2001 From: Fraser Cormack Date: Thu, 4 Jul 2024 13:36:54 +0100 Subject: [PATCH] [SYCL] Clarify a few work-group attribute docs (#14417) Declaring upfront that `reqd_work_group_size` has the same semantics between OpenCL and SYCL is misleading because the dimensionality is different. The docs for `work_group_size_hint` were better in clarifying the difference, so those have been copied and adapted. Describe OpenCL upfront as it's simpler, and won't be buried by the wordier SYCL documentation. Try and re-enforce the SYCL dimensionality rules to aid people more familiar with OpenCL and other similar languages. Clarify that the `max_work_group_size` attribute (at least currently) behaves like `reqd_work_group_size` but we require *all* dimensions to be supplied. --- clang/include/clang/Basic/AttrDocs.td | 77 +++++++++++++++------------ 1 file changed, 42 insertions(+), 35 deletions(-) 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.