[SYCL] Deprecate [[intel::reqd_work_group_size]] attribute spelling

clang/include/clang/Basic/AttrDocs.td

@@ -2526,14 +2526,14 @@ of the kernel the attribute is applied to.
   // argument to the attribute is the one which increments fastest.
   struct func {
     [[intel::num_simd_work_items(4)]]
-    [[intel::reqd_work_group_size(7, 4, 64)]]
+    [[sycl::reqd_work_group_size(7, 4, 64)]]
     void operator()() const {}
   };
 
   // Note, '8' is evenly divisible by '8'; in SYCL, the last
   // argument to the attribute is the one which increments fastest.
   struct bar {
-    [[intel::reqd_work_group_size(1, 1, 8)]]
+    [[sycl::reqd_work_group_size(1, 1, 8)]]
     [[intel::num_simd_work_items(8)]]
     void operator()() const {}
   };
@@ -2659,6 +2659,10 @@ As an Intel extension, the ``[[intel::reqd_work_group_size(X, Y, Z)]]``
 spelling is supported with the same semantics as the
 ``[[sycl::reqd_work_group_size(X, Y, Z)]]`` spelling.
 
+The ``[[intel::reqd_work_group_size(X, Y, Z)]]`` attribute spelling
+is deprecated in favor of SYCL 2020 attribute spelling
+``[[sycl::reqd_work_group_size]]``.
+
 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.
@@ -2667,30 +2671,6 @@ In OpenCL C, this attribute is available with the GNU spelling
 
   __kernel __attribute__((reqd_work_group_size(8, 16, 32))) void test() {}
 
-  }];
-}
-
-def WorkGroupSizeHintAttrDocs : Documentation {
-  let Category = DocCatFunction;
-  let Content = [{
-Applies to a device function/lambda function. Hint to the compiler on the work-
-group size most likely to be used when launching the kernel at runtime. Each
-argument must be a nonzero integral constant expression whose value is
-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 GNU spelling is deprecated in SYCL mode.
-
-.. code-block:: c++
-
-  [[sycl::work_group_size_hint(4, 4, 4)]] void foo() {}
-
-  class Foo {
-  public:
-    [[sycl::work_group_size_hint(2, 2, 2)]] void operator()() const {}
-  };
-
 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, and in SYCL, the last argument is the index that
@@ -2714,14 +2694,14 @@ in the ``reqd_work_group_size`` attribute.
   // argument to the attribute is the one which increments fastest.
   struct func {
     [[intel::num_simd_work_items(4)]]
-    [[intel::reqd_work_group_size(7, 4, 64)]]
+    [[sycl::reqd_work_group_size(7, 4, 64)]]
     void operator()() const {}
   };
 
   // Note, '8' is evenly divisible by '8'; in SYCL, the last
   // argument to the attribute is the one which increments fastest.
   struct bar {
-    [[intel::reqd_work_group_size(1, 1, 8)]]
+    [[sycl::reqd_work_group_size(1, 1, 8)]]
     [[intel::num_simd_work_items(8)]]
     void operator()() const {}
   };
@@ -2765,6 +2745,30 @@ in the ``reqd_work_group_size`` attribute.
   }];
 }
 
