[SYCL] handle group size if sycl or not

Author: dklochkov-emb

clang/lib/Sema/SemaDeclAttr.cpp

@@ -2859,6 +2859,25 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
                  WorkGroupAttr(S.Context, AL, WGSize[0], WGSize[1], WGSize[2]));
 }
 
+// Handles reqd_sub_group_size and sub_group_size.
+static void handleSubGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
+  if (!AL.checkExactlyNumArgs(S, 1))
+    return;
+
+  uint32_t WGSize[1];
+  Expr *E = AL.getArgAsExpr(0);
+  if (!S.checkUInt32Argument(AL, E, WGSize[0], 0,
+                             /*StrictlyUnsigned=*/true))
+    return;
+  if (WGSize[0] == 0) {
+    S.Diag(AL.getLoc(), diag::err_attribute_argument_is_zero)
+        << AL << E->getSourceRange();
+    return;
+  }
+  D->addAttr(::new (S.Context)
+                 IntelReqdSubGroupSizeAttr(S.Context, AL, E));
+}
+
 static void handleVecTypeHint(Sema &S, Decl *D, const ParsedAttr &AL) {
   // Given attribute is deprecated without replacement in SYCL 2020 mode.
   // Ignore the attribute in SYCL 2020.
@@ -6801,7 +6820,11 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL,
     S.SYCL().handleSYCLIntelMaxWorkGroupsPerMultiprocessor(D, AL);
     break;
   case ParsedAttr::AT_IntelReqdSubGroupSize:
-    S.SYCL().handleIntelReqdSubGroupSizeAttr(D, AL);
+    // TODO: deprecate other spellings and change if to assert
+    if (S.getLangOpts().SYCLIsDevice || S.getLangOpts().SYCLIsHost)
+      S.SYCL().handleIntelReqdSubGroupSizeAttr(D, AL);
+    else
+      handleSubGroupSize(S, D, AL);
     break;
   case ParsedAttr::AT_IntelNamedSubGroupSize:
     S.SYCL().handleIntelNamedSubGroupSizeAttr(D, AL);

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

@@ -32,33 +32,33 @@ int main() {
   });
   return 0;
 }
-// expected-note@+2 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// 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(16)]] SYCL_EXTERNAL void B();
-// expected-note@+2 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// 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(16)]] void A() // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}} 
 { 
 }
 
-// expected-note@+2 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// 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(16)]] SYCL_EXTERNAL void B();
 
-// expected-note@+2 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// 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(16)]] SYCL_EXTERNAL void B() { // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
   A();
 }
-// expected-note@+3 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// expected-note@+3 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}}
 // expected-warning@+2{{attribute 'intel::reqd_sub_group_size' is deprecated}}
 // 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}}
 
 // expected-note@+5 {{conflicting attribute is here}}
 // expected-error@+4 {{conflicting attributes applied to a SYCL kernel}}
 // expected-warning@+3 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
-// expected-note@+2 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// 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(4)]] __attribute__((sycl_device)) void sg_size4() {
   sg_size2();
@@ -69,38 +69,38 @@ int main() {
 // 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}}
-// expected-note@+2 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// 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("foo")]] void func() {} // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'const char[4]'}}
-// expected-note@+2 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// 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(-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}}
 
 // Diagnostic is emitted because the arguments mismatch.
-// expected-note@+2 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// 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 quux();  // expected-note {{previous attribute is here}}
-// expected-note@+2 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// 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(100)]] void quux(); // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}} expected-note {{previous attribute is here}}
-// expected-note@+2 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// 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(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.
 [[intel::reqd_sub_group_size]] void quibble(); // expected-error {{'reqd_sub_group_size' attribute takes one argument}}
 
 // No diagnostic is emitted because the arguments match.
-// expected-note@+2 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// 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();
-// expected-note@+2 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// 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() {} // 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.
-// expected-note@+2 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// 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(); // OK
 
@@ -110,7 +110,7 @@ template <typename Ty>
 // expected-error@+5{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
 // expected-error@+4 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}}
 // expected-error@+3 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}}
-// expected-note@+2 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// 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(Ty{})]] void func() {}
 
@@ -129,21 +129,21 @@ void test() {
 int foo1();
 // expected-error@+4{{expression is not an integral constant expression}}
 // expected-note@+3{{non-constexpr function 'foo1' cannot be used in a constant expression}}
-// expected-note@+2 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// 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(foo1() + 12)]] void func1();
 
 // Test that checks expression is a constant expression.
 constexpr int bar1() { return 0; }
-// expected-note@+2 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// 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(bar1() + 12)]] void func2(); // OK
 
 // Test that checks template parameter support on member function of class template.
 template <int SIZE>
 class KernelFunctor {
 public:
-  // expected-note@+3 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+  // expected-note@+3 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}}
   // expected-warning@+2{{attribute 'intel::reqd_sub_group_size' is deprecated}}
   // expected-error@+1{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
   [[intel::reqd_sub_group_size(SIZE)]] void operator()() {}
@@ -157,20 +157,20 @@ int check() {
 
 // Test that checks template parameter support on function.
 template <int N>
-// expected-note@+4 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// expected-note@+4 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}}
 // expected-warning@+3{{attribute 'intel::reqd_sub_group_size' is deprecated}}
 // 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() {}
 
 template <int N>
-// expected-note@+3 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// expected-note@+3 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}}
 // expected-warning@+2{{attribute 'intel::reqd_sub_group_size' is deprecated}}
 // 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}}
 
 template <int N>
-// expected-note@+2{{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+// 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(N)]] void func4() {} // expected-warning {{attribute 'reqd_sub_group_size' is already applied with different arguments}}
 

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

@@ -10,20 +10,26 @@
 
 #include "Inputs/sycl.hpp"
 
+// expected-note@+4 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}}
+// expected-warning@+3 {{attribute 'intel::sub_group_size' is deprecated}}
 // expected-error@+2 {{'named_sub_group_size' and '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}}
 // expected-note@+1 {{conflicting attribute is here}}
 [[intel::named_sub_group_size(primary)]] [[intel::sub_group_size(1)]] void f2();
 
+// expected-note@+3 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}}
+// expected-warning@+2 {{attribute 'intel::sub_group_size' is deprecated}}
 // 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}}
 [[intel::named_sub_group_size(primary)]] void f3();
 
 // expected-note@+1 {{conflicting attribute is here}}
 [[intel::named_sub_group_size(primary)]] void f4();
+// expected-note@+3 {{did you mean to use 'sycl::reqd_sub_group_size' instead?}}
+// expected-warning@+2 {{attribute 'intel::sub_group_size' is deprecated}}
 // expected-error@+1 {{'sub_group_size' and 'named_sub_group_size' attributes are not compatible}}
 [[intel::sub_group_size(1)]] void f4();
 
@@ -115,8 +121,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 +165,7 @@ void calls_kernel_5() {
 
 // Don't diag with the old sub-group-size.
 void calls_kernel_6() {
-  // expected-note@+2 {{did you mean to use 'intel::reqd_sub_group_size' instead?}}
+  // 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}}
   sycl::kernel_single_task<class Kernel6>([]() [[intel::reqd_sub_group_size(10)]] { // #Kernel6
     NoAttrExternalNotDefined();