From 68c459a7fc3c24d867a63deb0f45804d746cb815 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Wed, 16 Feb 2022 05:42:24 -0800
Subject: [PATCH 01/18] [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 https://github.com/intel/llvm/pull/5392

This patch fixes them.

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp               | 45 +++++++++++++++++++
 .../SemaSYCL/intel-max-work-group-size.cpp    | 34 ++++++++++++++
 .../redeclaration-attribute-propagation.cpp   | 11 ++---
 3 files changed, 82 insertions(+), 8 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 186cdf508fcad..0ba529d1f1050 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/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
diff --git a/clang/test/SemaSYCL/intel-max-work-group-size.cpp b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
index 8c8088a105813..fdeca3546ccda 100644
--- a/clang/test/SemaSYCL/intel-max-work-group-size.cpp
+++ b/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}}
diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp
index ccdbfbfb9f20e..169fc1e4ede53 100644
--- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp
+++ b/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(); });

From 651869ce3086c718da223745c82e67a90dc72837 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Fri, 18 Feb 2022 11:44:47 -0800
Subject: [PATCH 02/18] address review comments

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp               |  6 +++++-
 .../intel-max-work-group-size-device.cpp      |  2 +-
 .../SemaSYCL/intel-max-work-group-size.cpp    | 16 +++++++++-------
 .../redeclaration-attribute-propagation.cpp   |  2 +-
 .../SemaSYCL/reqd-sub-group-size-host.cpp     | 19 +++++++++++++++++++
 ...eqd-work-group-size-device-direct-prop.cpp |  7 ++++---
 .../SemaSYCL/reqd-work-group-size-device.cpp  |  7 ++++---
 7 files changed, 43 insertions(+), 16 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 0ba529d1f1050..acfe93e9b8c18 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3259,6 +3259,7 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) {
           (getExprValue(AL.getArgAsExpr(2), Ctx) <= *A->getZDimVal()))) {
       S.Diag(AL.getLoc(), diag::err_conflicting_sycl_function_attributes)
           << AL << A->getSpelling();
+      S.Diag(A->getLocation(), diag::note_conflicting_attribute);
       Result &= false;
     }
   }
@@ -3272,6 +3273,7 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) {
            getExprValue(A->getZDim(), Ctx)))) {
       S.Diag(AL.getLoc(), diag::err_conflicting_sycl_function_attributes)
           << AL << A->getSpelling();
+      S.Diag(A->getLocation(), diag::note_conflicting_attribute);
       Result &= false;
     }
   }
@@ -3537,7 +3539,7 @@ static bool checkWorkGroupSizeAttrValues(const Expr *LHS, const Expr *RHS) {
 
   // 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());
+  return LHSCE->getResultAsAPSInt() > RHSCE->getResultAsAPSInt();
 }
 
 void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
@@ -3583,6 +3585,7 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
 	checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), ZDim)) {
       Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes)
          << CI << DeclAttr->getSpelling();
+      Diag(DeclAttr->getLocation(), diag::note_conflicting_attribute);
       return;
     }
   }
@@ -3657,6 +3660,7 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr(
         checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getZDim())) {
       Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes)
          << DeclAttr << A.getSpelling();
+      Diag(A.getLoc(), diag::note_conflicting_attribute);
       return nullptr;
     }
   }