+def WorkGroupSizeHintAttrDocs : Documentation {
+  let Category = DocCatFunction;
+  let Content = [{
+Applies to a device function/lambda function. Hint to the compiler on the work-
+group size most likely to be used when launching the kernel at runtime. Each
+argument must be a nonzero integral constant expression whose value is
+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 GNU spelling is deprecated in SYCL mode.
+
+.. code-block:: c++
+
+  [[sycl::work_group_size_hint(4, 4, 4)]] void foo() {}
+
+  class Foo {
+  public:
+    [[sycl::work_group_size_hint(2, 2, 2)]] void operator()() const {}
+  };
+
+  }];
+}
+
 def SYCLIntelMaxWorkGroupSizeAttrDocs : Documentation {
   let Category = DocCatFunction;
   let Heading = "intel::max_work_group_size";

clang/lib/Sema/SemaDeclAttr.cpp

@@ -337,6 +337,15 @@ void Sema::CheckDeprecatedSYCLAttributeSpelling(const ParsedAttr &A,
       return;
     }
 
+    // Deprecate [[intel::reqd_work_group_size]] attribute spelling in favor
+    // of SYCL 2020 attribute spelling [[sycl::reqd_work_group_size]].
+    if (A.hasScope() && A.getScopeName()->isStr("intel")) {
+      if (A.getKind() == ParsedAttr::AT_ReqdWorkGroupSize) {
+        DiagnoseDeprecatedAttribute(A, "sycl", NewName);
+        return;
+      }
+    }
+
     // All GNU-style spellings are deprecated in favor of a C++-style spelling.
     if (A.getSyntax() == ParsedAttr::AS_GNU) {
       // Note: we cannot suggest an automatic fix-it because GNU-style

clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp

@@ -12,17 +12,17 @@ class Functor {
 
 class Functor1 {
 public:
-  [[intel::reqd_sub_group_size(2), intel::reqd_work_group_size(64, 32, 32)]] void operator()() const {}
+  [[intel::reqd_sub_group_size(2), sycl::reqd_work_group_size(64, 32, 32)]] void operator()() const {}
 };
 
 template <int SIZE, int SIZE1, int SIZE2>
 class Functor2 {
 public:
-  [[intel::reqd_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() const {}
+  [[sycl::reqd_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() const {}
 };
 
 template <int N, int N1, int N2>
-[[intel::reqd_work_group_size(N, N1, N2)]] void func() {}
+[[sycl::reqd_work_group_size(N, N1, N2)]] void func() {}
 
 int main() {
   q.submit([&](handler &h) {

clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp

@@ -67,7 +67,7 @@ struct TRIFuncObjGood2 {
 };
 
 struct TRIFuncObjGood3 {
-  [[intel::reqd_work_group_size(1)]]
+  [[sycl::reqd_work_group_size(1)]]
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
@@ -85,7 +85,7 @@ struct TRIFuncObjGood5 {
 };
 
 struct TRIFuncObjGood6 {
-  [[intel::reqd_work_group_size(4, 1, 1)]]
+  [[sycl::reqd_work_group_size(4, 1, 1)]]
   [[intel::max_global_work_dim(3)]] void
   operator()() const {}
 };
@@ -119,7 +119,7 @@ void TRIFuncObjBad::operator()() const {}
 // attributes when merging, so the test compiles without
 // any diagnostic when it shouldn't.
 struct TRIFuncObjBad1 {
-  [[intel::reqd_work_group_size(4, 4, 4)]] void
+  [[sycl::reqd_work_group_size(4, 4, 4)]] void
   operator()() const;
 };
 
@@ -166,7 +166,9 @@ struct TRIFuncObjBad5 {
 };
 
 struct TRIFuncObjBad6 {
-  [[intel::reqd_work_group_size(4)]]   // expected-error{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}}
+  [[intel::reqd_work_group_size(4)]]   // expected-error{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}} \
+                                       // expected-warning {{attribute 'intel::reqd_work_group_size' is deprecated}} \
+                                       // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}}
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
@@ -184,7 +186,7 @@ struct TRIFuncObjBad8 {
   operator()() const;
 };
 
-[[intel::reqd_work_group_size(4, 4, 4)]] // expected-error{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}}
+[[sycl::reqd_work_group_size(4, 4, 4)]] // expected-error{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}}
 void TRIFuncObjBad8::operator()() const {}
 
 struct TRIFuncObjBad9 {
@@ -201,7 +203,7 @@ void TRIFuncObjBad9::operator()() const {}
 struct TRIFuncObjBad10 {
   // expected-error@+2{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}}
   // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}}
-  [[intel::reqd_work_group_size(-4, 1)]]
+  [[sycl::reqd_work_group_size(-4, 1)]]
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
@@ -219,7 +221,7 @@ struct TRIFuncObjBad12 {
 };
 
 struct TRIFuncObjBad13 {
-  [[intel::reqd_work_group_size(4)]]
+  [[sycl::reqd_work_group_size(4)]]
   [[intel::max_global_work_dim(-2)]] // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}}
   void operator()() const {}
 };

clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp

@@ -10,7 +10,7 @@ queue q;
 // expected-no-diagnostics
 class Functor {
 public:
-  [[intel::reqd_work_group_size(4)]] void operator()() const {}
+  [[sycl::reqd_work_group_size(4)]] void operator()() const {}
 };
 
 void bar() {
@@ -21,15 +21,14 @@ void bar() {
 }
 
 #else
-[[intel::reqd_work_group_size(4)]] void f4x1x1() {} // expected-note {{conflicting attribute is here}}
+[[sycl::reqd_work_group_size(4)]] void f4x1x1() {} // expected-note {{conflicting attribute is here}}
 // expected-note@-1 {{conflicting attribute is here}}
-[[intel::reqd_work_group_size(32)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}}
+[[sycl::reqd_work_group_size(32)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}}
+[[sycl::reqd_work_group_size(16)]] void f16x1x1() {}      // expected-note {{conflicting attribute is here}}
+[[sycl::reqd_work_group_size(16, 16)]] void f16x16x1() {} // expected-note {{conflicting attribute is here}}
 
-[[intel::reqd_work_group_size(16)]] void f16x1x1() {}      // expected-note {{conflicting attribute is here}}
-[[intel::reqd_work_group_size(16, 16)]] void f16x16x1() {} // expected-note {{conflicting attribute is here}}
-
-[[intel::reqd_work_group_size(32, 32)]] void f32x32x1() {}      // expected-note {{conflicting attribute is here}}
-[[intel::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}}
+[[sycl::reqd_work_group_size(32, 32)]] void f32x32x1() {}      // expected-note {{conflicting attribute is here}}
+[[sycl::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}}
 
 #ifdef TRIGGER_ERROR
 class Functor32 {
@@ -43,33 +42,34 @@ class Functor32 {
 class Functor33 {
 public:
   // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}}
-  [[intel::reqd_work_group_size(32, -4)]] void operator()() const {}
+  [[sycl::reqd_work_group_size(32, -4)]] void operator()() const {}
 };
 
 class Functor30 {
 public:
   // expected-warning@+1 2{{implicit conversion changes signedness: 'int' to 'unsigned long long'}}
-  [[intel::reqd_work_group_size(30, -30, -30)]] void operator()() const {}
+  [[sycl::reqd_work_group_size(30, -30, -30)]] void operator()() const {}
 };
 
 class Functor16 {
 public:
-  [[intel::reqd_work_group_size(16)]] void operator()() const {}
+  [[sycl::reqd_work_group_size(16)]] void operator()() const {}
 };
 
 class Functor64 {
 public:
-  [[intel::reqd_work_group_size(64, 64)]] void operator()() const {}
+  [[sycl::reqd_work_group_size(64, 64)]] void operator()() const {}
 };
 
 class Functor16x16x16 {
 public:
-  [[intel::reqd_work_group_size(16, 16, 16)]] void operator()() const {}
+  [[intel::reqd_work_group_size(16, 16, 16)]] void operator()() const {} // expected-warning {{attribute 'intel::reqd_work_group_size' is deprecated}} \
+                                                                         // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}}
 };
 
 class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}}
 public:
-  [[intel::reqd_work_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}}
+  [[sycl::reqd_work_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}}
     f4x1x1();
   }
 };
@@ -107,7 +107,7 @@ int main() {
     Functor30 f30;
     h.single_task<class kernel_name6>(f30);
 
-    h.single_task<class kernel_name7>([]() [[intel::reqd_work_group_size(32, 32, 32)]] {
+    h.single_task<class kernel_name7>([]() [[sycl::reqd_work_group_size(32, 32, 32)]] {
       f32x32x32();
     });
 #ifdef TRIGGER_ERROR
@@ -130,7 +130,7 @@ int main() {
     });
 
     // expected-error@+1 {{expected variable name or 'this' in lambda capture list}}
-    h.single_task<class kernel_name12>([[intel::reqd_work_group_size(32, 32, 32)]][]() {
+    h.single_task<class kernel_name12>([[sycl::reqd_work_group_size(32, 32, 32)]][]() {
       f32x32x32();
     });
 

clang/test/SemaSYCL/intel-reqd-work-group-size-host.cpp

@@ -1,13 +1,13 @@
 // RUN: %clang_cc1 -fsycl-is-host -Wno-sycl-2017-compat -fsyntax-only -verify %s
 // expected-no-diagnostics
 
-[[intel::reqd_work_group_size(4)]] void f4x1x1() {}
+[[sycl::reqd_work_group_size(4)]] void f4x1x1() {}
 
-[[intel::reqd_work_group_size(16)]] void f16x1x1() {}
+[[sycl::reqd_work_group_size(16)]] void f16x1x1() {}
 
-[[intel::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {}
+[[sycl::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {}
 
 class Functor64 {
 public:
-  [[intel::reqd_work_group_size(64, 64, 64)]] void operator()() const {}
+  [[sycl::reqd_work_group_size(64, 64, 64)]] void operator()() const {}
 };