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

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, this patch checks if values of reqd_work_group_size arguments are
equal or less than values of max_work_group_size attribute arguments.

Some of the test cases were missed during refactoring work with PGA function
attribute [[intel::max_work_group_size()]] on intel#5392

This patch fixes them.

Signed-off-by: Soumi Manna <[email protected]>

Author: smanna12

clang/lib/Sema/SemaDeclAttr.cpp

@@ -3523,6 +3523,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 *LHS, const Expr *RHS) {
+  // If any of the operand is still value dependent, we can't test anything.
+  const auto *LHSCE = dyn_cast<ConstantExpr>(LHS);
+  const auto *RHSCE = dyn_cast<ConstantExpr>(RHS);
+
+  if (!LHSCE || !RHSCE)
+    return false;
+
+  // Otherwise, check if value of reqd_work_group_size argument is
+  // equal or less than value of max_work_group_size attribute argument.
+  return !(LHSCE->getResultAsAPSInt() <= RHSCE->getResultAsAPSInt());
+}
+
 void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
                                             const AttributeCommonInfo &CI,
                                             Expr *XDim, Expr *YDim,
@@ -3556,6 +3573,20 @@ 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.
+  if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
+    if (checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), XDim),
+	checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), YDim),
+	checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), ZDim)) {
+      Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes)
+         << CI << DeclAttr->getSpelling();
+      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.
@@ -3616,6 +3647,20 @@ 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.
+  if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
+    if (checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getXDim()),
+        checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), A.getYDim()),
+        checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getZDim())) {
+      Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes)
+         << DeclAttr << A.getSpelling();
+      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.cpp

@@ -30,6 +30,12 @@ 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(16, 16, 16)]] void operator()() const; //expected-error{{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
+  [[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}}
+};
+
 // 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 +65,31 @@ 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)]]
+[[intel::max_work_group_size(16, 16, 16)]]  //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)]]
+[[sycl::reqd_work_group_size(64, 64, 64)]] void f13() {} // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}
+
+[[intel::max_work_group_size(16, 16, 16)]] void f14();
+[[sycl::reqd_work_group_size(64, 64, 64)]] void f14(); // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}

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();
+[[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}}
@@ -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(); });