diff --git a/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp b/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp
index 70b8553f915ff..4ed0981b472e3 100644
--- a/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp
@@ -35,7 +35,7 @@ struct Func {
 
 #ifdef TRIGGER_ERROR
 struct DAFuncObj {
-  [[intel::max_work_group_size(4, 4, 4)]]
+  [[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?}}
diff --git a/clang/test/SemaSYCL/intel-max-work-group-size.cpp b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
index fdeca3546ccda..fe5cd924e2115 100644
--- a/clang/test/SemaSYCL/intel-max-work-group-size.cpp
+++ b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
@@ -32,8 +32,10 @@ class Functor {
 
 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}}
+  [[sycl::reqd_work_group_size(64, 64, 64)]] // expected-note {{conflicting attribute is here}}
+  [[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)]] // expected-note {{conflicting attribute is here}}
+  [[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.
@@ -71,8 +73,8 @@ void instantiate() {
 // 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}}
+[[sycl::reqd_work_group_size(64, 64, 64)]]  // expected-note {{conflicting attribute is here}}
+[[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();
@@ -88,8 +90,8 @@ void f9() {}
 [[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)]] // expected-note {{conflicting attribute is here}}
+[[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();
+[[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}}
diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp
index 169fc1e4ede53..cb3a00ae89e88 100644
--- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp
+++ b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp
@@ -20,7 +20,7 @@ func1();
 
 #else
 //second case - expect error
-[[intel::max_work_group_size(4, 4, 4)]] 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
diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-host.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-host.cpp
index ab11ef9a19ce9..f18a8edc2009a 100644
--- a/clang/test/SemaSYCL/reqd-sub-group-size-host.cpp
+++ b/clang/test/SemaSYCL/reqd-sub-group-size-host.cpp
@@ -7,3 +7,22 @@ class Functor {
 public:
   [[intel::reqd_sub_group_size(16)]] void operator()() {}
 };
+
+// RUN: %clang_cc1 -fsycl-is-host -triple x86_64-pc-linux-gnu -fsyntax-only -verify %s -DSYCLHOST
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fsyntax-only -verify %s
+
+#ifdef SYCLHOST
+// expected-no-diagnostics
+#endif
+
+void foo()
+{
+#ifndef SYCLHOST
+// expected-warning@+2 {{'reqd_sub_group_size' attribute ignored}}
+#endif
+  class Functor {
+public:
+  [[intel::reqd_sub_group_size(16)]] void operator()() {}
+};
+}
+
diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp
index 88b8ab6a3a3a1..1b28b759b110f 100644
--- a/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp
+++ b/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp
@@ -26,9 +26,10 @@ __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
 
diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp
index 700f861c7fce9..86634a00aad9a 100644
--- a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp
+++ b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp
@@ -49,9 +49,10 @@ 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 {

From 369ae6bf9d5912123546937e9f9e503d0f0c2332 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Fri, 18 Feb 2022 11:47:14 -0800
Subject: [PATCH 03/18] Remove redundant change

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 .../SemaSYCL/reqd-sub-group-size-host.cpp     | 19 -------------------
 1 file changed, 19 deletions(-)

diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-host.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-host.cpp
index f18a8edc2009a..ab11ef9a19ce9 100644
--- a/clang/test/SemaSYCL/reqd-sub-group-size-host.cpp
+++ b/clang/test/SemaSYCL/reqd-sub-group-size-host.cpp
@@ -7,22 +7,3 @@ class Functor {
 public:
   [[intel::reqd_sub_group_size(16)]] void operator()() {}
 };
-
-// RUN: %clang_cc1 -fsycl-is-host -triple x86_64-pc-linux-gnu -fsyntax-only -verify %s -DSYCLHOST
-// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fsyntax-only -verify %s
-
-#ifdef SYCLHOST
-// expected-no-diagnostics
-#endif
-
-void foo()
-{
-#ifndef SYCLHOST
-// expected-warning@+2 {{'reqd_sub_group_size' attribute ignored}}
-#endif
-  class Functor {
-public:
-  [[intel::reqd_sub_group_size(16)]] void operator()() {}
-};
-}
-

From 5b68889275ed765caed979413bf6b26da8d99113 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Tue, 22 Feb 2022 09:14:08 -0800
Subject: [PATCH 04/18] address review comments

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/include/clang/Basic/DiagnosticSemaKinds.td       |  2 +-
 clang/lib/Sema/SemaDeclAttr.cpp                        | 10 +++++-----
 .../SemaSYCL/redeclaration-attribute-propagation.cpp   |  2 +-
 3 files changed, 7 insertions(+), 7 deletions(-)

diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 6b634ef0d5335..0fcae19fe6e62 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -11600,7 +11600,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<
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index acfe93e9b8c18..f736777781611 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3258,7 +3258,7 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) {
           (getExprValue(AL.getArgAsExpr(1), Ctx) <= *A->getYDimVal()) &&
           (getExprValue(AL.getArgAsExpr(2), Ctx) <= *A->getZDimVal()))) {
       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;
     }
@@ -3272,7 +3272,7 @@ 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;
     }
@@ -3584,8 +3584,8 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
 	checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), YDim),
 	checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), ZDim)) {
       Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes)
-         << CI << DeclAttr->getSpelling();
-      Diag(DeclAttr->getLocation(), diag::note_conflicting_attribute);
+         << CI << DeclAttr;
+      Diag(DeclAttr->getLoc(), diag::note_conflicting_attribute);
       return;
     }
   }
@@ -3659,7 +3659,7 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr(
         checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), A.getYDim()),
         checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getZDim())) {
       Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes)
-         << DeclAttr << A.getSpelling();
+         << DeclAttr << &A;
       Diag(A.getLoc(), diag::note_conflicting_attribute);
       return nullptr;
     }
diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp
index cb3a00ae89e88..437b3c86ed1c0 100644
--- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp
+++ b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp
@@ -31,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}}

From e7570fce8cf1dcf900494475c845fc3e142810cd Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Wed, 23 Feb 2022 10:24:39 -0800
Subject: [PATCH 05/18] Address review comments

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp                   | 12 ++++++------
 clang/test/SemaSYCL/intel-max-work-group-size.cpp |  6 ++++++
 2 files changed, 12 insertions(+), 6 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index f736777781611..e8cbbc13513b9 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3529,17 +3529,17 @@ static bool InvalidWorkGroupSizeAttrs(const Expr *MGValue, const Expr *XDim,
 // 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) {
+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 *LHSCE = dyn_cast<ConstantExpr>(LHS);
-  const auto *RHSCE = dyn_cast<ConstantExpr>(RHS);
+  const auto *RWGSCE = dyn_cast<ConstantExpr>(RWGS);
+  const auto *MWGSCE = dyn_cast<ConstantExpr>(MWGS);
 
-  if (!LHSCE || !RHSCE)
+  if (!RWGSCE || !MWGSCE)
     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();
+  // greater than value of max_work_group_size attribute argument.
+  return RWGSCE->getResultAsAPSInt() > MWGSCE->getResultAsAPSInt();
 }
 
 void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
