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

Author: frasercrmck

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 <int N, int N1, int N2>
   [[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.