[FPGA][SYCL] Fix max_work_group_size and reqd_work_group_size attribute arguments check

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -11597,7 +11597,7 @@ def err_sycl_non_constant_array_type : Error<
 def err_conflicting_sycl_kernel_attributes : Error<
   "conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function">;
 def err_conflicting_sycl_function_attributes : Error<
-   "%0 attribute conflicts with '%1' attribute">;
+   "%0 attribute conflicts with %1 attribute">;
 def err_sycl_function_attribute_mismatch : Error<
   "SYCL kernel without %0 attribute can't call a function with this attribute">;
 def err_sycl_x_y_z_arguments_must_be_one : Error<

clang/lib/Sema/SemaDeclAttr.cpp

@@ -3268,12 +3268,31 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) {
 
   ASTContext &Ctx = S.getASTContext();
 
+  // 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
+  // increments the fastest.
+  //
+  // [[sycl::reqd_work_group_size]] and [[cl::reqd_work_group_size]] are
+  // available in SYCL modes and follow the SYCL rules.
+  // __attribute__((reqd_work_group_size)) is only available in OpenCL mode
+  // and follows the OpenCL rules.
   if (const auto *A = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
-    if (!((getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getXDimVal()) &&
-          (getExprValue(AL.getArgAsExpr(1), Ctx) <= *A->getYDimVal()) &&
-          (getExprValue(AL.getArgAsExpr(2), Ctx) <= *A->getZDimVal()))) {
+    bool CheckFirstArgument =
+        S.getLangOpts().OpenCL
+            ? getExprValue(AL.getArgAsExpr(0), Ctx) > *A->getZDimVal()
+            : getExprValue(AL.getArgAsExpr(0), Ctx) > *A->getXDimVal();
+    bool CheckSecondArgument =
+        getExprValue(AL.getArgAsExpr(1), Ctx) > *A->getYDimVal();
+    bool CheckThirdArgument =
+        S.getLangOpts().OpenCL
+            ? getExprValue(AL.getArgAsExpr(2), Ctx) > *A->getXDimVal()
+            : getExprValue(AL.getArgAsExpr(2), Ctx) > *A->getZDimVal();
+
+    if (CheckFirstArgument || CheckSecondArgument || CheckThirdArgument) {
       S.Diag(AL.getLoc(), diag::err_conflicting_sycl_function_attributes)
-          << AL << A->getSpelling();
+          << AL << A;
+      S.Diag(A->getLocation(), diag::note_conflicting_attribute);
       Result &= false;
     }
   }
@@ -3286,7 +3305,8 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) {
           (getExprValue(AL.getArgAsExpr(2), Ctx) >=
            getExprValue(A->getZDim(), Ctx)))) {
       S.Diag(AL.getLoc(), diag::err_conflicting_sycl_function_attributes)
-          << AL << A->getSpelling();
+          << AL << A;
+      S.Diag(A->getLocation(), diag::note_conflicting_attribute);
       Result &= false;
     }
   }
@@ -3562,6 +3582,23 @@ static bool InvalidWorkGroupSizeAttrs(const Expr *MGValue, const Expr *XDim,
            ZDimExpr->getResultAsAPSInt() != 1));
 }
 
+// If the [[intel::max_work_group_size(X, Y, Z)]] attribute is specified on
+// a declaration along with [[sycl::reqd_work_group_size(X1, Y1, Z1)]]
+// attribute, check to see if values of reqd_work_group_size arguments are
+// equal or less than values of max_work_group_size attribute arguments.
+static bool checkWorkGroupSizeAttrValues(const Expr *RWGS, const Expr *MWGS) {
+  // If any of the operand is still value dependent, we can't test anything.
+  const auto *RWGSCE = dyn_cast<ConstantExpr>(RWGS);
+  const auto *MWGSCE = dyn_cast<ConstantExpr>(MWGS);
+
+  if (!RWGSCE || !MWGSCE)
+    return false;
+
+  // Otherwise, check if value of reqd_work_group_size argument is
+  // greater than value of max_work_group_size attribute argument.
+  return RWGSCE->getResultAsAPSInt() > MWGSCE->getResultAsAPSInt();
+}
+
 void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
                                             const AttributeCommonInfo &CI,
                                             Expr *XDim, Expr *YDim,