diff --git a/clang/test/SemaSYCL/intel-max-work-group-size.cpp b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
index fe5cd924e2115..f5db13ad25d00 100644
--- a/clang/test/SemaSYCL/intel-max-work-group-size.cpp
+++ b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
@@ -95,3 +95,9 @@ void f9() {}
 
 [[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}}
+
+[[intel::max_work_group_size(2, 2, 2)]] void f15();
+[[sycl::reqd_work_group_size(2, 2, 2)]] void f15(); // OK
+
+[[sycl::reqd_work_group_size(4, 4, 4)]]
+[[intel::max_work_group_size(4, 4, 4)]] void f16() {} // OK

From 2cb7c5c06dde53b118b6eca9561d1d3b9b9a5b3b Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Fri, 25 Feb 2022 11:46:02 -0800
Subject: [PATCH 06/18] add reqd_work_group_size spellings and orderings

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp               | 62 ++++++++++++++++---
 .../intel-max-global-work-dim-device.cpp      | 12 ++--
 .../intel-max-work-group-size-device.cpp      |  1 +
 .../SemaSYCL/intel-max-work-group-size.cpp    | 21 +++++--
 4 files changed, 77 insertions(+), 19 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index e8cbbc13513b9..3ed78fef10def 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3253,10 +3253,24 @@ 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.
+  //
+  // The arguments to max_work_group_size are ordered based on which index
+  // increments the fastest. In SYCL, the last argument is the index that
+  // increments the fastest.
   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->getXDimVal()
+			      : getExprValue(AL.getArgAsExpr(2), Ctx) <= *A->getXDimVal();
+    bool checkSecondArgument = getExprValue(AL.getArgAsExpr(1), Ctx) <= *A->getYDimVal();
+    bool checkThirdArgument = S.getLangOpts().OpenCL
+                              ? getExprValue(AL.getArgAsExpr(3), Ctx) <= *A->getZDimVal()
+			      : getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getZDimVal();
+
+    if (!(checkFirstArgument &&  checkSecondArgument && checkThirdArgument)) {
       S.Diag(AL.getLoc(), diag::err_conflicting_sycl_function_attributes)
           << AL << A;
       S.Diag(A->getLocation(), diag::note_conflicting_attribute);
@@ -3579,10 +3593,25 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
   // 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.
+  //
+  // The arguments to max_work_group_size are ordered based on which index
+  // increments the fastest. In SYCL, the last argument is the index that
+  // increments the fastest.
   if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
-    if (checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), XDim),
-	checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), YDim),
-	checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), ZDim)) {
+    bool checkFirstArgument = getLangOpts().OpenCL
+                              ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), XDim)
+                              : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), XDim);
+    bool checkSecondArgument = checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), YDim);
+    bool checkThirdArgument = getLangOpts().OpenCL
+                              ? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), ZDim)
+                              : checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), ZDim);
+
+    if (checkFirstArgument &&  checkSecondArgument && checkThirdArgument) {
       Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes)
          << CI << DeclAttr;
       Diag(DeclAttr->getLoc(), diag::note_conflicting_attribute);
@@ -3654,10 +3683,25 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr(
   // 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.
+  //
+  // The arguments to max_work_group_size are ordered based on which index
+  // increments the fastest. In SYCL, the last argument is the index that
+  // increments the fastest.
   if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
-    if (checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getXDim()),
-        checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), A.getYDim()),
-        checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getZDim())) {
+    bool checkFirstArgument = getLangOpts().OpenCL
+	                      ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getXDim())
+			      : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getXDim());
+    bool checkSecondArgument = checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), A.getYDim());
+    bool checkThirdArgument = getLangOpts().OpenCL
+	                      ? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getZDim())
+			      : checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getZDim());
+
+    if (checkFirstArgument &&  checkSecondArgument && checkThirdArgument) {
       Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes)
          << DeclAttr << &A;
       Diag(A.getLoc(), diag::note_conflicting_attribute);
