[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 the 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.
@@ -2714,14 +2718,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 {}
   };

clang/lib/Sema/SemaDeclAttr.cpp

@@ -337,6 +337,14 @@ void Sema::CheckDeprecatedSYCLAttributeSpelling(const ParsedAttr &A,
       return;
     }
 
+    // Deprecate [[intel::reqd_work_group_size]] attribute spelling in favor
+    // of the SYCL 2020 attribute spelling [[sycl::reqd_work_group_size]].
+    if (A.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && A.hasScope() &&
+        A.getScopeName()->isStr("intel")) {
+      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 {}
 };