@@ -3595,6 +3632,40 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
   if (!XDim || !YDim || !ZDim)
     return;
 
+  // If the [[intel::max_work_group_size(X, Y, Z)]] attribute is specified on
+  // a declaration along with [[sycl::reqd_work_group_size(X1, Y1, Z1)]]
+  // attribute, check to see if values of reqd_work_group_size arguments are
+  // equal or less than values of max_work_group_size attribute arguments.
+  //
+  // 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
+  // increments the fastest.
+  //
+  // [[sycl::reqd_work_group_size]] and [[cl::reqd_work_group_size]] are
+  // available in SYCL modes and follow the SYCL rules.
+  // __attribute__((reqd_work_group_size)) is only available in OpenCL mode
+  // and follows the OpenCL rules.
+  if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
+    bool CheckFirstArgument =
+        getLangOpts().OpenCL
+            ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), ZDim)
+            : checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), XDim);
+    bool CheckSecondArgument =
+        checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), YDim);
+    bool CheckThirdArgument =
+        getLangOpts().OpenCL
+            ? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), XDim)
+            : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), ZDim);
+
+    if (CheckFirstArgument || CheckSecondArgument || CheckThirdArgument) {
+      Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes)
+          << CI << DeclAttr;
+      Diag(DeclAttr->getLoc(), diag::note_conflicting_attribute);
+      return;
+    }
+  }
+
   // If the declaration has a SYCLIntelMaxWorkGroupSizeAttr, check to see if
   // the attribute holds equal values to (1, 1, 1) in case the value of
   // SYCLIntelMaxGlobalWorkDimAttr equals to 0.
@@ -3655,6 +3726,40 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr(
       return nullptr;
   }
 
+  // If the [[intel::max_work_group_size(X, Y, Z)]] attribute is specified on
+  // a declaration along with [[sycl::reqd_work_group_size(X1, Y1, Z1)]]
+  // attribute, check to see if values of reqd_work_group_size arguments are
+  // equal or less than values of max_work_group_size attribute arguments.
+  //
+  // 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
+  // increments the fastest.
+  //
+  // [[sycl::reqd_work_group_size]] and [[cl::reqd_work_group_size]] are
+  // available in SYCL modes and follow the SYCL rules.
+  // __attribute__((reqd_work_group_size)) is only available in OpenCL mode
+  // and follows the OpenCL rules.
+  if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
+    bool CheckFirstArgument =
+        getLangOpts().OpenCL
+            ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getZDim())
+            : checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getXDim());
+    bool CheckSecondArgument =
+        checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), A.getYDim());
+    bool CheckThirdArgument =
+        getLangOpts().OpenCL
+            ? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getXDim())
+            : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getZDim());
+
+    if (CheckFirstArgument || CheckSecondArgument || CheckThirdArgument) {
+      Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes)
+          << DeclAttr << &A;
+      Diag(A.getLoc(), diag::note_conflicting_attribute);
+      return nullptr;
+    }
+  }
+
   // If the declaration has a SYCLIntelMaxWorkGroupSizeAttr,
   // check to see if the attribute holds equal values to
   // (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr

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