diff --git a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
index 2a9d375a010b0..c154357509755 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
@@ -62,8 +62,8 @@ struct TRIFuncObjGood1 {
 
 struct TRIFuncObjGood2 {
   [[intel::max_global_work_dim(3)]]
-  [[intel::max_work_group_size(8, 1, 1)]]
-  [[sycl::reqd_work_group_size(4, 1, 1)]] void
+  [[intel::max_work_group_size(8, 1, 8)]]
+  [[sycl::reqd_work_group_size(4, 1, 8)]] void
   operator()() const {}
 };
 
@@ -300,8 +300,8 @@ int main() {
     // CHECK-NEXT:  value: Int 1
     // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
-    // CHECK-NEXT:  value: Int 1
-    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    // CHECK-NEXT:  value: Int 8
+    // CHECK-NEXT:  IntegerLiteral{{.*}}8{{$}}
     // CHECK:       ReqdWorkGroupSizeAttr {{.*}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 4
@@ -310,8 +310,8 @@ int main() {
     // CHECK-NEXT:  value: Int 1
     // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
-    // CHECK-NEXT:  value: Int 1
-    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    // CHECK-NEXT:  value: Int 8
+    // CHECK-NEXT:  IntegerLiteral{{.*}}8{{$}}
 
     h.single_task<class test_kernel5>(TRIFuncObjGood3());
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel5
diff --git a/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp b/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp
index 4ed0981b472e3..b86ce667576f2 100644
--- a/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp
@@ -41,6 +41,7 @@ struct DAFuncObj {
                                         // expected-note{{did you mean to use 'sycl::reqd_work_group_size' instead?}}
   void operator()() const {}
 };
+
 #endif // TRIGGER_ERROR
 
 int main() {
diff --git a/clang/test/SemaSYCL/intel-max-work-group-size.cpp b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
index f5db13ad25d00..430e017d8ab6e 100644
--- a/clang/test/SemaSYCL/intel-max-work-group-size.cpp
+++ b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
@@ -96,8 +96,21 @@ void f9() {}
 [[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}}
 
-[[intel::max_work_group_size(2, 2, 2)]] void f15();
-[[sycl::reqd_work_group_size(2, 2, 2)]] void f15(); // OK
+[[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 f16() {} // OK
 
-[[sycl::reqd_work_group_size(4, 4, 4)]]
-[[intel::max_work_group_size(4, 4, 4)]] void f16() {} // OK
+[[intel::max_work_group_size(1, 1, 16)]] void f17();
+[[sycl::reqd_work_group_size(16)]] void f17(); // OK
+
+[[sycl::reqd_work_group_size(16)]]
+[[intel::max_work_group_size(1, 1, 16)]] void f18() {} // OK
+
+[[sycl::reqd_work_group_size(16, 16)]]
+[[intel::max_work_group_size(1, 16, 16)]] void f19() {} // OK
+
+[[intel::max_work_group_size(2, 3, 7)]] void f20();
+[[sycl::reqd_work_group_size(7, 3, 2)]] void f20(); // OK
+
+[[intel::max_work_group_size(1, 2, 3)]] // expected-note {{conflicting attribute is here}}
+[[sycl::reqd_work_group_size(1, 2, 3)]] void f21() {}; // expected-error {{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}

From 5413321fd3dfb674b6a8c3a05ca479e9735f4615 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Fri, 25 Feb 2022 11:56:32 -0800
Subject: [PATCH 07/18] Fix formats

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp | 33 ++++++++++++++++++---------------
 1 file changed, 18 insertions(+), 15 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 3ed78fef10def..e2343faec53ad 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3263,12 +3263,13 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) {
   // increments the fastest.
   if (const auto *A = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
     bool checkFirstArgument = S.getLangOpts().OpenCL
-                              ? getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getXDimVal()
-			      : getExprValue(AL.getArgAsExpr(2), Ctx) <= *A->getXDimVal();
-    bool checkSecondArgument = getExprValue(AL.getArgAsExpr(1), Ctx) <= *A->getYDimVal();
+              ? getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getXDimVal()
+	      : getExprValue(AL.getArgAsExpr(2), Ctx) <= *A->getXDimVal();
+    bool checkSecondArgument =
+	 getExprValue(AL.getArgAsExpr(1), Ctx) <= *A->getYDimVal();
     bool checkThirdArgument = S.getLangOpts().OpenCL
-                              ? getExprValue(AL.getArgAsExpr(3), Ctx) <= *A->getZDimVal()
-			      : getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getZDimVal();
+              ? getExprValue(AL.getArgAsExpr(3), Ctx) <= *A->getZDimVal()
+	      : getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getZDimVal();
 
     if (!(checkFirstArgument &&  checkSecondArgument && checkThirdArgument)) {
       S.Diag(AL.getLoc(), diag::err_conflicting_sycl_function_attributes)
@@ -3604,12 +3605,13 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
   // increments the fastest.
   if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
     bool checkFirstArgument = getLangOpts().OpenCL
-                              ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), XDim)
-                              : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), XDim);
-    bool checkSecondArgument = checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), YDim);
+              ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), XDim)
+              : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), XDim);
+    bool checkSecondArgument =
+         checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), YDim);
     bool checkThirdArgument = getLangOpts().OpenCL
-                              ? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), ZDim)
-                              : checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), ZDim);
+              ? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), ZDim)
+              : checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), ZDim);
 
     if (checkFirstArgument &&  checkSecondArgument && checkThirdArgument) {
       Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes)
@@ -3694,12 +3696,13 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr(
   // increments the fastest.
   if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
     bool checkFirstArgument = getLangOpts().OpenCL
-	                      ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getXDim())
-			      : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getXDim());
-    bool checkSecondArgument = checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), A.getYDim());
+	      ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getXDim())
+	      : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getXDim());
+    bool checkSecondArgument =
+	 checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), A.getYDim());
     bool checkThirdArgument = getLangOpts().OpenCL
-	                      ? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getZDim())
-			      : checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getZDim());
+	      ? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getZDim())
+	      : checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getZDim());
 
     if (checkFirstArgument &&  checkSecondArgument && checkThirdArgument) {
       Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes)

