Deprecate intel reqd sub group size

clang/lib/Sema/SemaSYCL.cpp

@@ -3167,7 +3167,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
 //        // code
 //      }
 //
-//      [[intel::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const
+//      [[sycl::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const
 //      {
 //        // code
 //      }

clang/lib/Sema/SemaSYCLDeclAttr.cpp

@@ -132,6 +132,13 @@ void SemaSYCL::checkDeprecatedSYCLAttributeSpelling(const ParsedAttr &A,
     return;
   }
 
+  // Additionally, diagnose deprecated [[intel::reqd_sub_group_size]] spelling
+  if (A.getKind() == ParsedAttr::AT_IntelReqdSubGroupSize && A.getScopeName() &&
+      A.getScopeName()->isStr("intel")) {
+    diagnoseDeprecatedAttribute(A, "sycl", "reqd_sub_group_size");
+    return;
+  }
+
   // Diagnose SYCL 2020 spellings in later SYCL modes.
   if (getLangOpts().getSYCLVersion() >= LangOptions::SYCL_2020) {
     // All attributes in the cl vendor namespace are deprecated in favor of a

clang/test/CodeGenSYCL/kernel-op-calls.cpp

@@ -11,7 +11,7 @@ class Functor1 {
 public:
  Functor1(){}
 
- [[intel::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const {}
+ [[sycl::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const {}
 
  [[sycl::work_group_size_hint(1, 2, 3)]] void operator()(sycl::id<2> id) const {}
 

clang/test/CodeGenSYCL/reqd-sub-group-size.cpp

@@ -7,25 +7,25 @@ queue q;
 
 class Functor16 {
 public:
-  [[intel::reqd_sub_group_size(16)]] void operator()() const {}
+  [[sycl::reqd_sub_group_size(16)]] void operator()() const {}
 };
 
 template <int SIZE>
 class Functor2 {
 public:
-  [[intel::reqd_sub_group_size(SIZE)]] void operator()() const {}
+  [[sycl::reqd_sub_group_size(SIZE)]] void operator()() const {}
 };
 
 template <int N>
-[[intel::reqd_sub_group_size(N)]] void func() {}
+[[sycl::reqd_sub_group_size(N)]] void func() {}
 
 int main() {
   q.submit([&](handler &h) {
     Functor16 f16;
     h.single_task<class kernel_name1>(f16);
 
     h.single_task<class kernel_name3>(
-        []() [[intel::reqd_sub_group_size(4)]]{});
+        []() [[sycl::reqd_sub_group_size(4)]]{});
 
     Functor2<2> f2;
     h.single_task<class kernel_name4>(f2);

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

@@ -7,12 +7,12 @@ queue q;
 
 class Functor {
 public:
-  [[intel::reqd_sub_group_size(4), cl::reqd_work_group_size(32, 16, 16)]] void operator()() const {}
+  [[sycl::reqd_sub_group_size(4), cl::reqd_work_group_size(32, 16, 16)]] void operator()() const {}
 };
 
 class Functor1 {
 public:
-  [[intel::reqd_sub_group_size(2), sycl::reqd_work_group_size(64, 32, 32)]] void operator()() const {}
+  [[sycl::reqd_sub_group_size(2), sycl::reqd_work_group_size(64, 32, 32)]] void operator()() const {}
 };
 
 template <int SIZE, int SIZE1, int SIZE2>

clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp

@@ -6,7 +6,7 @@ template <typename T> class Fobj {
 public:
   Fobj() {}
   void operator()() const {
-    auto L0 = []() [[intel::reqd_sub_group_size(4)]]{};
+    auto L0 = []() [[sycl::reqd_sub_group_size(4)]]{};
     L0();
   }
 };

clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp

@@ -9,14 +9,14 @@ queue q;
 
 class Functor16 {
 public:
-  [[intel::reqd_sub_group_size(16)]] void operator()() const {}
+  [[sycl::reqd_sub_group_size(16)]] void operator()() const {}
 };
 
 // Test that checks template parameter support on member function of class template.
 template <int SIZE>
 class KernelFunctor {
 public:
-  [[intel::reqd_sub_group_size(SIZE)]] void operator()() const {}
+  [[sycl::reqd_sub_group_size(SIZE)]] void operator()() const {}
 };
 
 // Test that checks template parameter support on function.
@@ -35,7 +35,7 @@ class KernelFunctor {
 // CHECK-NEXT: NonTypeTemplateParmDecl
 // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 12
 template <int N>
-[[intel::reqd_sub_group_size(N)]] void func() {}
+[[sycl::reqd_sub_group_size(N)]] void func() {}
 
 int main() {
   q.submit([&](handler &h) {
@@ -52,14 +52,14 @@ int main() {
     // CHECK-NEXT: ConstantExpr {{.*}} 'int'
     // CHECK-NEXT: value: Int 2
     // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}
-    h.single_task<class kernel_name3>([]() [[intel::reqd_sub_group_size(2)]] {});
+    h.single_task<class kernel_name3>([]() [[sycl::reqd_sub_group_size(2)]] {});
 
     // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5
     // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
     // CHECK-NEXT: ConstantExpr {{.*}} 'int'
     // CHECK-NEXT: value: Int 6
     // CHECK-NEXT: IntegerLiteral{{.*}}6{{$}}
-    h.single_task<class kernel_name5>([]() [[intel::reqd_sub_group_size(6)]] {});
+    h.single_task<class kernel_name5>([]() [[sycl::reqd_sub_group_size(6)]] {});
 
     // CHECK: FunctionDecl {{.*}}kernel_name_6
     // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
@@ -79,8 +79,8 @@ int main() {
         // CHECK-NEXT: value: Int 8
         // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}
         // CHECK-NOT: IntelReqdSubGroupSizeAttr
-        []() [[intel::reqd_sub_group_size(8),
-               intel::reqd_sub_group_size(8)]] {});
+        []() [[sycl::reqd_sub_group_size(8),
+               sycl::reqd_sub_group_size(8)]] {});
   });
   func<12>();
   return 0;

clang/test/SemaSYCL/reqd-sub-group-size-host.cpp

@@ -1,9 +1,9 @@
 // RUN: %clang_cc1 -fsycl-is-host -fsyntax-only -verify %s
 // expected-no-diagnostics
 
-[[intel::reqd_sub_group_size(8)]] void fun() {}
+[[sycl::reqd_sub_group_size(8)]] void fun() {}
 
 class Functor {
 public:
-  [[intel::reqd_sub_group_size(16)]] void operator()() {}
+  [[sycl::reqd_sub_group_size(16)]] void operator()() {}
 };

clang/test/SemaSYCL/reqd-sub-group-size.cpp

@@ -1,6 +1,6 @@
 // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -verify -pedantic %s
 
-// The test checks functionality of [[intel::reqd_sub_group_size()]] attribute on SYCL kernel.
+// The test checks functionality of [[sycl::reqd_sub_group_size()]] attribute on SYCL kernel and [[intel::reqd_sub_group_size()]] is deprecated.
 
 
 #include "sycl.hpp" //clang/test/SemaSYCL/Inputs/sycl.hpp
@@ -32,44 +32,51 @@ int main() {
   });
   return 0;
 }
-[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B();
-[[intel::reqd_sub_group_size(16)]] void A() // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}} 
+
+[[sycl::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B();
+[[sycl::reqd_sub_group_size(16)]] void A() // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}} 
 { 
 }
 
-[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B() { // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
+[[sycl::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B();
+[[sycl::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B() { // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
   A();
 }
+
 // expected-note@+1 {{conflicting attribute is here}}
-[[intel::reqd_sub_group_size(2)]] void sg_size2() {} // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
+[[sycl::reqd_sub_group_size(2)]] void sg_size2() {} // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
 
 // expected-note@+3 {{conflicting attribute is here}}
 // expected-error@+2 {{conflicting attributes applied to a SYCL kernel}}
 // expected-warning@+1 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
-[[intel::reqd_sub_group_size(4)]] __attribute__((sycl_device)) void sg_size4() {
+[[sycl::reqd_sub_group_size(4)]] __attribute__((sycl_device)) void sg_size4() {
   sg_size2();
 }
 
 // Test that checks support and functionality of reqd_sub_group_size attribute support on function.
 
 // Tests for incorrect argument values for Intel reqd_sub_group_size attribute.
-[[intel::reqd_sub_group_size]] void one() {}         // expected-error {{'reqd_sub_group_size' attribute takes one argument}}
-[[intel::reqd_sub_group_size(5)]] int a;             // expected-error{{'reqd_sub_group_size' attribute only applies to functions}}
-[[intel::reqd_sub_group_size("foo")]] void func() {} // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'const char[4]'}}
-[[intel::reqd_sub_group_size(-1)]] void func1() {}   // expected-error{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
-[[intel::reqd_sub_group_size(0, 1)]] void arg() {}   // expected-error{{'reqd_sub_group_size' attribute takes one argument}}
+[[sycl::reqd_sub_group_size]] void one() {}         // expected-error {{'reqd_sub_group_size' attribute takes one argument}}
+[[sycl::reqd_sub_group_size(5)]] int a;             // expected-error{{'reqd_sub_group_size' attribute only applies to functions}}
+[[sycl::reqd_sub_group_size("foo")]] void func() {} // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'const char[4]'}}
+[[sycl::reqd_sub_group_size(-1)]] void func1() {}   // expected-error{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
+[[sycl::reqd_sub_group_size(0, 1)]] void arg() {}   // expected-error{{'reqd_sub_group_size' attribute takes one argument}}
 
 // Diagnostic is emitted because the arguments mismatch.
-[[intel::reqd_sub_group_size(12)]] void quux();  // expected-note {{previous attribute is here}}
-[[intel::reqd_sub_group_size(100)]] void quux(); // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}} expected-note {{previous attribute is here}}
+[[sycl::reqd_sub_group_size(12)]] void quux();  // expected-note {{previous attribute is here}}
+[[sycl::reqd_sub_group_size(100)]] void quux(); // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}} expected-note {{previous attribute is here}}
 [[sycl::reqd_sub_group_size(200)]] void quux();  // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}}
 
 // Make sure there's at least one argument passed.
 [[sycl::reqd_sub_group_size]] void quibble(); // expected-error {{'reqd_sub_group_size' attribute takes one argument}}
 
 // No diagnostic is emitted because the arguments match.
+[[sycl::reqd_sub_group_size(12)]] void same();
+[[sycl::reqd_sub_group_size(12)]] void same() {} // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
+
+// expected-note@+2 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}}
+// expected-warning@+1{{attribute 'intel::reqd_sub_group_size' is deprecated}}
 [[intel::reqd_sub_group_size(12)]] void same();
-[[intel::reqd_sub_group_size(12)]] void same() {} // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
 
 // No diagnostic because the attributes are synonyms with identical behavior.
 [[sycl::reqd_sub_group_size(12)]] void same(); // OK
@@ -80,7 +87,7 @@ template <typename Ty>
 // expected-error@+3{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
 // expected-error@+2 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}}
 // expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}}
-[[intel::reqd_sub_group_size(Ty{})]] void func() {}
+[[sycl::reqd_sub_group_size(Ty{})]] void func() {}
 
 struct S {};
 void test() {
@@ -97,18 +104,18 @@ void test() {
 int foo1();
 // expected-error@+2{{expression is not an integral constant expression}}
 // expected-note@+1{{non-constexpr function 'foo1' cannot be used in a constant expression}}
-[[intel::reqd_sub_group_size(foo1() + 12)]] void func1();
+[[sycl::reqd_sub_group_size(foo1() + 12)]] void func1();
 
 // Test that checks expression is a constant expression.
 constexpr int bar1() { return 0; }
-[[intel::reqd_sub_group_size(bar1() + 12)]] void func2(); // OK
+[[sycl::reqd_sub_group_size(bar1() + 12)]] void func2(); // OK
 
 // Test that checks template parameter support on member function of class template.
 template <int SIZE>
 class KernelFunctor {
 public:
   // expected-error@+1{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
-  [[intel::reqd_sub_group_size(SIZE)]] void operator()() {}
+  [[sycl::reqd_sub_group_size(SIZE)]] void operator()() {}
 };
 
 int check() {
@@ -121,14 +128,14 @@ int check() {
 template <int N>
 // expected-error@+2{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
 // expected-warning@+1 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
-[[intel::reqd_sub_group_size(N)]] void func3() {}
+[[sycl::reqd_sub_group_size(N)]] void func3() {}
 
 template <int N>
 // expected-warning@+1 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
-[[intel::reqd_sub_group_size(4)]] void func4(); // expected-note {{previous attribute is here}}
+[[sycl::reqd_sub_group_size(4)]] void func4(); // expected-note {{previous attribute is here}}
 
 template <int N>
-[[intel::reqd_sub_group_size(N)]] void func4() {} // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}}
+[[sycl::reqd_sub_group_size(N)]] void func4() {} // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}}
 
 int check1() {
   // no error expected

clang/test/SemaSYCL/sub-group-size.cpp

@@ -3,29 +3,29 @@
 // RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-default-sub-group-size=10 -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -verify=expected,primary %s
 
 // Validate the semantic analysis checks for the interaction betwen the
-// named_sub_group_size and sub_group_size attributes. These are not able to be
+// named_sub_group_size and reqd_sub_group_size attributes. These are not able to be
 // combined, and require that they only be applied to non-sycl-kernel/
 // non-sycl-device functions if they match the kernel they are being called
 // from.
 
 #include "Inputs/sycl.hpp"
 
-// expected-error@+2 {{'named_sub_group_size' and 'sub_group_size' attributes are not compatible}}
+// expected-error@+2 {{'named_sub_group_size' and 'reqd_sub_group_size' attributes are not compatible}}
 // expected-note@+1 {{conflicting attribute is here}}
-[[intel::sub_group_size(1)]] [[intel::named_sub_group_size(automatic)]] void f1();
-// expected-error@+2 {{'sub_group_size' and 'named_sub_group_size' attributes are not compatible}}
+[[sycl::reqd_sub_group_size(1)]] [[intel::named_sub_group_size(automatic)]] void f1();
+// expected-error@+2 {{'reqd_sub_group_size' and 'named_sub_group_size' attributes are not compatible}}
 // expected-note@+1 {{conflicting attribute is here}}
-[[intel::named_sub_group_size(primary)]] [[intel::sub_group_size(1)]] void f2();
+[[intel::named_sub_group_size(primary)]] [[sycl::reqd_sub_group_size(1)]] void f2();
 
 // expected-note@+1 {{conflicting attribute is here}}
-[[intel::sub_group_size(1)]] void f3();
-// expected-error@+1 {{'named_sub_group_size' and 'sub_group_size' attributes are not compatible}}
+[[sycl::reqd_sub_group_size(1)]] void f3();
+// expected-error@+1 {{'named_sub_group_size' and 'reqd_sub_group_size' attributes are not compatible}}
 [[intel::named_sub_group_size(primary)]] void f3();
 
 // expected-note@+1 {{conflicting attribute is here}}
 [[intel::named_sub_group_size(primary)]] void f4();
-// expected-error@+1 {{'sub_group_size' and 'named_sub_group_size' attributes are not compatible}}
-[[intel::sub_group_size(1)]] void f4();
+// expected-error@+1 {{'reqd_sub_group_size' and 'named_sub_group_size' attributes are not compatible}}
+[[sycl::reqd_sub_group_size(1)]] void f4();
 
 // expected-note@+1 {{previous attribute is here}}
 [[intel::named_sub_group_size(automatic)]] void f5();
@@ -115,8 +115,14 @@ void calls_kernel_3() {
   });
 }
 
+// expected-note@+2 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}}
+// expected-warning@+1{{attribute 'intel::sub_group_size' is deprecated}}
 [[intel::sub_group_size(10)]] void AttrFunc2() {}                           // #AttrFunc2
+// expected-note@+2 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}}
+// expected-warning@+1{{attribute 'intel::sub_group_size' is deprecated}}
 [[intel::sub_group_size(10)]] SYCL_EXTERNAL void AttrExternalDefined2() {}  // #AttrExternalDefined2
+// expected-note@+2 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}}
+// expected-warning@+1{{attribute 'intel::sub_group_size' is deprecated}}
 [[intel::sub_group_size(10)]] SYCL_EXTERNAL void AttrExternalNotDefined2(); // #AttrExternalNotDefined2
 
 void calls_kernel_4() {
@@ -153,7 +159,7 @@ void calls_kernel_5() {
 
 // Don't diag with the old sub-group-size.
 void calls_kernel_6() {
-  sycl::kernel_single_task<class Kernel6>([]() [[intel::reqd_sub_group_size(10)]] { // #Kernel6
+  sycl::kernel_single_task<class Kernel6>([]() [[sycl::reqd_sub_group_size(10)]] { // #Kernel6
     NoAttrExternalNotDefined();
   });
 }

clang/test/SemaSYCL/sycl-attr-warn-non-kernel.cpp

@@ -7,7 +7,7 @@
 [[sycl::reqd_work_group_size(16)]] void f1(){ // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
 }
 
-[[intel::reqd_sub_group_size(12)]] void f3(){ // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
+[[sycl::reqd_sub_group_size(12)]] void f3(){ // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
 }
 
 [[sycl::reqd_work_group_size(16)]] void f4(){ // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
@@ -41,7 +41,7 @@ class Functor16x16x16 {
 
 class FunctorSubGroupSize4 {
 public:
-  [[intel::reqd_sub_group_size(4)]] void operator()() const{}  // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
+  [[sycl::reqd_sub_group_size(4)]] void operator()() const{}  // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
 };
 
 class Functor8 {

sycl/test-e2e/AOT/reqd-sg-size.cpp

@@ -43,7 +43,7 @@ template <size_t... Ns> struct SubgroupDispatcher {
           accessor acc{buf, cgh};
           cgh.parallel_for<kernel_name<size>>(
               nd_range<1>(1, 1),
-              [=](auto item) [[intel::reqd_sub_group_size(size)]] {
+              [=](auto item) [[sycl::reqd_sub_group_size(size)]] {
                 acc[0] = item.get_sub_group().get_max_local_range()[0];
               });
         });

sycl/test-e2e/ESIMD/named_barriers/allocate_barrier_InvokeSimd.cpp

@@ -32,7 +32,7 @@
 #ifdef IMPL_SUBGROUP
 #define SUBGROUP_ATTR
 #else
-#define SUBGROUP_ATTR [[intel::reqd_sub_group_size(VL)]]
+#define SUBGROUP_ATTR [[sycl::reqd_sub_group_size(VL)]]
 #endif
 
 using namespace sycl;

sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/array_input_sort.cpp

@@ -78,7 +78,7 @@ void RunSortOverGroupArray(sycl::queue &Q, const std::vector<T> &DataToSort,
                                                        CGH);
 
        CGH.parallel_for(
-           NDRange, [=](sycl::nd_item<Dims> id) [[intel::reqd_sub_group_size(
+           NDRange, [=](sycl::nd_item<Dims> id) [[sycl::reqd_sub_group_size(
                         ReqSubGroupSize)]] {
              const size_t GlobalLinearID = id.get_global_linear_id();
              using RadixSorterT = oneapi_exp::radix_sorters::group_sorter<

sycl/test-e2e/GroupAlgorithm/SYCL2020/group_sort/group_and_joint_sort.cpp

@@ -115,7 +115,7 @@ void RunJointSort(sycl::queue &Q, const std::vector<T> &DataToSort,
 
        CGH.parallel_for<KernelNameJoint<IntWrapper<Dims>,
                                         UseGroupWrapper<UseGroup>, T, Compare>>(
-           NDRange, [=](sycl::nd_item<Dims> ID) [[intel::reqd_sub_group_size(
+           NDRange, [=](sycl::nd_item<Dims> ID) [[sycl::reqd_sub_group_size(
                         ReqSubGroupSize)]] {
              auto Group = [&]() {
                if constexpr (UseGroup == UseGroupT::SubGroup)
@@ -282,7 +282,7 @@ void RunSortOVerGroup(sycl::queue &Q, const std::vector<T> &DataToSort,
 
        CGH.parallel_for<KernelNameOverGroup<
            IntWrapper<Dims>, UseGroupWrapper<UseGroup>, T, Compare>>(
-           NDRange, [=](sycl::nd_item<Dims> id) [[intel::reqd_sub_group_size(
+           NDRange, [=](sycl::nd_item<Dims> id) [[sycl::reqd_sub_group_size(
                         ReqSubGroupSize)]] {
              const size_t GlobalLinearID = id.get_global_linear_id();