@@ -35,12 +35,14 @@ struct Func {
 
 #ifdef TRIGGER_ERROR
 struct DAFuncObj {
-  [[intel::max_work_group_size(4, 4, 4)]]
-  [[cl::reqd_work_group_size(8, 8, 4)]] // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} \
-                                        // expected-warning{{attribute 'cl::reqd_work_group_size' is deprecated}} \
-                                        // expected-note{{did you mean to use 'sycl::reqd_work_group_size' instead?}}
-  void operator()() const {}
+  [[intel::max_work_group_size(4, 4, 4)]] // expected-note {{conflicting attribute is here}}
+  [[cl::reqd_work_group_size(8, 8, 4)]]   // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} \
+                                         // expected-warning{{attribute 'cl::reqd_work_group_size' is deprecated}} \
+                                         // expected-note{{did you mean to use 'sycl::reqd_work_group_size' instead?}}
+  void
+  operator()() const {}
 };
+
 #endif // TRIGGER_ERROR
 
 int main() {

clang/test/SemaSYCL/intel-max-work-group-size.cpp

@@ -30,6 +30,11 @@ class Functor {
   [[intel::max_work_group_size(16, 16, 16)]] [[intel::max_work_group_size(32, 32, 32)]] void operator()(int) const; // expected-warning {{attribute 'max_work_group_size' is already applied with different arguments}} expected-note {{previous attribute is here}}
 };
 
+class FunctorC {
+public:
+  [[sycl::reqd_work_group_size(64, 64, 64)]] [[intel::max_work_group_size(64, 64, 64)]] void operator()() const;
+  [[sycl::reqd_work_group_size(64, 64, 64)]] [[intel::max_work_group_size(16, 16, 16)]] void operator()(int) const; // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} expected-note {{conflicting attribute is here}}
+};
 // Ensure that template arguments behave appropriately based on instantiations.
 template <int N>
 [[intel::max_work_group_size(N, 1, 1)]] void f6(); // #f6
@@ -59,3 +64,48 @@ void instantiate() {
   // expected-note@#f7prev {{previous attribute is here}}
   f7<2, 2, 2>(); // expected-note {{in instantiation}}
 }
+
+// If the [[intel::max_work_group_size(X, Y, Z)]] attribute is specified on
+// a declaration along with [[sycl::reqd_work_group_size(X1, Y1, Z1)]]
+// attribute, check to see if values of reqd_work_group_size arguments are
+// equal or less than values coming from max_work_group_size attribute.
+[[sycl::reqd_work_group_size(64, 64, 64)]] // expected-note {{conflicting attribute is here}}
+[[intel::max_work_group_size(64, 16, 64)]] // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
+void
+f9() {}
+
+[[intel::max_work_group_size(4, 4, 4)]] void f10();
+[[sycl::reqd_work_group_size(2, 2, 2)]] void f10(); // OK
+
+[[sycl::reqd_work_group_size(2, 2, 2)]] [[intel::max_work_group_size(4, 4, 4)]] void f11() {} // OK
+
+// FIXME: We do not have support yet for checking
+// reqd_work_group_size and max_work_group_size
+// attributes when merging, so the test compiles without
+// any diagnostic when it shouldn't.
+[[sycl::reqd_work_group_size(64, 64, 64)]] void f12();
+[[intel::max_work_group_size(16, 16, 16)]] void f12(); // expected error but now OK.
+
+[[intel::max_work_group_size(16, 16, 16)]]      // expected-note {{conflicting attribute is here}}
+[[sycl::reqd_work_group_size(16, 64, 16)]] void // expected-error {{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}
+f13() {}
+
+[[intel::max_work_group_size(16, 16, 16)]] void f14(); // expected-note {{conflicting attribute is here}}
+[[sycl::reqd_work_group_size(64, 64, 64)]] void f14(); // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}
+
+[[cl::reqd_work_group_size(1, 2, 3)]] // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \
+		                      // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}}
+[[intel::max_work_group_size(1, 2, 3)]] void
+f15() {} // OK
+
+[[intel::max_work_group_size(2, 3, 7)]] void f16(); // expected-note {{conflicting attribute is here}}
+[[sycl::reqd_work_group_size(7, 3, 2)]] void f16(); // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}
+
+[[intel::max_work_group_size(1, 2, 3)]] [[sycl::reqd_work_group_size(1, 2, 3)]] void f17(){}; // OK
+
+[[sycl::reqd_work_group_size(16)]]            // expected-note {{conflicting attribute is here}}
+[[intel::max_work_group_size(1, 1, 16)]] void // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
+f18();
+
+[[intel::max_work_group_size(16, 16, 1)]] void f19();
+[[sycl::reqd_work_group_size(16, 16)]] void f19(); // OK

clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp

@@ -20,13 +20,8 @@ func1();
 
 #else
 //second case - expect error
-[[intel::max_work_group_size(4, 4, 4)]] // expected-note {{conflicting attribute is here}}
-void
-func2();
-
-[[sycl::reqd_work_group_size(8, 8, 8)]] // expected-note {{conflicting attribute is here}}
-void
-func2() {}
+[[intel::max_work_group_size(4, 4, 4)]] void func2();   // expected-note {{conflicting attribute is here}}
+[[sycl::reqd_work_group_size(8, 8, 8)]] void func2() {} // expected-error {{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}
 
 //third case - expect error
 [[sycl::reqd_work_group_size(4, 4, 4)]] // expected-note {{conflicting attribute is here}}
@@ -36,7 +31,7 @@ func3();
 [[sycl::reqd_work_group_size(1, 1, 1)]] // expected-note {{conflicting attribute is here}}
 void
 // expected-warning@+1 {{attribute 'reqd_work_group_size' is already applied with different arguments}}
-func3() {} // expected-error {{'reqd_work_group_size' attribute conflicts with ''reqd_work_group_size'' attribute}}
+func3() {} // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
 
 // fourth case - expect warning.
 [[intel::max_work_group_size(4, 4, 4)]] void func4();   // expected-note {{previous attribute is here}}
@@ -77,7 +72,7 @@ int main() {
 
 #else
     h.single_task<class test_kernel2>(
-        []() { func2(); }); // expected-error {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}}
+        []() { func2(); });
 
     h.single_task<class test_kernel3>(
         []() { func3(); });

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

@@ -26,9 +26,11 @@ __attribute__((reqd_work_group_size(4))) void four_yet_again(); // expected-erro
 class Functor32 {
 public:
   // expected-note@+3{{conflicting attribute is here}}
-  // expected-warning@+2{{attribute 'reqd_work_group_size' is already applied with different arguments}}
-  // expected-error@+1{{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
-  [[sycl::reqd_work_group_size(32, 1, 1)]] [[sycl::reqd_work_group_size(1, 1, 32)]] void operator()() const {}
+  // expected-warning@+3{{attribute 'reqd_work_group_size' is already applied with different arguments}}
+  // expected-error@+2{{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
+  [[sycl::reqd_work_group_size(32, 1, 1)]] // expected-note {{conflicting attribute is here}}
+  [[sycl::reqd_work_group_size(1, 1, 32)]] void
+  operator()() const {}
 };
 #endif // TRIGGER_ERROR
 

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

@@ -49,9 +49,11 @@ class Functor16 {
 class Functor32 {
 public:
   // expected-note@+3{{conflicting attribute is here}}
-  // expected-warning@+2{{attribute 'reqd_work_group_size' is already applied with different arguments}}
-  // expected-error@+1{{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
-  [[sycl::reqd_work_group_size(32, 1, 1)]] [[sycl::reqd_work_group_size(1, 1, 32)]] void operator()() const {}
+  // expected-warning@+3{{attribute 'reqd_work_group_size' is already applied with different arguments}}
+  // expected-error@+2 {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
+  [[sycl::reqd_work_group_size(32, 1, 1)]] // expected-note {{conflicting attribute is here}}
+  [[sycl::reqd_work_group_size(1, 1, 32)]] void
+  operator()() const {}
 };
 #endif
 class Functor16x16x16 {