From 238abe13cc18775d1c218d5d92fdb936835f57af Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Fri, 25 Feb 2022 12:03:36 -0800
Subject: [PATCH 08/18] Fix formats

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp | 12 ++++++------
 1 file changed, 6 insertions(+), 6 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index e2343faec53ad..9c0a7024a0201 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3266,12 +3266,12 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) {
               ? getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getXDimVal()
 	      : getExprValue(AL.getArgAsExpr(2), Ctx) <= *A->getXDimVal();
     bool checkSecondArgument =
-	 getExprValue(AL.getArgAsExpr(1), Ctx) <= *A->getYDimVal();
+	      getExprValue(AL.getArgAsExpr(1), Ctx) <= *A->getYDimVal();
     bool checkThirdArgument = S.getLangOpts().OpenCL
               ? getExprValue(AL.getArgAsExpr(3), Ctx) <= *A->getZDimVal()
 	      : getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getZDimVal();
 
-    if (!(checkFirstArgument &&  checkSecondArgument && checkThirdArgument)) {
+    if (!(checkFirstArgument && checkSecondArgument && checkThirdArgument)) {
       S.Diag(AL.getLoc(), diag::err_conflicting_sycl_function_attributes)
           << AL << A;
       S.Diag(A->getLocation(), diag::note_conflicting_attribute);
@@ -3608,12 +3608,12 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
               ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), XDim)
               : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), XDim);
     bool checkSecondArgument =
-         checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), YDim);
+              checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), YDim);
     bool checkThirdArgument = getLangOpts().OpenCL
               ? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), ZDim)
               : checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), ZDim);
 
-    if (checkFirstArgument &&  checkSecondArgument && checkThirdArgument) {
+    if (checkFirstArgument && checkSecondArgument && checkThirdArgument) {
       Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes)
          << CI << DeclAttr;
       Diag(DeclAttr->getLoc(), diag::note_conflicting_attribute);
@@ -3699,12 +3699,12 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr(
 	      ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getXDim())
 	      : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getXDim());
     bool checkSecondArgument =
-	 checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), A.getYDim());
+	      checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), A.getYDim());
     bool checkThirdArgument = getLangOpts().OpenCL
 	      ? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getZDim())
 	      : checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getZDim());
 
-    if (checkFirstArgument &&  checkSecondArgument && checkThirdArgument) {
+    if (checkFirstArgument && checkSecondArgument && checkThirdArgument) {
       Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes)
          << DeclAttr << &A;
       Diag(A.getLoc(), diag::note_conflicting_attribute);

From 05e13035114ca31022ea57b943a448b53897992e Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Sun, 27 Feb 2022 07:25:07 -0800
Subject: [PATCH 09/18] Address review comments

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp               | 28 +++++++++----------
 .../intel-max-global-work-dim-device.cpp      | 12 ++++----
 .../SemaSYCL/intel-max-work-group-size.cpp    | 23 +++++----------
 3 files changed, 27 insertions(+), 36 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 9c0a7024a0201..ed45f2390440e 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3263,13 +3263,13 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) {
   // increments the fastest.
   if (const auto *A = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
     bool checkFirstArgument = S.getLangOpts().OpenCL
-              ? getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getXDimVal()
-	      : getExprValue(AL.getArgAsExpr(2), Ctx) <= *A->getXDimVal();
+              ? 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(3), Ctx) <= *A->getZDimVal()
-	      : getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getZDimVal();
+              ? 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)
@@ -3605,15 +3605,15 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
   // increments the fastest.
   if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
     bool checkFirstArgument = getLangOpts().OpenCL
-              ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), XDim)
-              : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), XDim);
+              ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), ZDim)
+              : checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), XDim);
     bool checkSecondArgument =
               checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), YDim);
     bool checkThirdArgument = getLangOpts().OpenCL
-              ? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), ZDim)
-              : checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), ZDim);
+              ? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), XDim)
+              : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), ZDim);
 
-    if (checkFirstArgument && checkSecondArgument && checkThirdArgument) {
+    if (checkFirstArgument || checkSecondArgument || checkThirdArgument) {
       Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes)
          << CI << DeclAttr;
       Diag(DeclAttr->getLoc(), diag::note_conflicting_attribute);
@@ -3696,15 +3696,15 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr(
   // increments the fastest.
   if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
     bool checkFirstArgument = getLangOpts().OpenCL
-	      ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getXDim())
-	      : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getXDim());
+	      ? 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.getZDim())
-	      : checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getZDim());
+	      ? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getXDim())
+	      : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getZDim());
 
-    if (checkFirstArgument && checkSecondArgument && checkThirdArgument) {
+    if (checkFirstArgument || checkSecondArgument || checkThirdArgument) {
       Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes)
          << DeclAttr << &A;
       Diag(A.getLoc(), diag::note_conflicting_attribute);
diff --git a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
index c154357509755..2a9d375a010b0 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
@@ -62,8 +62,8 @@ struct TRIFuncObjGood1 {
 
 struct TRIFuncObjGood2 {
   [[intel::max_global_work_dim(3)]]
-  [[intel::max_work_group_size(8, 1, 8)]]
-  [[sycl::reqd_work_group_size(4, 1, 8)]] void
+  [[intel::max_work_group_size(8, 1, 1)]]
+  [[sycl::reqd_work_group_size(4, 1, 1)]] void
   operator()() const {}
 };
 
@@ -300,8 +300,8 @@ int main() {
     // CHECK-NEXT:  value: Int 1
     // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
-    // CHECK-NEXT:  value: Int 8
-    // CHECK-NEXT:  IntegerLiteral{{.*}}8{{$}}
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
     // CHECK:       ReqdWorkGroupSizeAttr {{.*}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 4
@@ -310,8 +310,8 @@ int main() {
     // CHECK-NEXT:  value: Int 1
     // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
-    // CHECK-NEXT:  value: Int 8
-    // CHECK-NEXT:  IntegerLiteral{{.*}}8{{$}}
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
 
     h.single_task<class test_kernel5>(TRIFuncObjGood3());
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel5
diff --git a/clang/test/SemaSYCL/intel-max-work-group-size.cpp b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
index 430e017d8ab6e..3121436e56801 100644
--- a/clang/test/SemaSYCL/intel-max-work-group-size.cpp
+++ b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
@@ -74,7 +74,7 @@ void instantiate() {
 // 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(16, 16, 16)]]  // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
+[[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();
@@ -91,26 +91,17 @@ void f9() {}
 [[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(64, 64, 64)]] void f13() {} // expected-error {{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}
+[[sycl::reqd_work_group_size(16, 64, 16)]] 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(); // 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 f16() {} // OK
+[[intel::max_work_group_size(1, 2, 3)]] void f15() {} // OK
 
-[[intel::max_work_group_size(1, 1, 16)]] void f17();
-[[sycl::reqd_work_group_size(16)]] void f17(); // 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}}
 
-[[sycl::reqd_work_group_size(16)]]
-[[intel::max_work_group_size(1, 1, 16)]] void f18() {} // OK
-
-[[sycl::reqd_work_group_size(16, 16)]]
-[[intel::max_work_group_size(1, 16, 16)]] void f19() {} // OK
-
-[[intel::max_work_group_size(2, 3, 7)]] void f20();
-[[sycl::reqd_work_group_size(7, 3, 2)]] void f20(); // OK
-
-[[intel::max_work_group_size(1, 2, 3)]] // expected-note {{conflicting attribute is here}}
-[[sycl::reqd_work_group_size(1, 2, 3)]] void f21() {}; // 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

From fcc265450136c148f0827144ecd9e61ab323d88d Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Sun, 27 Feb 2022 07:29:58 -0800
Subject: [PATCH 10/18] Update comments

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp | 12 ------------
 1 file changed, 12 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index ed45f2390440e..cfc4204375c6d 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3257,10 +3257,6 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) {
   // 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.
-  //
-  // The arguments to max_work_group_size are ordered based on which index
-  // increments the fastest. In SYCL, the last argument is the index that
-  // increments the fastest.
   if (const auto *A = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
     bool checkFirstArgument = S.getLangOpts().OpenCL
               ? getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getZDimVal()
@@ -3599,10 +3595,6 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
   // 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.
-  //
-  // The arguments to max_work_group_size are ordered based on which index
-  // increments the fastest. In SYCL, the last argument is the index that
-  // increments the fastest.
   if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
     bool checkFirstArgument = getLangOpts().OpenCL
               ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), ZDim)
@@ -3690,10 +3682,6 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr(
   // 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.
-  //
-  // The arguments to max_work_group_size are ordered based on which index
-  // increments the fastest. In SYCL, the last argument is the index that
-  // increments the fastest.
   if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
     bool checkFirstArgument = getLangOpts().OpenCL
 	      ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getZDim())

From 741832eccece0dd1471609fe2c5369dfcec4aed5 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Sun, 27 Feb 2022 08:07:33 -0800
Subject: [PATCH 11/18] add comments for different spellings of
 reqqd_work_group attribute

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp | 15 +++++++++++++++
 1 file changed, 15 insertions(+)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index cfc4204375c6d..9b5eb2142fa84 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3257,6 +3257,11 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) {
   // 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>()) {
     bool checkFirstArgument = S.getLangOpts().OpenCL
               ? getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getZDimVal()
@@ -3595,6 +3600,11 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
   // 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)
@@ -3682,6 +3692,11 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr(
   // 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())

From 113294ec4557847f83c6e0a1519e2df99d99f6e2 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Mon, 28 Feb 2022 06:21:20 -0800
Subject: [PATCH 12/18] Add test cases for default values of
 reqd_work_group_size

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/test/SemaSYCL/intel-max-work-group-size.cpp | 6 ++++++
 1 file changed, 6 insertions(+)

diff --git a/clang/test/SemaSYCL/intel-max-work-group-size.cpp b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
index 3121436e56801..35e4b23f78d60 100644
--- a/clang/test/SemaSYCL/intel-max-work-group-size.cpp
+++ b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
@@ -105,3 +105,9 @@ void f9() {}
 
 [[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 f18();  // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
+
+[[intel::max_work_group_size(16, 16, 1)]] void f19();
+[[sycl::reqd_work_group_size(16, 16)]] void f19(); // OK

From a8c96aac53e7127a9b42b0b7d6c11b480883db6c Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Tue, 1 Mar 2022 09:24:42 -0800
Subject: [PATCH 13/18] Fix Coding style guideline

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp | 22 +++++++++++-----------
 1 file changed, 11 insertions(+), 11 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 85d5c67c3cb1c..d742b11c04f12 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3278,16 +3278,16 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) {
   // __attribute__((reqd_work_group_size)) is only available in OpenCL mode
   // and follows the OpenCL rules.
   if (const auto *A = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
-    bool checkFirstArgument = S.getLangOpts().OpenCL
+    bool CheckFirstArgument = S.getLangOpts().OpenCL
               ? getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getZDimVal()
 	      : getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getXDimVal();
-    bool checkSecondArgument =
+    bool CheckSecondArgument =
 	      getExprValue(AL.getArgAsExpr(1), Ctx) <= *A->getYDimVal();
-    bool checkThirdArgument = S.getLangOpts().OpenCL
+    bool CheckThirdArgument = S.getLangOpts().OpenCL
               ? getExprValue(AL.getArgAsExpr(2), Ctx) <= *A->getXDimVal()
 	      : getExprValue(AL.getArgAsExpr(2), Ctx) <= *A->getZDimVal();
 
-    if (!(checkFirstArgument && checkSecondArgument && checkThirdArgument)) {
+    if (!(CheckFirstArgument && CheckSecondArgument && CheckThirdArgument)) {
       S.Diag(AL.getLoc(), diag::err_conflicting_sycl_function_attributes)
           << AL << A;
       S.Diag(A->getLocation(), diag::note_conflicting_attribute);
@@ -3621,16 +3621,16 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
   // __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
+    bool CheckFirstArgument = getLangOpts().OpenCL
               ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), ZDim)
               : checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), XDim);
-    bool checkSecondArgument =
+    bool CheckSecondArgument =
               checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), YDim);
-    bool checkThirdArgument = getLangOpts().OpenCL
+    bool CheckThirdArgument = getLangOpts().OpenCL
               ? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), XDim)
               : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), ZDim);
 
-    if (checkFirstArgument || checkSecondArgument || checkThirdArgument) {
+    if (CheckFirstArgument || CheckSecondArgument || CheckThirdArgument) {
       Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes)
          << CI << DeclAttr;
       Diag(DeclAttr->getLoc(), diag::note_conflicting_attribute);
@@ -3713,16 +3713,16 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr(
   // __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
+    bool CheckFirstArgument = getLangOpts().OpenCL
 	      ? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getZDim())
 	      : checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getXDim());
-    bool checkSecondArgument =
+    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) {
+    if (CheckFirstArgument || CheckSecondArgument || CheckThirdArgument) {
       Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes)
          << DeclAttr << &A;
       Diag(A.getLoc(), diag::note_conflicting_attribute);

From 154d81fd10cb8de055567bbe9a52ddc6115af85f Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Tue, 1 Mar 2022 09:31:26 -0800
Subject: [PATCH 14/18] Fix Coding style guideline

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index d742b11c04f12..8954c74499ca0 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3718,7 +3718,7 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr(
 	      : checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getXDim());
     bool CheckSecondArgument =
 	      checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), A.getYDim());
-    bool checkThirdArgument = getLangOpts().OpenCL
+    bool CheckThirdArgument = getLangOpts().OpenCL
 	      ? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getXDim())
 	      : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getZDim());
 

From 6f224d06c2c81066ea9fda6b185ff3207d2ee630 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Wed, 2 Mar 2022 07:39:06 -0800
Subject: [PATCH 15/18] Fix test

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/test/SemaSYCL/intel-max-work-group-size.cpp | 8 +++-----
 1 file changed, 3 insertions(+), 5 deletions(-)

diff --git a/clang/test/SemaSYCL/intel-max-work-group-size.cpp b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
index 35e4b23f78d60..8421a5debbaf8 100644
--- a/clang/test/SemaSYCL/intel-max-work-group-size.cpp
+++ b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
@@ -32,12 +32,10 @@ class Functor {
 
 class FunctorC {
 public:
-  [[sycl::reqd_work_group_size(64, 64, 64)]] // expected-note {{conflicting attribute is here}}
-  [[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)]] // expected-note {{conflicting attribute is here}}
-  [[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}}
-};
+  [[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

From 4fb6c9b64d43e679227d1e2fe102fb029f197391 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Wed, 2 Mar 2022 07:41:16 -0800
Subject: [PATCH 16/18] remove exrtra space

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/test/SemaSYCL/intel-max-work-group-size.cpp | 1 -
 1 file changed, 1 deletion(-)

diff --git a/clang/test/SemaSYCL/intel-max-work-group-size.cpp b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
index 8421a5debbaf8..abc2c2ec0e6e6 100644
--- a/clang/test/SemaSYCL/intel-max-work-group-size.cpp
+++ b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
@@ -34,7 +34,6 @@ 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>

From 330fc3b9fd0f2ae4dc26871fd718faa12f15c9f7 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Wed, 2 Mar 2022 08:42:12 -0800
Subject: [PATCH 17/18] address review comments and format issues

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp               | 56 ++++++++++---------
 .../intel-max-work-group-size-device.cpp      |  9 +--
 .../SemaSYCL/intel-max-work-group-size.cpp    | 23 ++++----
 .../redeclaration-attribute-propagation.cpp   |  2 +-
 ...eqd-work-group-size-device-direct-prop.cpp |  3 +-
 .../SemaSYCL/reqd-work-group-size-device.cpp  |  3 +-
 6 files changed, 53 insertions(+), 43 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 8954c74499ca0..20dce226a9f3d 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3278,16 +3278,18 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) {
   // __attribute__((reqd_work_group_size)) is only available in OpenCL mode
   // and follows the OpenCL rules.
   if (const auto *A = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
-    bool CheckFirstArgument = S.getLangOpts().OpenCL
-              ? getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getZDimVal()
-	      : getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getXDimVal();
+    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();
+        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)) {
+    if (CheckFirstArgument || CheckSecondArgument || CheckThirdArgument) {
       S.Diag(AL.getLoc(), diag::err_conflicting_sycl_function_attributes)
           << AL << A;
       S.Diag(A->getLocation(), diag::note_conflicting_attribute);
@@ -3621,18 +3623,20 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
   // __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 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);
+        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(CI.getLoc(), diag::err_conflicting_sycl_function_attributes)
+          << CI << DeclAttr;
       Diag(DeclAttr->getLoc(), diag::note_conflicting_attribute);
       return;
     }
@@ -3713,18 +3717,20 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr(
   // __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 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());
+        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;
+          << DeclAttr << &A;
       Diag(A.getLoc(), diag::note_conflicting_attribute);
       return nullptr;
     }
diff --git a/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp b/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp
index b86ce667576f2..a07324db6ced4 100644
--- a/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp
@@ -36,10 +36,11 @@ struct Func {
 #ifdef TRIGGER_ERROR
 struct DAFuncObj {
   [[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 {}
+  [[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
diff --git a/clang/test/SemaSYCL/intel-max-work-group-size.cpp b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
index abc2c2ec0e6e6..6bbee30f9e983 100644
--- a/clang/test/SemaSYCL/intel-max-work-group-size.cpp
+++ b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
@@ -65,20 +65,19 @@ void instantiate() {
   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() {}
+[[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
+[[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
@@ -88,23 +87,25 @@ void f9() {}
 [[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 f13() {} // expected-error {{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}
+[[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(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
+[[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 f18();  // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
+[[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
diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp
index 437b3c86ed1c0..af8730ea98700 100644
--- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp
+++ b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp
@@ -20,7 +20,7 @@ func1();
 
 #else
 //second case - expect error
-[[intel::max_work_group_size(4, 4, 4)]] void func2(); // expected-note {{conflicting attribute is here}}
+[[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
diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp
index 1b28b759b110f..8bc681b5b690d 100644
--- a/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp
+++ b/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp
@@ -29,7 +29,8 @@ class Functor32 {
   // 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 {}
+  [[sycl::reqd_work_group_size(1, 1, 32)]] void
+  operator()() const {}
 };
 #endif // TRIGGER_ERROR
 
diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp
index 86634a00aad9a..a914cd661d0b0 100644
--- a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp
+++ b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp
@@ -52,7 +52,8 @@ class Functor32 {
   // 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 {}
+  [[sycl::reqd_work_group_size(1, 1, 32)]] void
+  operator()() const {}
 };
 #endif
 class Functor16x16x16 {

From f9406eefa31f8fa2c09397dd4bffc7f1375ed979 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Wed, 2 Mar 2022 08:51:06 -0800
Subject: [PATCH 18/18] Fix Format errors

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp                     | 2 +-
 clang/test/SemaSYCL/intel-max-work-group-size.cpp   | 4 ++--
 clang/test/SemaSYCL/reqd-work-group-size-device.cpp | 2 +-
 3 files changed, 4 insertions(+), 4 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 20dce226a9f3d..2cb051e33bc56 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3635,7 +3635,7 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
             : checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), ZDim);
 
     if (CheckFirstArgument || CheckSecondArgument || CheckThirdArgument) {
-       Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes)
+      Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes)
           << CI << DeclAttr;
       Diag(DeclAttr->getLoc(), diag::note_conflicting_attribute);
       return;
diff --git a/clang/test/SemaSYCL/intel-max-work-group-size.cpp b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
index 6bbee30f9e983..524884b00675c 100644
--- a/clang/test/SemaSYCL/intel-max-work-group-size.cpp
+++ b/clang/test/SemaSYCL/intel-max-work-group-size.cpp
@@ -86,7 +86,7 @@ f9() {}
 [[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}}
+[[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() {}
 
@@ -103,7 +103,7 @@ f15() {} // OK
 
 [[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}}
+[[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();
 
diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp
index a914cd661d0b0..634257a4aebe1 100644
--- a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp
+++ b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp
@@ -51,7 +51,7 @@ class Functor32 {
   // expected-note@+3{{conflicting attribute is here}}
   // 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(32, 1, 1)]] // expected-note {{conflicting attribute is here}}
   [[sycl::reqd_work_group_size(1, 1, 32)]] void
   operator()() const {}
 };