From 3614707bf42d9cdd86d2a8876b5e375c97c5ab96 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Mon, 8 Mar 2021 20:29:42 -0800
Subject: [PATCH 01/29] [SYCL][FPGA] Refactor of
 [[intel::max_global_work_dim()]] attribute

This patch

1. refactors FPGA function attribute using https://github.com/intel/llvm/pull/3224:
   [[intel::max_global_work_dim()]] attribute to better fit for community standards.
2. refactors the way we handle duplicate attributes and
   mutually exclusive attributes logic when present on a given declaration.
3. handles redeclarations or template instantiations properly.
4. adds tests
5. adds new test cases where the value of 'max_global_work_dim' attribute
   equals to 0, we shall ensure that if max_work_group_size and
   reqd_work_group_size attributes exist, they hold equal values (1, 1, 1).

   Before the new refactoring patch, we silently accepeted this test case:

   struct TRIFuncObjBad {
     [[intel::max_work_group_size(8, 8, 8)]]
     [[cl::reqd_work_group_size(4, 4, 4)]]
     [[intel::max_global_work_dim(0)]]
      void operator()() const {}
   };

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/include/clang/Sema/Sema.h               |  18 +-
 clang/lib/Sema/SemaDecl.cpp                   |   2 +
 clang/lib/Sema/SemaDeclAttr.cpp               | 106 +++++++-
 .../lib/Sema/SemaTemplateInstantiateDecl.cpp  |  12 +-
 .../intel-max-global-work-dim-device.cpp      | 237 +++++++++++++++++-
 ...ice-intel-max-global-work-dim-template.cpp |   9 +
 6 files changed, 351 insertions(+), 33 deletions(-)

diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index c4c4a8dae34a6..299c92e1bfbd8 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -10234,6 +10234,10 @@ class Sema final {
                                 Expr *E);
   SYCLIntelLoopFuseAttr *
   MergeSYCLIntelLoopFuseAttr(Decl *D, const SYCLIntelLoopFuseAttr &A);
+  void AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D, const AttributeCommonInfo &CI,
+                                        Expr *E);
+  SYCLIntelMaxGlobalWorkDimAttr *MergeSYCLIntelMaxGlobalWorkDimAttr(
+      Decl *D, const SYCLIntelMaxGlobalWorkDimAttr &A);
 
   /// AddAlignedAttr - Adds an aligned attribute to a particular declaration.
   void AddAlignedAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E,
@@ -13100,20 +13104,6 @@ void Sema::addIntelSingleArgAttr(Decl *D, const AttributeCommonInfo &CI,
         return;
       }
     }
-    if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim) {
-      if (ArgInt < 0) {
-        Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer)
-            << CI << /*non-negative*/ 1;
-        return;
-      }
-    }
-    if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim) {
-      if (ArgInt > 3) {
-        Diag(E->getBeginLoc(), diag::err_attribute_argument_out_of_range)
-            << CI << 0 << 3 << E->getSourceRange();
-        return;
-      }
-    }
     if (CI.getParsedKind() == ParsedAttr::AT_IntelFPGAPrivateCopies) {
       if (ArgInt < 0) {
         Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer)
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 6c7631f655810..18f1628efe3fe 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -2626,6 +2626,8 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D,
     NewAttr = S.MergeSYCLIntelSchedulerTargetFmaxMhzAttr(D, *A);
   else if (const auto *A = dyn_cast<SYCLIntelNoGlobalWorkOffsetAttr>(Attr))
     NewAttr = S.MergeSYCLIntelNoGlobalWorkOffsetAttr(D, *A);
+  else if (const auto *A = dyn_cast<SYCLIntelMaxGlobalWorkDimAttr>(Attr))
+    NewAttr = S.MergeSYCLIntelMaxGlobalWorkDimAttr(D, *A);
   else if (Attr->shouldInheritEvenIfAlreadyPresent() || !DeclHasAttr(D, Attr))
     NewAttr = cast<InheritableAttr>(Attr->clone(S.Context));
 
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index b67e650afa0be..e341bcb0a4d41 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3430,23 +3430,103 @@ static void handleSYCLIntelSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D,
 }
 
 // Handles max_global_work_dim.
-static void handleMaxGlobalWorkDimAttr(Sema &S, Decl *D, const ParsedAttr &A) {
-  if (D->isInvalidDecl())
-    return;
-
-  Expr *E = A.getArgAsExpr(0);
+void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
+                                            const AttributeCommonInfo &CI,
+                                            Expr *E) {
+  if (!E->isValueDependent()) {
+    // Validate that we have an integer constant expression and then store the
+    // converted constant expression into the semantic attribute so that we
+    // don't have to evaluate it again later.
+    llvm::APSInt ArgVal;
+    ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal);
+    if (Res.isInvalid())
+      return;
+    E = Res.get();
 
-  if (!checkWorkGroupSizeValues(S, D, A)) {
-    D->setInvalidDecl();
-    return;
+    // This attribute requires a non-negative value.
+    if (ArgVal < 0) {
+      Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer)
+          << CI << /*non-negative*/ 1;
+      return;
+    }
+    // This attribute must be in the range [0, 3].
+    if (ArgVal > 3) {
+      Diag(E->getBeginLoc(), diag::err_attribute_argument_out_of_range)
+          << CI << 0 << 3 << E->getSourceRange();
+      return;
+    }
+    // Check to see if there's a duplicate attribute with different values
+    // already applied to the declaration.
+    if (const auto *DeclAttr =
+            D->getAttr<SYCLIntelMaxGlobalWorkDimAttr>()) {
+      // If the other attribute argument is instantiation dependent, we won't
+      // have converted it to a constant expression yet and thus we test
+      // whether this is a null pointer.
+      const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue());
+      if (DeclExpr && ArgVal != DeclExpr->getResultAsAPSInt()) {
+        Diag(CI.getLoc(), diag::warn_duplicate_attribute) << CI;
+        Diag(DeclAttr->getLoc(), diag::note_previous_attribute);
+        return;
+      }
+    }
+    // If the declaration has an [[intel::reqd_work_group_size()]] or
+    // [[cl::reqd_work_group_size()]] attribute, check to see if they
+    // hold equal values (1, 1, 1) in case the value of
+    // [[intel::max_global_work_dim()]] attribute equals to 0.
+    if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
+      Optional<llvm::APSInt> XDimVal = DeclAttr->getXDimVal(Context);
+      Optional<llvm::APSInt> YDimVal = DeclAttr->getYDimVal(Context);
+      Optional<llvm::APSInt> ZDimVal = DeclAttr->getZDimVal(Context);
+      if ((ArgVal == 0) && (*XDimVal != 1 || *YDimVal != 1 || *ZDimVal != 1)) {
+        Diag(DeclAttr->getLocation(), diag::err_sycl_x_y_z_arguments_must_be_one)
+            << DeclAttr << CI;
+        return;
+      }
+    }
+    // If the declaration has an [[intel::max_work_group_size()]]
+    // attribute, check to see if it holds equal values (1, 1, 1) in
+    // case the value of [[intel::max_global_work_dim()]] attribute equals to 0.
+    if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
+      Optional<llvm::APSInt> XDimVal = DeclAttr->getXDimVal(Context);
+      Optional<llvm::APSInt> YDimVal = DeclAttr->getYDimVal(Context);
+      Optional<llvm::APSInt> ZDimVal = DeclAttr->getZDimVal(Context);
+      if ((ArgVal == 0) && (*XDimVal != 1 || *YDimVal != 1 || *ZDimVal != 1)) {
+        Diag(DeclAttr->getLocation(), diag::err_sycl_x_y_z_arguments_must_be_one)
+            << DeclAttr << CI;
+        return;
+      }
+    }
   }
 
-  if (D->getAttr<SYCLIntelMaxGlobalWorkDimAttr>())
-    S.Diag(A.getLoc(), diag::warn_duplicate_attribute) << A;
+  D->addAttr(::new (Context)
+                 SYCLIntelMaxGlobalWorkDimAttr(Context, CI, E));
+}
 
-  S.CheckDeprecatedSYCLAttributeSpelling(A);
+SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr(
+    Decl *D, const SYCLIntelMaxGlobalWorkDimAttr &A) {
+  // Check to see if there's a duplicate attribute with different values
+  // already applied to the declaration.
+  if (const auto *DeclAttr =
+          D->getAttr<SYCLIntelMaxGlobalWorkDimAttr>()) {
+    const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue());
+    const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getValue());
+    if (DeclExpr && MergeExpr &&
+        DeclExpr->getResultAsAPSInt() != MergeExpr->getResultAsAPSInt()) {
+      Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A;
+      Diag(A.getLoc(), diag::note_previous_attribute);
+      return nullptr;
+    }
+  }
+  return ::new (Context)
+      SYCLIntelMaxGlobalWorkDimAttr(Context, A, A.getValue());
+}
 
-  S.addIntelSingleArgAttr<SYCLIntelMaxGlobalWorkDimAttr>(D, A, E);
+static void handleSYCLIntelMaxGlobalWorkDimAttr(Sema &S, Decl *D,
+                                                const ParsedAttr &AL) {
+  S.CheckDeprecatedSYCLAttributeSpelling(AL);
+
+  Expr *E = AL.getArgAsExpr(0);
+  S.AddSYCLIntelMaxGlobalWorkDimAttr(D, AL, E);
 }
 
 // Handles [[intel::loop_fuse]] and [[intel::loop_fuse_independent]].
@@ -9129,7 +9209,7 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
     handleSYCLIntelSchedulerTargetFmaxMhzAttr(S, D, AL);
     break;
   case ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim:
-    handleMaxGlobalWorkDimAttr(S, D, AL);
+    handleSYCLIntelMaxGlobalWorkDimAttr(S, D, AL);
     break;
   case ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset:
     handleSYCLIntelNoGlobalWorkOffsetAttr(S, D, AL);
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index 1be77cc032795..257862eb76100 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -683,6 +683,16 @@ static void instantiateSYCLIntelNoGlobalWorkOffsetAttr(
     S.AddSYCLIntelNoGlobalWorkOffsetAttr(New, *A, Result.getAs<Expr>());
 }
 
+static void instantiateSYCLIntelMaxGlobalWorkDimAttr(
+    Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
+    const SYCLIntelMaxGlobalWorkDimAttr *A, Decl *New) {
+  EnterExpressionEvaluationContext Unevaluated(
+      S, Sema::ExpressionEvaluationContext::ConstantEvaluated);
+  ExprResult Result = S.SubstExpr(A->getValue(), TemplateArgs);
+  if (!Result.isInvalid())
+    S.AddSYCLIntelMaxGlobalWorkDimAttr(New, *A, Result.getAs<Expr>());
+}
+
 template <typename AttrName>
 static void instantiateIntelSYCLFunctionAttr(
     Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
@@ -892,7 +902,7 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
     }
     if (const auto *SYCLIntelMaxGlobalWorkDim =
             dyn_cast<SYCLIntelMaxGlobalWorkDimAttr>(TmplAttr)) {
-      instantiateIntelSYCLFunctionAttr<SYCLIntelMaxGlobalWorkDimAttr>(
+      instantiateSYCLIntelMaxGlobalWorkDimAttr(
           *this, TemplateArgs, SYCLIntelMaxGlobalWorkDim, New);
       continue;
     }
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 48e9305d6ecbf..a8b554d6320d5 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
@@ -33,6 +33,20 @@ struct Func {
   [[intelfpga::max_global_work_dim(2)]] void operator()() const {}
 };
 
+//Checking of duplicate argument values
+[[intel::max_global_work_dim(2)]] void bar();
+[[intel::max_global_work_dim(2)]] void bar() {} // OK
+
+//Checking of different argument values
+[[intel::max_global_work_dim(2)]] void baz();  // expected-note {{previous attribute is here}}
+[[intel::max_global_work_dim(1)]] void baz();  // expected-warning {{attribute 'max_global_work_dim' is already applied with different arguments}}
+
+// Checks correctness of mutual usage of different work_group_size attributes:
+// reqd_work_group_size, max_work_group_size and max_global_work_dim.
+// In case the value of 'max_global_work_dim' attribute equals to 0 we shall
+// ensure that if max_work_group_size and reqd_work_group_size attributes exist,
+// they hold equal values (1, 1, 1).
+
 struct TRIFuncObjGood1 {
   [[intel::max_global_work_dim(0)]]
   [[intel::max_work_group_size(1, 1, 1)]]
@@ -47,7 +61,49 @@ struct TRIFuncObjGood2 {
   operator()() const {}
 };
 
+struct TRIFuncObjGood3 {
+  [[intel::reqd_work_group_size(1, 1, 1)]]
+  [[intel::max_global_work_dim(0)]] void
+  operator()() const {}
+};
+
+struct TRIFuncObjGood4 {
+  [[cl::reqd_work_group_size(1, 1, 1)]]
+  [[intel::max_global_work_dim(0)]] void
+  operator()() const {}
+};
+
+struct TRIFuncObjGood5 {
+  [[intel::max_work_group_size(1, 1, 1)]]
+  [[intel::max_global_work_dim(0)]] void
+  operator()() const {}
+};
+
+struct TRIFuncObjGood6 {
+  [[intel::reqd_work_group_size(4, 1, 1)]]
+  [[intel::max_global_work_dim(3)]] void
+  operator()() const {}
+};
+
+struct TRIFuncObjGood7 {
+  [[cl::reqd_work_group_size(4, 1, 1)]]
+  [[intel::max_global_work_dim(3)]] void
+  operator()() const {}
+};
+
+struct TRIFuncObjGood8 {
+  [[intel::max_work_group_size(8, 1, 1)]]
+  [[intel::max_global_work_dim(3)]] void
+  operator()() const {}
+};
+
 #ifdef TRIGGER_ERROR
+// Checks correctness of mutual usage of different work_group_size attributes:
+// reqd_work_group_size, max_work_group_size and max_global_work_dim.
+// In case the value of 'max_global_work_dim' attribute equals to 0 we shall
+// ensure that if max_work_group_size and reqd_work_group_size attributes exist,
+// they hold equal values (1, 1, 1).
+
 struct TRIFuncObjBad {
   [[intel::max_global_work_dim(0)]]
   [[intel::max_work_group_size(8, 8, 8)]] // expected-error{{'max_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
@@ -55,6 +111,59 @@ struct TRIFuncObjBad {
   void
   operator()() const {}
 };
+
+struct TRIFuncObjBad1 {
+  [[intel::max_work_group_size(8, 8, 8)]] // expected-error{{'max_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+  [[intel::max_global_work_dim(0)]] void
+  operator()() const {}
+};
+
+struct TRIFuncObjBad2 {
+  [[cl::reqd_work_group_size(4, 4, 4)]]   // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+  [[intel::max_global_work_dim(0)]] void
+  operator()() const {}
+};
+
+struct TRIFuncObjBad3 {
+  [[intel::reqd_work_group_size(4, 4, 4)]]   // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+  [[intel::max_global_work_dim(0)]] void
+  operator()() const {}
+};
+
+// Tests for incorrect argument values for Intel FPGA function attributes:
+// reqd_work_group_size, max_work_group_size and max_global_work_dim.
+
+struct TRIFuncObjBad4 {
+  // expected-error@+2{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+  // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}}
+  [[intel::reqd_work_group_size(4, 4, -4)]]
+  [[intel::max_global_work_dim(0)]] void
+  operator()() const {}
+};
+
+struct TRIFuncObjBad5 {
+  [[intel::max_work_group_size(4, 4, 4.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}}
+  [[intel::max_global_work_dim(0)]] void
+  operator()() const {}
+};
+
+struct TRIFuncObjBad6 {
+  [[cl::reqd_work_group_size(0, 4, 4)]] // expected-error{{'reqd_work_group_size' attribute must be greater than 0}}
+  [[intel::max_global_work_dim(0)]] void
+  operator()() const {}
+};
+
+struct TRIFuncObjBad7 {
+  [[intel::reqd_work_group_size(4, 4, 4)]]
+  [[intel::max_global_work_dim(-2)]] // expected-error{{'max_global_work_dim' attribute requires a non-negative integral compile time constant expression}}
+  void operator()() const {}
+};
+
+struct TRIFuncObjBad8 {
+  [[intel::max_work_group_size(4, 4, 4)]]
+  [[intel::max_global_work_dim(4.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}}
+  void operator()() const {}
+};
 #endif // TRIGGER_ERROR
 
 int main() {
@@ -137,19 +246,137 @@ int main() {
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 3
     // CHECK-NEXT:  IntegerLiteral{{.*}}3{{$}}
+
+    h.single_task<class test_kernel5>(TRIFuncObjGood3());
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel5
+    // CHECK:       ReqdWorkGroupSizeAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 0
+    // CHECK-NEXT:  IntegerLiteral{{.*}}0{{$}}
+
+    h.single_task<class test_kernel6>(TRIFuncObjGood4());
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel6
+    // CHECK:       ReqdWorkGroupSizeAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 0
+    // CHECK-NEXT:  IntegerLiteral{{.*}}0{{$}}
+
+    h.single_task<class test_kernel7>(TRIFuncObjGood5());
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel7
+    // CHECK:       SYCLIntelMaxWorkGroupSizeAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 0
+    // CHECK-NEXT:  IntegerLiteral{{.*}}0{{$}}
+
+    h.single_task<class test_kernel8>(TRIFuncObjGood6());
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel8
+    // CHECK:       ReqdWorkGroupSizeAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 3
+    // CHECK-NEXT:  IntegerLiteral{{.*}}3{{$}}
+
+    h.single_task<class test_kernel9>(TRIFuncObjGood7());
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel9
+    // CHECK:       ReqdWorkGroupSizeAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 3
+    // CHECK-NEXT:  IntegerLiteral{{.*}}3{{$}}
+
+    h.single_task<class test_kernel10>(TRIFuncObjGood8());
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel10
+    // CHECK:       SYCLIntelMaxWorkGroupSizeAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 8
+    // CHECK-NEXT:  IntegerLiteral{{.*}}8{{$}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 1
+    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 3
+    // CHECK-NEXT:  IntegerLiteral{{.*}}3{{$}}
 #ifdef TRIGGER_ERROR
     [[intel::max_global_work_dim(1)]] int Var = 0; // expected-error{{'max_global_work_dim' attribute only applies to functions}}
 
-    h.single_task<class test_kernel6>(
+    h.single_task<class test_kernel11>(
         []() [[intel::max_global_work_dim(-8)]]{}); // expected-error{{'max_global_work_dim' attribute requires a non-negative integral compile time constant expression}}
 
-    h.single_task<class test_kernel7>(
-        []() [[intel::max_global_work_dim(3),
+    h.single_task<class test_kernel2>(
+        []() [[intel::max_global_work_dim(3),      // expected-note {{previous attribute is here}}
                intel::max_global_work_dim(2)]]{}); // expected-warning{{attribute 'max_global_work_dim' is already applied with different arguments}}
 
-    h.single_task<class test_kernel8>(TRIFuncObjBad());
+    h.single_task<class test_kernel3>(TRIFuncObjBad());
+
+    h.single_task<class test_kernel4>(TRIFuncObjBad1());
+
+    h.single_task<class test_kernel5>(TRIFuncObjBad2());
+
+    h.single_task<class test_kernel6>(TRIFuncObjBad3());
+
+    h.single_task<class test_kernel7>(TRIFuncObjBad4());
+
+    h.single_task<class test_kernel8>(TRIFuncObjBad5());
+
+    h.single_task<class test_kernel9>(TRIFuncObjBad6());
+
+    h.single_task<class test_kerne20>(TRIFuncObjBad7());
+
+    h.single_task<class test_kerne21>(TRIFuncObjBad8());
 
-    h.single_task<class test_kernel9>(
+    h.single_task<class test_kernel22>(
         []() [[intel::max_global_work_dim(4)]]{}); // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}}
 #endif // TRIGGER_ERROR
   });
diff --git a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
index 87e45364994f1..9ef4defc4ee42 100644
--- a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
+++ b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
@@ -61,11 +61,20 @@ template <int N>
 // expected-error@+1{{'max_global_work_dim' attribute requires a non-negative integral compile time constant expression}}
 [[intel::max_global_work_dim(N)]] void func3() {}
 
+// Test that checks template instantiations for different argument values.
+template <int size>
+[[intel::max_global_work_dim(1)]] void func4(); // expected-note {{previous attribute is here}}
+
+template <int size>
+[[intel::max_global_work_dim(size)]] void func4() {} // expected-warning {{attribute 'max_global_work_dim' is already applied with different arguments}}
+
 int check() {
   // no error expected
   func3<3>();
   //expected-note@+1{{in instantiation of function template specialization 'func3<-1>' requested here}}
   func3<-1>();
+  //expected-note@+1 {{in instantiation of function template specialization 'func4<2>' requested here}}
+  func4<2>();
   return 0;
 }
 

From bda9eef27661f580273d0415ec83013cfa19ce88 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Mon, 8 Mar 2021 20:49:37 -0800
Subject: [PATCH 02/29] Fix errors

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/include/clang/Sema/Sema.h                |  5 +++--
 clang/lib/Sema/SemaDeclAttr.cpp                | 15 +++++++--------
 clang/lib/Sema/SemaTemplateInstantiateDecl.cpp |  4 ++--
 3 files changed, 12 insertions(+), 12 deletions(-)

diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 299c92e1bfbd8..fe3f74bbbf117 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -10236,8 +10236,9 @@ class Sema final {
   MergeSYCLIntelLoopFuseAttr(Decl *D, const SYCLIntelLoopFuseAttr &A);
   void AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D, const AttributeCommonInfo &CI,
                                         Expr *E);
-  SYCLIntelMaxGlobalWorkDimAttr *MergeSYCLIntelMaxGlobalWorkDimAttr(
-      Decl *D, const SYCLIntelMaxGlobalWorkDimAttr &A);
+  SYCLIntelMaxGlobalWorkDimAttr *
+  MergeSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
+                                     const SYCLIntelMaxGlobalWorkDimAttr &A);
 
   /// AddAlignedAttr - Adds an aligned attribute to a particular declaration.
   void AddAlignedAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E,
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index e341bcb0a4d41..030fcbba84c79 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3457,8 +3457,7 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
     }
     // Check to see if there's a duplicate attribute with different values
     // already applied to the declaration.
-    if (const auto *DeclAttr =
-            D->getAttr<SYCLIntelMaxGlobalWorkDimAttr>()) {
+    if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxGlobalWorkDimAttr>()) {
       // If the other attribute argument is instantiation dependent, we won't
       // have converted it to a constant expression yet and thus we test
       // whether this is a null pointer.
@@ -3478,7 +3477,8 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
       Optional<llvm::APSInt> YDimVal = DeclAttr->getYDimVal(Context);
       Optional<llvm::APSInt> ZDimVal = DeclAttr->getZDimVal(Context);
       if ((ArgVal == 0) && (*XDimVal != 1 || *YDimVal != 1 || *ZDimVal != 1)) {
-        Diag(DeclAttr->getLocation(), diag::err_sycl_x_y_z_arguments_must_be_one)
+        Diag(DeclAttr->getLocation(),
+             diag::err_sycl_x_y_z_arguments_must_be_one)
             << DeclAttr << CI;
         return;
       }
@@ -3491,23 +3491,22 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
       Optional<llvm::APSInt> YDimVal = DeclAttr->getYDimVal(Context);
       Optional<llvm::APSInt> ZDimVal = DeclAttr->getZDimVal(Context);
       if ((ArgVal == 0) && (*XDimVal != 1 || *YDimVal != 1 || *ZDimVal != 1)) {
-        Diag(DeclAttr->getLocation(), diag::err_sycl_x_y_z_arguments_must_be_one)
+	Diag(DeclAttr->getLocation(),
+             diag::err_sycl_x_y_z_arguments_must_be_one)
             << DeclAttr << CI;
         return;
       }
     }
   }
 
-  D->addAttr(::new (Context)
-                 SYCLIntelMaxGlobalWorkDimAttr(Context, CI, E));
+  D->addAttr(::new (Context) SYCLIntelMaxGlobalWorkDimAttr(Context, CI, E));
 }
 
 SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr(
     Decl *D, const SYCLIntelMaxGlobalWorkDimAttr &A) {
   // Check to see if there's a duplicate attribute with different values
   // already applied to the declaration.
-  if (const auto *DeclAttr =
-          D->getAttr<SYCLIntelMaxGlobalWorkDimAttr>()) {
+  if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxGlobalWorkDimAttr>()) {
     const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue());
     const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getValue());
     if (DeclExpr && MergeExpr &&
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index 257862eb76100..8769a4ae566ab 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -902,8 +902,8 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
     }
     if (const auto *SYCLIntelMaxGlobalWorkDim =
             dyn_cast<SYCLIntelMaxGlobalWorkDimAttr>(TmplAttr)) {
-      instantiateSYCLIntelMaxGlobalWorkDimAttr(
-          *this, TemplateArgs, SYCLIntelMaxGlobalWorkDim, New);
+      instantiateSYCLIntelMaxGlobalWorkDimAttr(*this, TemplateArgs,
+                                               SYCLIntelMaxGlobalWorkDim, New);
       continue;
     }
     if (const auto *SYCLIntelLoopFuse =

From a56f4491dfe727840a65c45e27685598d1c27987 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Mon, 8 Mar 2021 20:59:42 -0800
Subject: [PATCH 03/29] Fix errors

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 030fcbba84c79..d8b79192234d8 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3491,7 +3491,7 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
       Optional<llvm::APSInt> YDimVal = DeclAttr->getYDimVal(Context);
       Optional<llvm::APSInt> ZDimVal = DeclAttr->getZDimVal(Context);
       if ((ArgVal == 0) && (*XDimVal != 1 || *YDimVal != 1 || *ZDimVal != 1)) {
-	Diag(DeclAttr->getLocation(),
+        Diag(DeclAttr->getLocation(),
              diag::err_sycl_x_y_z_arguments_must_be_one)
             << DeclAttr << CI;
         return;

From 86509bfff96415eeed4dc773bf2441c7ebb69a86 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Tue, 9 Mar 2021 10:48:34 -0800
Subject: [PATCH 04/29] address review comments

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp               | 46 +++++++++----------
 .../intel-max-global-work-dim-device.cpp      | 12 ++---
 2 files changed, 27 insertions(+), 31 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index d8b79192234d8..a396c1861a846 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3468,34 +3468,30 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
         return;
       }
     }
-    // If the declaration has an [[intel::reqd_work_group_size()]] or
-    // [[cl::reqd_work_group_size()]] attribute, check to see if they
-    // hold equal values (1, 1, 1) in case the value of
-    // [[intel::max_global_work_dim()]] attribute equals to 0.
-    if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
-      Optional<llvm::APSInt> XDimVal = DeclAttr->getXDimVal(Context);
-      Optional<llvm::APSInt> YDimVal = DeclAttr->getYDimVal(Context);
-      Optional<llvm::APSInt> ZDimVal = DeclAttr->getZDimVal(Context);
-      if ((ArgVal == 0) && (*XDimVal != 1 || *YDimVal != 1 || *ZDimVal != 1)) {
-        Diag(DeclAttr->getLocation(),
-             diag::err_sycl_x_y_z_arguments_must_be_one)
-            << DeclAttr << CI;
-        return;
+    // Checks correctness of mutual usage of different work_group_size
+    // attributes: reqd_work_group_size, max_work_group_size and
+    // max_global_work_dim.
+    // In case the value of 'max_global_work_dim' attribute equals to 0
+    // we shall ensure that if max_work_group_size and reqd_work_group_size
+    // attributes exist, they hold equal values (1, 1, 1).
+    auto Check = [this, &CI](const auto *A) {
+      Optional<llvm::APSInt> XDimVal = A->getXDimVal(Context);
+      Optional<llvm::APSInt> YDimVal = A->getYDimVal(Context);
+      Optional<llvm::APSInt> ZDimVal = A->getZDimVal(Context);
+      if (*XDimVal != 1 || *YDimVal != 1 || *ZDimVal != 1) {
+        Diag(A->getLocation(), diag::err_sycl_x_y_z_arguments_must_be_one)
+            << A << CI;
+        return false;
       }
-    }
-    // If the declaration has an [[intel::max_work_group_size()]]
-    // attribute, check to see if it holds equal values (1, 1, 1) in
-    // case the value of [[intel::max_global_work_dim()]] attribute equals to 0.
+      return true;
+    };
     if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
-      Optional<llvm::APSInt> XDimVal = DeclAttr->getXDimVal(Context);
-      Optional<llvm::APSInt> YDimVal = DeclAttr->getYDimVal(Context);
-      Optional<llvm::APSInt> ZDimVal = DeclAttr->getZDimVal(Context);
-      if ((ArgVal == 0) && (*XDimVal != 1 || *YDimVal != 1 || *ZDimVal != 1)) {
-        Diag(DeclAttr->getLocation(),
-             diag::err_sycl_x_y_z_arguments_must_be_one)
-            << DeclAttr << CI;
+      if ((ArgVal == 0) && !Check(DeclAttr))
+        return;
+    }
+    if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
+      if ((ArgVal == 0) && !Check(DeclAttr))
         return;
-      }
     }
   }
 
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 a8b554d6320d5..1afbeb0556c0f 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
@@ -33,11 +33,11 @@ struct Func {
   [[intelfpga::max_global_work_dim(2)]] void operator()() const {}
 };
 
-//Checking of duplicate argument values
+// Checking of duplicate argument values.
 [[intel::max_global_work_dim(2)]] void bar();
 [[intel::max_global_work_dim(2)]] void bar() {} // OK
 
-//Checking of different argument values
+// Checking of different argument values.
 [[intel::max_global_work_dim(2)]] void baz();  // expected-note {{previous attribute is here}}
 [[intel::max_global_work_dim(1)]] void baz();  // expected-warning {{attribute 'max_global_work_dim' is already applied with different arguments}}
 
@@ -62,7 +62,7 @@ struct TRIFuncObjGood2 {
 };
 
 struct TRIFuncObjGood3 {
-  [[intel::reqd_work_group_size(1, 1, 1)]]
+  [[intel::reqd_work_group_size(1)]]
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
@@ -125,7 +125,7 @@ struct TRIFuncObjBad2 {
 };
 
 struct TRIFuncObjBad3 {
-  [[intel::reqd_work_group_size(4, 4, 4)]]   // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+  [[intel::reqd_work_group_size(4)]]   // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
@@ -136,7 +136,7 @@ struct TRIFuncObjBad3 {
 struct TRIFuncObjBad4 {
   // expected-error@+2{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
   // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}}
-  [[intel::reqd_work_group_size(4, 4, -4)]]
+  [[intel::reqd_work_group_size(-4, 1)]]
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
@@ -154,7 +154,7 @@ struct TRIFuncObjBad6 {
 };
 
 struct TRIFuncObjBad7 {
-  [[intel::reqd_work_group_size(4, 4, 4)]]
+  [[intel::reqd_work_group_size(4)]]
   [[intel::max_global_work_dim(-2)]] // expected-error{{'max_global_work_dim' attribute requires a non-negative integral compile time constant expression}}
   void operator()() const {}
 };

From 165f537475e24ad932f4b7d5a1d247c46f039fcf Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Tue, 9 Mar 2021 11:09:59 -0800
Subject: [PATCH 05/29] Update comments

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

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index a396c1861a846..bf5d82724a4c6 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3468,12 +3468,10 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
         return;
       }
     }
-    // Checks correctness of mutual usage of different work_group_size
-    // attributes: reqd_work_group_size, max_work_group_size and
-    // max_global_work_dim.
-    // In case the value of 'max_global_work_dim' attribute equals to 0
-    // we shall ensure that if max_work_group_size and reqd_work_group_size
-    // attributes exist, they hold equal values (1, 1, 1).
+    // If the declaration has a SYCLIntelMaxWorkGroupSizeAttr or
+    // ReqdWorkGroupSizeAttr, check to see if they hold equal values
+    // (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr
+    // equals to 0.
     auto Check = [this, &CI](const auto *A) {
       Optional<llvm::APSInt> XDimVal = A->getXDimVal(Context);
       Optional<llvm::APSInt> YDimVal = A->getYDimVal(Context);

From 1f82fa3c639433b7e542304ac3c6f645f2aebfe1 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Tue, 9 Mar 2021 12:03:32 -0800
Subject: [PATCH 06/29] remove paren

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

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index bf5d82724a4c6..061ec167a4993 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3484,11 +3484,11 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
       return true;
     };
     if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
-      if ((ArgVal == 0) && !Check(DeclAttr))
+      if (ArgVal == 0 && !Check(DeclAttr))
         return;
     }
     if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
-      if ((ArgVal == 0) && !Check(DeclAttr))
+      if (ArgVal == 0 && !Check(DeclAttr))
         return;
     }
   }

From 9e6a71e930c4506f736df19e4ea72f6ded2b3a95 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Wed, 10 Mar 2021 10:00:26 -0800
Subject: [PATCH 07/29] Remove generic template function
 addIntelSingleArgAttr(), remove -fsycl from test and redundant codes

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/include/clang/Sema/Sema.h                 |  2 --
 clang/lib/Sema/SemaDeclAttr.cpp                 | 15 ---------------
 clang/lib/Sema/SemaTemplateInstantiateDecl.cpp  | 11 -----------
 clang/test/SemaSYCL/intel-fpga-global-const.cpp |  2 +-
 4 files changed, 1 insertion(+), 29 deletions(-)

diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 30109ee8063de..f510ac6e23504 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -10207,8 +10207,6 @@ class Sema final {
   void AddIntelFPGABankBitsAttr(Decl *D, const AttributeCommonInfo &CI,
                                 Expr **Exprs, unsigned Size);
   template <typename AttrType>
-  void addIntelSingleArgAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E);
-  template <typename AttrType>
   void addIntelTripleArgAttr(Decl *D, const AttributeCommonInfo &CI,
                              Expr *XDimExpr, Expr *YDimExpr, Expr *ZDimExpr);
   void AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI,
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index ddf6e4563007f..23d80ecc04481 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3041,21 +3041,6 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) {
 
   ASTContext &Ctx = S.getASTContext();
 
-  if (AL.getKind() == ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim) {
-    ArrayRef<const Expr *> Dims;
-    Attr *B = nullptr;
-    if (const auto *B = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>())
-      Dims = B->dimensions();
-    else if (const auto *B = D->getAttr<ReqdWorkGroupSizeAttr>())
-      Dims = B->dimensions();
-    if (B) {
-      Result &=
-          checkZeroDim(B, getExprValue(Dims[0], Ctx),
-                       getExprValue(Dims[1], Ctx), getExprValue(Dims[2], Ctx));
-    }
-    return Result;
-  }
-
   if (AL.getKind() == ParsedAttr::AT_SYCLIntelMaxWorkGroupSize)
     S.CheckDeprecatedSYCLAttributeSpelling(AL);
 
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index 3762cedecf910..52ecadbc0231d 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -693,17 +693,6 @@ static void instantiateSYCLIntelMaxGlobalWorkDimAttr(
     S.AddSYCLIntelMaxGlobalWorkDimAttr(New, *A, Result.getAs<Expr>());
 }
 
-template <typename AttrName>
-static void instantiateIntelSYCLFunctionAttr(
-    Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
-    const AttrName *Attr, Decl *New) {
-  EnterExpressionEvaluationContext Unevaluated(
-      S, Sema::ExpressionEvaluationContext::ConstantEvaluated);
-  ExprResult Result = S.SubstExpr(Attr->getValue(), TemplateArgs);
-  if (!Result.isInvalid())
-    S.addIntelSingleArgAttr<AttrName>(New, *Attr, Result.getAs<Expr>());
-}
-
 static void instantiateIntelFPGAPrivateCopiesAttr(
     Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
     const IntelFPGAPrivateCopiesAttr *A, Decl *New) {
diff --git a/clang/test/SemaSYCL/intel-fpga-global-const.cpp b/clang/test/SemaSYCL/intel-fpga-global-const.cpp
index 90b5615b70570..e3b4f5687d5bd 100644
--- a/clang/test/SemaSYCL/intel-fpga-global-const.cpp
+++ b/clang/test/SemaSYCL/intel-fpga-global-const.cpp
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s
 
 // Test that checks global constant variable (which allows the redeclaration) since
 // IntelFPGAConstVar is one of the subjects listed for [[intel::max_replicates()]] attribute.

From a17f9cdcf4bc2eb7dccc74e328e2ad685a167a62 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Wed, 10 Mar 2021 18:50:18 -0800
Subject: [PATCH 08/29] 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      |  4 +--
 2 files changed, 18 insertions(+), 14 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 23d80ecc04481..7cd0194f1fc89 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3447,10 +3447,12 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
       // If the other attribute argument is instantiation dependent, we won't
       // have converted it to a constant expression yet and thus we test
       // whether this is a null pointer.
-      const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue());
-      if (DeclExpr && ArgVal != DeclExpr->getResultAsAPSInt()) {
-        Diag(CI.getLoc(), diag::warn_duplicate_attribute) << CI;
-        Diag(DeclAttr->getLoc(), diag::note_previous_attribute);
+      if (const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue())) {
+        if (ArgVal != DeclExpr->getResultAsAPSInt()) {
+	  Diag(CI.getLoc(), diag::warn_duplicate_attribute) << CI;
+          Diag(DeclAttr->getLoc(), diag::note_previous_attribute);
+        }
+        // If there is no mismatch, silently ignore duplicate attribute.
         return;
       }
     }
@@ -3478,7 +3480,6 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
         return;
     }
   }
-
   D->addAttr(::new (Context) SYCLIntelMaxGlobalWorkDimAttr(Context, CI, E));
 }
 
@@ -3487,15 +3488,18 @@ SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr(
   // Check to see if there's a duplicate attribute with different values
   // already applied to the declaration.
   if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxGlobalWorkDimAttr>()) {
-    const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue());
-    const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getValue());
-    if (DeclExpr && MergeExpr &&
-        DeclExpr->getResultAsAPSInt() != MergeExpr->getResultAsAPSInt()) {
-      Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A;
-      Diag(A.getLoc(), diag::note_previous_attribute);
-      return nullptr;
+    if (const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue())) {
+      if (const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getValue())) {
+	if (DeclExpr->getResultAsAPSInt() != MergeExpr->getResultAsAPSInt()) {
+          Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A;
+          Diag(A.getLoc(), diag::note_previous_attribute);
+        }
+        // If there is no mismatch, silently ignore duplicate attribute.
+        return nullptr;
+      }
     }
   }
+
   return ::new (Context)
       SYCLIntelMaxGlobalWorkDimAttr(Context, A, A.getValue());
 }
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 288c523b08b24..f92f3087cf415 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
@@ -33,9 +33,9 @@ struct Func {
   [[intelfpga::max_global_work_dim(2)]] void operator()() const {}
 };
 
-// Checking of duplicate argument values.
+// No diagnostic is thrown since arguments match. Duplicate attribute is silently ignored.
 [[intel::max_global_work_dim(2)]] void bar();
-[[intel::max_global_work_dim(2)]] void bar() {} // OK
+[[intel::max_global_work_dim(2)]] void bar() {}
 
 // Checking of different argument values.
 [[intel::max_global_work_dim(2)]] void baz();  // expected-note {{previous attribute is here}}

From 586a2d9e66277cb86aa01e9147cd761ac8898938 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Wed, 10 Mar 2021 18:58:49 -0800
Subject: [PATCH 09/29] Fix errors

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

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 7cd0194f1fc89..7f3bb97e4c0f6 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3449,7 +3449,7 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
       // whether this is a null pointer.
       if (const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue())) {
         if (ArgVal != DeclExpr->getResultAsAPSInt()) {
-	  Diag(CI.getLoc(), diag::warn_duplicate_attribute) << CI;
+          Diag(CI.getLoc(), diag::warn_duplicate_attribute) << CI;
           Diag(DeclAttr->getLoc(), diag::note_previous_attribute);
         }
         // If there is no mismatch, silently ignore duplicate attribute.
@@ -3490,7 +3490,7 @@ SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr(
   if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxGlobalWorkDimAttr>()) {
     if (const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue())) {
       if (const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getValue())) {
-	if (DeclExpr->getResultAsAPSInt() != MergeExpr->getResultAsAPSInt()) {
+        if (DeclExpr->getResultAsAPSInt() != MergeExpr->getResultAsAPSInt()) {
           Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A;
           Diag(A.getLoc(), diag::note_previous_attribute);
         }

From f172eb46641ae0c68cd572c1bf4ff728edae3737 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Sat, 20 Mar 2021 13:58:33 -0700
Subject: [PATCH 10/29] don't duplicate the attribute with the same arguments

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 .../intel-max-global-work-dim-device.cpp      | 47 ++++++++++++++-----
 1 file changed, 34 insertions(+), 13 deletions(-)

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 f92f3087cf415..2e21f1f642409 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
@@ -33,7 +33,7 @@ struct Func {
   [[intelfpga::max_global_work_dim(2)]] void operator()() const {}
 };
 
-// No diagnostic is thrown since arguments match. Duplicate attribute is silently ignored.
+// No diagnostic is thrown since arguments match.
 [[intel::max_global_work_dim(2)]] void bar();
 [[intel::max_global_work_dim(2)]] void bar() {}
 
@@ -348,35 +348,56 @@ int main() {
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 3
     // CHECK-NEXT:  IntegerLiteral{{.*}}3{{$}}
+
+    // Diagnostic is thrown since arguments mismatch.
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel11 'void ()'
+    // CHECK:       SYCLIntelMaxWorkGroupSizeAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 2
+    // CHECK-NEXT:  IntegerLiteral{{.*}}3{{$}}
+    h.single_task<class test_kernel1>(
+        []() [[intel::max_global_work_dim(3),      // expected-note {{previous attribute is here}}
+               intel::max_global_work_dim(2)]]{}); // expected-warning{{attribute 'max_global_work_dim' is already applied with different arguments}}
+
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel12 'void ()'
+    // CHECK:       SYCLIntelMaxWorkGroupSizeAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 2
+    // CHECK-NEXT:  IntegerLiteral{{.*}}2{{$}}
+    // No diagnostic is thrown since arguments match. Duplicate attribute is silently ignored.
+    h.single_task<class test_kernel2>(
+        []() [[intel::max_global_work_dim(2),
+               intel::max_global_work_dim(2)]]{}); // OK
+
 #ifdef TRIGGER_ERROR
     [[intel::max_global_work_dim(1)]] int Var = 0; // expected-error{{'max_global_work_dim' attribute only applies to functions}}
 
-    h.single_task<class test_kernel11>(
+    h.single_task<class test_kernel13>(
         []() [[intel::max_global_work_dim(-8)]]{}); // expected-error{{'max_global_work_dim' attribute requires a non-negative integral compile time constant expression}}
 
-    h.single_task<class test_kernel2>(
+    h.single_task<class test_kernel4>(
         []() [[intel::max_global_work_dim(3),      // expected-note {{previous attribute is here}}
                intel::max_global_work_dim(2)]]{}); // expected-warning{{attribute 'max_global_work_dim' is already applied with different arguments}}
 
-    h.single_task<class test_kernel3>(TRIFuncObjBad());
+    h.single_task<class test_kernel5>(TRIFuncObjBad());
 
-    h.single_task<class test_kernel4>(TRIFuncObjBad1());
+    h.single_task<class test_kernel6>(TRIFuncObjBad1());
 
-    h.single_task<class test_kernel5>(TRIFuncObjBad2());
+    h.single_task<class test_kernel7>(TRIFuncObjBad2());
 
-    h.single_task<class test_kernel6>(TRIFuncObjBad3());
+    h.single_task<class test_kernel8>(TRIFuncObjBad3());
 
-    h.single_task<class test_kernel7>(TRIFuncObjBad4());
+    h.single_task<class test_kernel9>(TRIFuncObjBad4());
 
-    h.single_task<class test_kernel8>(TRIFuncObjBad5());
+    h.single_task<class test_kernel20>(TRIFuncObjBad5());
 
-    h.single_task<class test_kernel9>(TRIFuncObjBad6());
+    h.single_task<class test_kernel21>(TRIFuncObjBad6());
 
-    h.single_task<class test_kerne20>(TRIFuncObjBad7());
+    h.single_task<class test_kerne22>(TRIFuncObjBad7());
 
-    h.single_task<class test_kerne21>(TRIFuncObjBad8());
+    h.single_task<class test_kerne23>(TRIFuncObjBad8());
 
-    h.single_task<class test_kernel22>(
+    h.single_task<class test_kernel24>(
         []() [[intel::max_global_work_dim(4)]]{}); // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}}
 #endif // TRIGGER_ERROR
   });

From 93e00bab73a5f61938b92173844ccf6df3d5702c Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Sat, 20 Mar 2021 15:06:26 -0700
Subject: [PATCH 11/29] Fix lit test failure

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 .../intel-max-global-work-dim-device.cpp      | 26 +++++++++----------
 1 file changed, 13 insertions(+), 13 deletions(-)

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 1a4027888ba5a..416e0903a8856 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
@@ -351,21 +351,21 @@ int main() {
 
     // Diagnostic is thrown since arguments mismatch.
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel11 'void ()'
-    // CHECK:       SYCLIntelMaxWorkGroupSizeAttr {{.*}}
+    // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
-    // CHECK-NEXT:  value: Int 2
+    // CHECK-NEXT:  value: Int 3
     // CHECK-NEXT:  IntegerLiteral{{.*}}3{{$}}
-    h.single_task<class test_kernel1>(
+    h.single_task<class test_kernel11>(
         []() [[intel::max_global_work_dim(3),      // expected-note {{previous attribute is here}}
                intel::max_global_work_dim(2)]]{}); // expected-warning{{attribute 'max_global_work_dim' is already applied with different arguments}}
 
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel12 'void ()'
-    // CHECK:       SYCLIntelMaxWorkGroupSizeAttr {{.*}}
+    // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 2
     // CHECK-NEXT:  IntegerLiteral{{.*}}2{{$}}
     // No diagnostic is thrown since arguments match. Duplicate attribute is silently ignored.
-    h.single_task<class test_kernel2>(
+    h.single_task<class test_kernel12>(
         []() [[intel::max_global_work_dim(2),
                intel::max_global_work_dim(2)]]{}); // OK
 
@@ -375,27 +375,27 @@ int main() {
     h.single_task<class test_kernel13>(
         []() [[intel::max_global_work_dim(-8)]]{}); // expected-error{{'max_global_work_dim' attribute requires a non-negative integral compile time constant expression}}
 
-    h.single_task<class test_kernel4>(
+    h.single_task<class test_kernell4>(
         []() [[intel::max_global_work_dim(3),      // expected-note {{previous attribute is here}}
                intel::max_global_work_dim(2)]]{}); // expected-warning{{attribute 'max_global_work_dim' is already applied with different arguments}}
 
-    h.single_task<class test_kernel5>(TRIFuncObjBad());
+    h.single_task<class test_kernel15>(TRIFuncObjBad());
 
-    h.single_task<class test_kernel6>(TRIFuncObjBad1());
+    h.single_task<class test_kernel16>(TRIFuncObjBad1());
 
-    h.single_task<class test_kernel7>(TRIFuncObjBad2());
+    h.single_task<class test_kernel17>(TRIFuncObjBad2());
 
-    h.single_task<class test_kernel8>(TRIFuncObjBad3());
+    h.single_task<class test_kernel18>(TRIFuncObjBad3());
 
-    h.single_task<class test_kernel9>(TRIFuncObjBad4());
+    h.single_task<class test_kernel19>(TRIFuncObjBad4());
 
     h.single_task<class test_kernel20>(TRIFuncObjBad5());
 
     h.single_task<class test_kernel21>(TRIFuncObjBad6());
 
-    h.single_task<class test_kerne22>(TRIFuncObjBad7());
+    h.single_task<class test_kernel22>(TRIFuncObjBad7());
 
-    h.single_task<class test_kerne23>(TRIFuncObjBad8());
+    h.single_task<class test_kernel23>(TRIFuncObjBad8());
 
     h.single_task<class test_kernel24>(
         []() [[intel::max_global_work_dim(4)]]{}); // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}}

From 28333310d07399a682a50bb42623a8144f4dee22 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Sat, 20 Mar 2021 15:14:37 -0700
Subject: [PATCH 12/29] Upadte test

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 .../intel-max-global-work-dim-device.cpp      | 36 +++++++------------
 1 file changed, 13 insertions(+), 23 deletions(-)

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 416e0903a8856..5c205460e65e0 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
@@ -349,55 +349,45 @@ int main() {
     // CHECK-NEXT:  value: Int 3
     // CHECK-NEXT:  IntegerLiteral{{.*}}3{{$}}
 
-    // Diagnostic is thrown since arguments mismatch.
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel11 'void ()'
     // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
-    // CHECK-NEXT:  value: Int 3
-    // CHECK-NEXT:  IntegerLiteral{{.*}}3{{$}}
-    h.single_task<class test_kernel11>(
-        []() [[intel::max_global_work_dim(3),      // expected-note {{previous attribute is here}}
-               intel::max_global_work_dim(2)]]{}); // expected-warning{{attribute 'max_global_work_dim' is already applied with different arguments}}
-
-    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel12 'void ()'
-    // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
-    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 2
     // CHECK-NEXT:  IntegerLiteral{{.*}}2{{$}}
     // No diagnostic is thrown since arguments match. Duplicate attribute is silently ignored.
-    h.single_task<class test_kernel12>(
+    h.single_task<class test_kernel11>(
         []() [[intel::max_global_work_dim(2),
                intel::max_global_work_dim(2)]]{}); // OK
 
 #ifdef TRIGGER_ERROR
     [[intel::max_global_work_dim(1)]] int Var = 0; // expected-error{{'max_global_work_dim' attribute only applies to functions}}
 
-    h.single_task<class test_kernel13>(
+    h.single_task<class test_kernel12>(
         []() [[intel::max_global_work_dim(-8)]]{}); // expected-error{{'max_global_work_dim' attribute requires a non-negative integral compile time constant expression}}
 
-    h.single_task<class test_kernell4>(
+    h.single_task<class test_kernell3>(
         []() [[intel::max_global_work_dim(3),      // expected-note {{previous attribute is here}}
                intel::max_global_work_dim(2)]]{}); // expected-warning{{attribute 'max_global_work_dim' is already applied with different arguments}}
 
-    h.single_task<class test_kernel15>(TRIFuncObjBad());
+    h.single_task<class test_kernel14>(TRIFuncObjBad());
 
-    h.single_task<class test_kernel16>(TRIFuncObjBad1());
+    h.single_task<class test_kernel15>(TRIFuncObjBad1());
 
-    h.single_task<class test_kernel17>(TRIFuncObjBad2());
+    h.single_task<class test_kernel16>(TRIFuncObjBad2());
 
-    h.single_task<class test_kernel18>(TRIFuncObjBad3());
+    h.single_task<class test_kernel17>(TRIFuncObjBad3());
 
-    h.single_task<class test_kernel19>(TRIFuncObjBad4());
+    h.single_task<class test_kernel18>(TRIFuncObjBad4());
 
-    h.single_task<class test_kernel20>(TRIFuncObjBad5());
+    h.single_task<class test_kernel19>(TRIFuncObjBad5());
 
-    h.single_task<class test_kernel21>(TRIFuncObjBad6());
+    h.single_task<class test_kernel20>(TRIFuncObjBad6());
 
-    h.single_task<class test_kernel22>(TRIFuncObjBad7());
+    h.single_task<class test_kernel21>(TRIFuncObjBad7());
 
-    h.single_task<class test_kernel23>(TRIFuncObjBad8());
+    h.single_task<class test_kernel22>(TRIFuncObjBad8());
 
-    h.single_task<class test_kernel24>(
+    h.single_task<class test_kernel23>(
         []() [[intel::max_global_work_dim(4)]]{}); // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}}
 #endif // TRIGGER_ERROR
   });

From 0be1b6210abbf855130185f0a2ca089c350721a3 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Sat, 20 Mar 2021 21:33:02 -0700
Subject: [PATCH 13/29] Update logic for dropping duplicate attribute with same
 arguments and adds tests

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp               |  6 ++-
 .../intel-max-global-work-dim-device.cpp      | 43 ++++++++-----------
 ...ice-intel-max-global-work-dim-template.cpp | 16 +++++++
 3 files changed, 38 insertions(+), 27 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 0e12ef0aebbfa..7a4c7e40f705e 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3452,7 +3452,8 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
           Diag(CI.getLoc(), diag::warn_duplicate_attribute) << CI;
           Diag(DeclAttr->getLoc(), diag::note_previous_attribute);
         }
-        // If there is no mismatch, silently ignore duplicate attribute.
+        // If there is no mismatch, drop any duplicate attributes.
+        D->dropAttr<SYCLIntelMaxGlobalWorkDimAttr>();
         return;
       }
     }
@@ -3494,7 +3495,8 @@ SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr(
           Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A;
           Diag(A.getLoc(), diag::note_previous_attribute);
         }
-        // If there is no mismatch, silently ignore duplicate attribute.
+        // If there is no mismatch, drop any duplicate attributes.
+        D->dropAttr<SYCLIntelMaxGlobalWorkDimAttr>();
         return nullptr;
       }
     }
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 5c205460e65e0..93f3a1c13685b 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
@@ -33,9 +33,12 @@ struct Func {
   [[intelfpga::max_global_work_dim(2)]] void operator()() const {}
 };
 
-// No diagnostic is thrown since arguments match.
-[[intel::max_global_work_dim(2)]] void bar();
-[[intel::max_global_work_dim(2)]] void bar() {}
+// No diagnostic is thrown since arguments match. Silently ignore duplicate attribute.
+[[intel::max_global_work_dim(1)]] void bar();
+[[intel::max_global_work_dim(1)]] void bar() {}
+
+[[intel::max_global_work_dim(2)]] void func_ignore();
+[[intel::max_global_work_dim(2)]] void func_ignore();
 
 // Checking of different argument values.
 [[intel::max_global_work_dim(2)]] void baz();  // expected-note {{previous attribute is here}}
@@ -349,45 +352,35 @@ int main() {
     // CHECK-NEXT:  value: Int 3
     // CHECK-NEXT:  IntegerLiteral{{.*}}3{{$}}
 
-    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel11 'void ()'
-    // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
-    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
-    // CHECK-NEXT:  value: Int 2
-    // CHECK-NEXT:  IntegerLiteral{{.*}}2{{$}}
-    // No diagnostic is thrown since arguments match. Duplicate attribute is silently ignored.
-    h.single_task<class test_kernel11>(
-        []() [[intel::max_global_work_dim(2),
-               intel::max_global_work_dim(2)]]{}); // OK
-
 #ifdef TRIGGER_ERROR
     [[intel::max_global_work_dim(1)]] int Var = 0; // expected-error{{'max_global_work_dim' attribute only applies to functions}}
 
-    h.single_task<class test_kernel12>(
+    h.single_task<class test_kernel11>(
         []() [[intel::max_global_work_dim(-8)]]{}); // expected-error{{'max_global_work_dim' attribute requires a non-negative integral compile time constant expression}}
 
-    h.single_task<class test_kernell3>(
+    h.single_task<class test_kernell2>(
         []() [[intel::max_global_work_dim(3),      // expected-note {{previous attribute is here}}
                intel::max_global_work_dim(2)]]{}); // expected-warning{{attribute 'max_global_work_dim' is already applied with different arguments}}
 
-    h.single_task<class test_kernel14>(TRIFuncObjBad());
+    h.single_task<class test_kernel13>(TRIFuncObjBad());
 
-    h.single_task<class test_kernel15>(TRIFuncObjBad1());
+    h.single_task<class test_kernel14>(TRIFuncObjBad1());
 
-    h.single_task<class test_kernel16>(TRIFuncObjBad2());
+    h.single_task<class test_kernel15>(TRIFuncObjBad2());
 
-    h.single_task<class test_kernel17>(TRIFuncObjBad3());
+    h.single_task<class test_kernel16>(TRIFuncObjBad3());
 
-    h.single_task<class test_kernel18>(TRIFuncObjBad4());
+    h.single_task<class test_kernel17>(TRIFuncObjBad4());
 
-    h.single_task<class test_kernel19>(TRIFuncObjBad5());
+    h.single_task<class test_kernel18>(TRIFuncObjBad5());
 
-    h.single_task<class test_kernel20>(TRIFuncObjBad6());
+    h.single_task<class test_kernel19>(TRIFuncObjBad6());
 
-    h.single_task<class test_kernel21>(TRIFuncObjBad7());
+    h.single_task<class test_kernel20>(TRIFuncObjBad7());
 
-    h.single_task<class test_kernel22>(TRIFuncObjBad8());
+    h.single_task<class test_kernel21>(TRIFuncObjBad8());
 
-    h.single_task<class test_kernel23>(
+    h.single_task<class test_kernel22>(
         []() [[intel::max_global_work_dim(4)]]{}); // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}}
 #endif // TRIGGER_ERROR
   });
diff --git a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
index adeec80e5ddb0..08b84826ebcca 100644
--- a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
+++ b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
@@ -68,6 +68,14 @@ template <int size>
 template <int size>
 [[intel::max_global_work_dim(size)]] void func4() {} // expected-warning {{attribute 'max_global_work_dim' is already applied with different arguments}}
 
+// Test that checks template instantiations for same argument values. We don't duplicate the attribute with the same arguments.
+template <int size>
+[[intel::max_global_work_dim(2)]] void func5();
+
+template <int size>
+[[intel::max_global_work_dim(size)]] void func5() {}
+
+
 int check() {
   // no error expected
   func3<3>();
@@ -75,6 +83,8 @@ int check() {
   func3<-1>();
   //expected-note@+1 {{in instantiation of function template specialization 'func4<2>' requested here}}
   func4<2>();
+  // no error expected
+  func5<2>();
   return 0;
 }
 
@@ -86,3 +96,9 @@ int check() {
 // CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}}
 // CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
 // CHECK-NEXT: IntegerLiteral{{.*}}3{{$}}
+
+// CHECK: FunctionDecl {{.*}} {{.*}} func5 'void ()'
+// CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}}
+// CHECK-NEXT: ConstantExpr {{.*}} 'int'
+// CHECK-NEXT: value: Int 2
+// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}

From afb684981fc233d0827c6f5315d57e4e2cb233ac Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Sun, 21 Mar 2021 06:49:03 -0700
Subject: [PATCH 14/29] Refcator duplicate attr logic for same argument and add
 test

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp               |  4 +--
 .../intel-max-global-work-dim-device.cpp      | 34 ++++++++++++-------
 2 files changed, 24 insertions(+), 14 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 7a4c7e40f705e..f06845e673428 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3451,10 +3451,10 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
         if (ArgVal != DeclExpr->getResultAsAPSInt()) {
           Diag(CI.getLoc(), diag::warn_duplicate_attribute) << CI;
           Diag(DeclAttr->getLoc(), diag::note_previous_attribute);
+          return;
         }
         // If there is no mismatch, drop any duplicate attributes.
         D->dropAttr<SYCLIntelMaxGlobalWorkDimAttr>();
-        return;
       }
     }
     // If the declaration has a SYCLIntelMaxWorkGroupSizeAttr or
@@ -3494,10 +3494,10 @@ SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr(
         if (DeclExpr->getResultAsAPSInt() != MergeExpr->getResultAsAPSInt()) {
           Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A;
           Diag(A.getLoc(), diag::note_previous_attribute);
+          return nullptr;
         }
         // If there is no mismatch, drop any duplicate attributes.
         D->dropAttr<SYCLIntelMaxGlobalWorkDimAttr>();
-        return nullptr;
       }
     }
   }
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 93f3a1c13685b..f253bd6d91277 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
@@ -352,35 +352,45 @@ int main() {
     // CHECK-NEXT:  value: Int 3
     // CHECK-NEXT:  IntegerLiteral{{.*}}3{{$}}
 
+    // Ignore duplicate attribute with same argument value.
+    h.single_task<class test_kernell1>(
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernell1 'void ()'
+    // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 3
+    // CHECK-NEXT:  IntegerLiteral{{.*}}3{{$}}
+        []() [[intel::max_global_work_dim(3),
+               intel::max_global_work_dim(3)]]{}); // Ok
+
 #ifdef TRIGGER_ERROR
     [[intel::max_global_work_dim(1)]] int Var = 0; // expected-error{{'max_global_work_dim' attribute only applies to functions}}
 
-    h.single_task<class test_kernel11>(
+    h.single_task<class test_kernel12>(
         []() [[intel::max_global_work_dim(-8)]]{}); // expected-error{{'max_global_work_dim' attribute requires a non-negative integral compile time constant expression}}
 
-    h.single_task<class test_kernell2>(
+    h.single_task<class test_kernell3>(
         []() [[intel::max_global_work_dim(3),      // expected-note {{previous attribute is here}}
                intel::max_global_work_dim(2)]]{}); // expected-warning{{attribute 'max_global_work_dim' is already applied with different arguments}}
 
-    h.single_task<class test_kernel13>(TRIFuncObjBad());
+    h.single_task<class test_kernel14>(TRIFuncObjBad());
 
-    h.single_task<class test_kernel14>(TRIFuncObjBad1());
+    h.single_task<class test_kernel15>(TRIFuncObjBad1());
 
-    h.single_task<class test_kernel15>(TRIFuncObjBad2());
+    h.single_task<class test_kernel16>(TRIFuncObjBad2());
 
-    h.single_task<class test_kernel16>(TRIFuncObjBad3());
+    h.single_task<class test_kernel17>(TRIFuncObjBad3());
 
-    h.single_task<class test_kernel17>(TRIFuncObjBad4());
+    h.single_task<class test_kernel18>(TRIFuncObjBad4());
 
-    h.single_task<class test_kernel18>(TRIFuncObjBad5());
+    h.single_task<class test_kernel19>(TRIFuncObjBad5());
 
-    h.single_task<class test_kernel19>(TRIFuncObjBad6());
+    h.single_task<class test_kernel20>(TRIFuncObjBad6());
 
-    h.single_task<class test_kernel20>(TRIFuncObjBad7());
+    h.single_task<class test_kernel21>(TRIFuncObjBad7());
 
-    h.single_task<class test_kernel21>(TRIFuncObjBad8());
+    h.single_task<class test_kernel22>(TRIFuncObjBad8());
 
-    h.single_task<class test_kernel22>(
+    h.single_task<class test_kernel23>(
         []() [[intel::max_global_work_dim(4)]]{}); // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}}
 #endif // TRIGGER_ERROR
   });

From 70f87ca59be0e83d76355ff78e037b8f6b68ac27 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Tue, 23 Mar 2021 20:00:01 -0700
Subject: [PATCH 15/29] Address review comments

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp                    |  6 ++----
 .../SemaSYCL/intel-max-global-work-dim-device.cpp  |  2 +-
 ...l-device-intel-max-global-work-dim-template.cpp | 14 ++++----------
 3 files changed, 7 insertions(+), 15 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index f06845e673428..ff67775b86487 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3451,10 +3451,9 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
         if (ArgVal != DeclExpr->getResultAsAPSInt()) {
           Diag(CI.getLoc(), diag::warn_duplicate_attribute) << CI;
           Diag(DeclAttr->getLoc(), diag::note_previous_attribute);
-          return;
         }
         // If there is no mismatch, drop any duplicate attributes.
-        D->dropAttr<SYCLIntelMaxGlobalWorkDimAttr>();
+        return;
       }
     }
     // If the declaration has a SYCLIntelMaxWorkGroupSizeAttr or
@@ -3494,10 +3493,9 @@ SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr(
         if (DeclExpr->getResultAsAPSInt() != MergeExpr->getResultAsAPSInt()) {
           Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A;
           Diag(A.getLoc(), diag::note_previous_attribute);
-          return nullptr;
         }
         // If there is no mismatch, drop any duplicate attributes.
-        D->dropAttr<SYCLIntelMaxGlobalWorkDimAttr>();
+        return nullptr;
       }
     }
   }
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 f253bd6d91277..938220de93442 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
@@ -33,7 +33,7 @@ struct Func {
   [[intelfpga::max_global_work_dim(2)]] void operator()() const {}
 };
 
-// No diagnostic is thrown since arguments match. Silently ignore duplicate attribute.
+// No diagnostic is thrown since arguments match.
 [[intel::max_global_work_dim(1)]] void bar();
 [[intel::max_global_work_dim(1)]] void bar() {}
 
diff --git a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
index 08b84826ebcca..facd0c17a46cb 100644
--- a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
+++ b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
@@ -68,14 +68,6 @@ template <int size>
 template <int size>
 [[intel::max_global_work_dim(size)]] void func4() {} // expected-warning {{attribute 'max_global_work_dim' is already applied with different arguments}}
 
-// Test that checks template instantiations for same argument values. We don't duplicate the attribute with the same arguments.
-template <int size>
-[[intel::max_global_work_dim(2)]] void func5();
-
-template <int size>
-[[intel::max_global_work_dim(size)]] void func5() {}
-
-
 int check() {
   // no error expected
   func3<3>();
@@ -83,11 +75,13 @@ int check() {
   func3<-1>();
   //expected-note@+1 {{in instantiation of function template specialization 'func4<2>' requested here}}
   func4<2>();
-  // no error expected
-  func5<2>();
   return 0;
 }
 
+// No diagnostic is thrown since arguments match. Duplicate attribute is silently ignored.
+[[intel::max_global_work_dim(2)]]
+[[intel::max_global_work_dim(2)]] void func5() {}
+
 // CHECK: FunctionDecl {{.*}} {{.*}} func3 'void ()'
 // CHECK: TemplateArgument integral 3
 // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}}

From 17166deabf49d5ea3c62901133e577779d131375 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Thu, 25 Mar 2021 11:04:23 -0700
Subject: [PATCH 16/29] address review comments

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp               |  4 +-
 .../intel-max-global-work-dim-device.cpp      |  4 +-
 ...ice-intel-max-global-work-dim-template.cpp | 67 ++++++++++++++++++-
 3 files changed, 68 insertions(+), 7 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index ff67775b86487..854fba17d97d0 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3452,7 +3452,7 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
           Diag(CI.getLoc(), diag::warn_duplicate_attribute) << CI;
           Diag(DeclAttr->getLoc(), diag::note_previous_attribute);
         }
-        // If there is no mismatch, drop any duplicate attributes.
+        // Drop the duplicate attribute.
         return;
       }
     }
@@ -3494,7 +3494,7 @@ SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr(
           Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A;
           Diag(A.getLoc(), diag::note_previous_attribute);
         }
-        // If there is no mismatch, drop any duplicate attributes.
+        // Do not add a duplicate attribute.
         return nullptr;
       }
     }
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 938220de93442..72e6d730ca13a 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
@@ -33,7 +33,7 @@ struct Func {
   [[intelfpga::max_global_work_dim(2)]] void operator()() const {}
 };
 
-// No diagnostic is thrown since arguments match.
+// No diagnostic is emitted because the arguments match.
 [[intel::max_global_work_dim(1)]] void bar();
 [[intel::max_global_work_dim(1)]] void bar() {}
 
@@ -45,7 +45,7 @@ struct Func {
 [[intel::max_global_work_dim(1)]] void baz();  // expected-warning {{attribute 'max_global_work_dim' is already applied with different arguments}}
 
 // Checks correctness of mutual usage of different work_group_size attributes:
-// reqd_work_group_size, max_work_group_size and max_global_work_dim.
+// reqd_work_group_size, max_work_group_size, and max_global_work_dim.
 // In case the value of 'max_global_work_dim' attribute equals to 0 we shall
 // ensure that if max_work_group_size and reqd_work_group_size attributes exist,
 // they hold equal values (1, 1, 1).
diff --git a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
index facd0c17a46cb..41784f5d51379 100644
--- a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
+++ b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
@@ -68,6 +68,55 @@ template <int size>
 template <int size>
 [[intel::max_global_work_dim(size)]] void func4() {} // expected-warning {{attribute 'max_global_work_dim' is already applied with different arguments}}
 
+// Checks correctness of mutual usage of different work_group_size attributes:
+// reqd_work_group_size, max_work_group_size, and max_global_work_dim.
+// In case the value of 'max_global_work_dim' attribute equals to 0 we shall
+// ensure that if max_work_group_size and reqd_work_group_size attributes exist,
+// they hold equal values (1, 1, 1).
+template <int N>
+// expected-error@+1{{'max_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+[[intel::max_work_group_size(N, N, N)]] void func5();
+template <int N>
+[[intel::max_global_work_dim(0)]] void func5();
+
+template <int N>
+// expected-error@+1{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+[[intel::reqd_work_group_size(N)]] void func6();
+template <int N>
+[[intel::max_global_work_dim(0)]] void func6();
+
+template <int N>
+// expected-error@+1{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+[[intel::reqd_work_group_size(N, N)]] void func7();
+template <int N>
+[[intel::max_global_work_dim(0)]] void func7();
+
+template <int N>
+// expected-error@+1{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+[[intel::reqd_work_group_size(N, N, N)]] void func8();
+template <int N>
+[[intel::max_global_work_dim(0)]] void func8();
+
+template <int N>
+[[intel::max_work_group_size(N, N, N)]] void func9();
+template <int N>
+[[intel::max_global_work_dim(0)]] void func9();
+
+template <int N>
+[[intel::reqd_work_group_size(N)]] void func10();
+template <int N>
+[[intel::max_global_work_dim(0)]] void func10();
+
+template <int N>
+[[intel::reqd_work_group_size(N, N)]] void func11();
+template <int N>
+[[intel::max_global_work_dim(0)]] void func11();
+
+template <int N>
+[[intel::reqd_work_group_size(N, N, N)]] void func12();
+template <int N>
+[[intel::max_global_work_dim(0)]] void func12();
+
 int check() {
   // no error expected
   func3<3>();
@@ -75,12 +124,24 @@ int check() {
   func3<-1>();
   //expected-note@+1 {{in instantiation of function template specialization 'func4<2>' requested here}}
   func4<2>();
+  //expected-note@+1 {{in instantiation of function template specialization 'func5<2>' requested here}}
+  func5<2>();
+  //expected-note@+1 {{in instantiation of function template specialization 'func6<2>' requested here}}
+  func6<2>();
+  //expected-note@+1 {{in instantiation of function template specialization 'func7<2>' requested here}}
+  func7<2>();
+  //expected-note@+1 {{in instantiation of function template specialization 'func8<2>' requested here}}
+  func8<2>();
+  func9<1>();  // OK
+  func10<1>(); // OK
+  func11<1>(); // OK
+  func12<1>(); // OK
   return 0;
 }
 
-// No diagnostic is thrown since arguments match. Duplicate attribute is silently ignored.
+// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored.
 [[intel::max_global_work_dim(2)]]
-[[intel::max_global_work_dim(2)]] void func5() {}
+[[intel::max_global_work_dim(2)]] void func13() {}
 
 // CHECK: FunctionDecl {{.*}} {{.*}} func3 'void ()'
 // CHECK: TemplateArgument integral 3
@@ -91,7 +152,7 @@ int check() {
 // CHECK-NEXT: NonTypeTemplateParmDecl {{.*}}
 // CHECK-NEXT: IntegerLiteral{{.*}}3{{$}}
 
-// CHECK: FunctionDecl {{.*}} {{.*}} func5 'void ()'
+// CHECK: FunctionDecl {{.*}} {{.*}} func13 'void ()'
 // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}}
 // CHECK-NEXT: ConstantExpr {{.*}} 'int'
 // CHECK-NEXT: value: Int 2

From 1ba0baaf8b1b7c3644ff40f759b656c4d1e03a0f Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Thu, 13 May 2021 13:54:14 -0700
Subject: [PATCH 17/29] handle dependent argument values and addressed review
 comments

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp               | 40 +++++++++++++++----
 ...ice-intel-max-global-work-dim-template.cpp | 35 ++++++----------
 2 files changed, 44 insertions(+), 31 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 7fe89e6463f9d..d6cfa24dc1691 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3525,26 +3525,50 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
     // ReqdWorkGroupSizeAttr, check to see if they hold equal values
     // (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr
     // equals to 0.
-    auto Check = [this, &CI](const auto *A) {
-      Optional<llvm::APSInt> XDimVal = A->getXDimVal(Context);
-      Optional<llvm::APSInt> YDimVal = A->getYDimVal(Context);
-      Optional<llvm::APSInt> ZDimVal = A->getZDimVal(Context);
-      if (*XDimVal != 1 || *YDimVal != 1 || *ZDimVal != 1) {
-        Diag(A->getLocation(), diag::err_sycl_x_y_z_arguments_must_be_one)
-            << A << CI;
-        return false;
+    auto Check = [this, &CI](const auto *A) -> bool {
+      Expr *XDimExpr = A->getXDim();
+      Expr *YDimExpr = A->getYDim();
+      Expr *ZDimExpr = A->getZDim();
+
+      if (!XDimExpr->isValueDependent() && !YDimExpr->isValueDependent() &&
+          !ZDimExpr->isValueDependent()) {
+        llvm::APSInt XDimVal, YDimVal, ZDimVal;
+        ExprResult XDim = VerifyIntegerConstantExpression(XDimExpr, &XDimVal);
+        ExprResult YDim = VerifyIntegerConstantExpression(YDimExpr, &YDimVal);
+        ExprResult ZDim = VerifyIntegerConstantExpression(ZDimExpr, &ZDimVal);
+
+        if (XDim.isInvalid())
+          return true;
+        XDimExpr = XDim.get();
+
+        if (YDim.isInvalid())
+          return true;
+        YDimExpr = YDim.get();
+
+        if (ZDim.isInvalid())
+          return true;
+        ZDimExpr = ZDim.get();
+
+        if (XDimVal != 1 || YDimVal != 1 || ZDimVal != 1) {
+	  Diag(A->getLocation(), diag::err_sycl_x_y_z_arguments_must_be_one)
+              << A << CI;
+          return false;
+        }
       }
       return true;
     };
+
     if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
       if (ArgVal == 0 && !Check(DeclAttr))
         return;
     }
+
     if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
       if (ArgVal == 0 && !Check(DeclAttr))
         return;
     }
   }
+
   D->addAttr(::new (Context) SYCLIntelMaxGlobalWorkDimAttr(Context, CI, E));
 }
 
diff --git a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
index 41784f5d51379..4b909b9b6d024 100644
--- a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
+++ b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
@@ -74,26 +74,22 @@ template <int size>
 // ensure that if max_work_group_size and reqd_work_group_size attributes exist,
 // they hold equal values (1, 1, 1).
 template <int N>
-// expected-error@+1{{'max_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
-[[intel::max_work_group_size(N, N, N)]] void func5();
+[[intel::max_work_group_size(N, N, N)]] void func5(); // expected-error {{'max_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
 template <int N>
 [[intel::max_global_work_dim(0)]] void func5();
 
 template <int N>
-// expected-error@+1{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
-[[intel::reqd_work_group_size(N)]] void func6();
+[[intel::reqd_work_group_size(N)]] void func6(); // expected-error {{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
 template <int N>
 [[intel::max_global_work_dim(0)]] void func6();
 
 template <int N>
-// expected-error@+1{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
-[[intel::reqd_work_group_size(N, N)]] void func7();
+[[intel::reqd_work_group_size(N, N)]] void func7(); // expected-error {{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
 template <int N>
 [[intel::max_global_work_dim(0)]] void func7();
 
 template <int N>
-// expected-error@+1{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
-[[intel::reqd_work_group_size(N, N, N)]] void func8();
+[[intel::reqd_work_group_size(N, N, N)]] void func8(); // expected-error {{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
 template <int N>
 [[intel::max_global_work_dim(0)]] void func8();
 
@@ -113,25 +109,18 @@ template <int N>
 [[intel::max_global_work_dim(0)]] void func11();
 
 template <int N>
-[[intel::reqd_work_group_size(N, N, N)]] void func12();
+[[cl::reqd_work_group_size(N, N, N)]] void func12();
 template <int N>
 [[intel::max_global_work_dim(0)]] void func12();
 
 int check() {
-  // no error expected
-  func3<3>();
-  //expected-note@+1{{in instantiation of function template specialization 'func3<-1>' requested here}}
-  func3<-1>();
-  //expected-note@+1 {{in instantiation of function template specialization 'func4<2>' requested here}}
-  func4<2>();
-  //expected-note@+1 {{in instantiation of function template specialization 'func5<2>' requested here}}
-  func5<2>();
-  //expected-note@+1 {{in instantiation of function template specialization 'func6<2>' requested here}}
-  func6<2>();
-  //expected-note@+1 {{in instantiation of function template specialization 'func7<2>' requested here}}
-  func7<2>();
-  //expected-note@+1 {{in instantiation of function template specialization 'func8<2>' requested here}}
-  func8<2>();
+  func3<3>();  // OK
+  func3<-1>(); // expected-note {{in instantiation of function template specialization 'func3<-1>' requested here}}
+  func4<2>();  // expected-note {{in instantiation of function template specialization 'func4<2>' requested here}}
+  func5<2>();  // expected-note {{in instantiation of function template specialization 'func5<2>' requested here}}
+  func6<2>();  // expected-note {{in instantiation of function template specialization 'func6<2>' requested here}}
+  func7<2>();  // expected-note {{in instantiation of function template specialization 'func7<2>' requested here}}
+  func8<2>();  // expected-note {{in instantiation of function template specialization 'func8<2>' requested here}}
   func9<1>();  // OK
   func10<1>(); // OK
   func11<1>(); // OK

From c4d5423dc1fec3164b9625d56135717bfc451e0c Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Thu, 13 May 2021 14:04:36 -0700
Subject: [PATCH 18/29] fix format issues

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 d6cfa24dc1691..70708a564fd4b 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3550,7 +3550,7 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
         ZDimExpr = ZDim.get();
 
         if (XDimVal != 1 || YDimVal != 1 || ZDimVal != 1) {
-	  Diag(A->getLocation(), diag::err_sycl_x_y_z_arguments_must_be_one)
+          Diag(A->getLocation(), diag::err_sycl_x_y_z_arguments_must_be_one)
               << A << CI;
           return false;
         }

From 8600280a6f69a0532c098d4ff07e5bcd6a915f8e Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Thu, 13 May 2021 14:14:00 -0700
Subject: [PATCH 19/29] fix wrong value

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

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 70708a564fd4b..e382ade0a1639 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3538,15 +3538,15 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
         ExprResult ZDim = VerifyIntegerConstantExpression(ZDimExpr, &ZDimVal);
 
         if (XDim.isInvalid())
-          return true;
+          return false;
         XDimExpr = XDim.get();
 
         if (YDim.isInvalid())
-          return true;
+          return false;
         YDimExpr = YDim.get();
 
         if (ZDim.isInvalid())
-          return true;
+          return false;
         ZDimExpr = ZDim.get();
 
         if (XDimVal != 1 || YDimVal != 1 || ZDimVal != 1) {

From 56a36ee8157046fbfd0e97edd4317167c7050c0b Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Fri, 14 May 2021 10:39:35 -0700
Subject: [PATCH 20/29] adress review comments

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

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index e382ade0a1639..a091754ee070a 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3539,15 +3539,12 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
 
         if (XDim.isInvalid())
           return false;
-        XDimExpr = XDim.get();
 
         if (YDim.isInvalid())
           return false;
-        YDimExpr = YDim.get();
 
         if (ZDim.isInvalid())
           return false;
-        ZDimExpr = ZDim.get();
 
         if (XDimVal != 1 || YDimVal != 1 || ZDimVal != 1) {
           Diag(A->getLocation(), diag::err_sycl_x_y_z_arguments_must_be_one)

From d8effc1e22d6d94eec34607e35ecce33ce7f1e17 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Mon, 24 May 2021 09:07:45 -0700
Subject: [PATCH 21/29] address review comments

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp               | 117 +++++++++++++-----
 .../intel-max-global-work-dim-device.cpp      |  56 ++++++++-
 2 files changed, 138 insertions(+), 35 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index a091754ee070a..b4651cd7473c5 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -2953,19 +2953,6 @@ static void handleWeakImportAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
 // they hold equal values (1, 1, 1).
 static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) {
   bool Result = true;
-  auto checkZeroDim = [&S, &AL](auto &A, size_t X, size_t Y, size_t Z,
-                                bool ReverseAttrs = false) -> bool {
-    if (X != 1 || Y != 1 || Z != 1) {
-      auto Diag =
-          S.Diag(AL.getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one);
-      if (ReverseAttrs)
-        Diag << AL << A;
-      else
-        Diag << A << AL;
-      return false;
-    }
-    return true;
-  };
 
   // Returns the unsigned constant integer value represented by
   // given expression.
@@ -2978,15 +2965,6 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) {
   if (AL.getKind() == ParsedAttr::AT_SYCLIntelMaxWorkGroupSize)
     S.CheckDeprecatedSYCLAttributeSpelling(AL);
 
-  if (const auto *A = D->getAttr<SYCLIntelMaxGlobalWorkDimAttr>()) {
-    if ((A->getValue()->getIntegerConstantExpr(Ctx)->getSExtValue()) == 0) {
-      Result &= checkZeroDim(A, getExprValue(AL.getArgAsExpr(0), Ctx),
-                             getExprValue(AL.getArgAsExpr(1), Ctx),
-                             getExprValue(AL.getArgAsExpr(2), Ctx),
-                             /*ReverseAttrs=*/true);
-    }
-  }
-
   if (const auto *A = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
     if (!((getExprValue(AL.getArgAsExpr(0), Ctx) <=
            getExprValue(A->getXDim(), Ctx)) &&
@@ -3082,6 +3060,20 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
         return;
       }
     }
+
+    if (const auto *A = D->getAttr<SYCLIntelMaxGlobalWorkDimAttr>()) {
+      int64_t MaxGlobalWorkDim =
+          A->getValue()->getIntegerConstantExpr(Ctx)->getSExtValue();
+
+      if (MaxGlobalWorkDim == 0 &&
+	  (XDimVal.getZExtValue() != 1 || YDimVal.getZExtValue() != 1 ||
+	   ZDimVal.getZExtValue() != 1)) {
+        S.Diag(AL.getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
+            << AL << A;
+        return;
+      }
+    }
+
     if (const auto *ExistingAttr = D->getAttr<WorkGroupAttr>()) {
       // Compare attribute arguments value and warn for a mismatch.
       if (ExistingAttr->getXDimVal(Ctx) != XDimVal ||
@@ -3481,6 +3473,27 @@ static void handleSYCLIntelSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D,
 }
 
 // Handles max_global_work_dim.
+// Returns a ArgCheck value; EqualToOne means all argument values are one,
+// NotEqualToOne means all argument values are not one, and Unknown means that
+// the argument values are invalid.
+enum class ArgCheck { Unknown, EqualToOne, NotEqualToOne };
+static ArgCheck AreArgValuesToOne(const Expr *E, const Expr *E1,
+		                  const Expr *E2) {
+  // If any of the expression is value dependent, test is invalid.
+  const auto *XDimExpr = dyn_cast<ConstantExpr>(E);
+  const auto *YDimExpr = dyn_cast<ConstantExpr>(E1);
+  const auto *ZDimExpr = dyn_cast<ConstantExpr>(E2);
+  if (!XDimExpr || !YDimExpr || !ZDimExpr)
+    return ArgCheck::Unknown;
+
+  // Check the argument values.
+  return (XDimExpr->getResultAsAPSInt() == 1 ||
+          YDimExpr->getResultAsAPSInt() == 1 ||
+          ZDimExpr->getResultAsAPSInt() == 1)
+          ? ArgCheck::EqualToOne
+          : ArgCheck::NotEqualToOne;
+}
+
 void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
                                             const AttributeCommonInfo &CI,
                                             Expr *E) {
@@ -3537,13 +3550,7 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
         ExprResult YDim = VerifyIntegerConstantExpression(YDimExpr, &YDimVal);
         ExprResult ZDim = VerifyIntegerConstantExpression(ZDimExpr, &ZDimVal);
 
-        if (XDim.isInvalid())
-          return false;
-
-        if (YDim.isInvalid())
-          return false;
-
-        if (ZDim.isInvalid())
+        if (XDim.isInvalid() || YDim.isInvalid() || ZDim.isInvalid())
           return false;
 
         if (XDimVal != 1 || YDimVal != 1 || ZDimVal != 1) {
@@ -3586,6 +3593,58 @@ SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr(
     }
   }
 
+  // If the declaration has a SYCLIntelMaxWorkGroupSizeAttr or
+  // ReqdWorkGroupSizeAttr, check to see if they hold equal values
+  // (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr
+  // equals to 0.
+  if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
+    const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getValue());
+    ArgCheck Results[] = {
+        AreArgValuesToOne(DeclAttr->getXDim(), DeclAttr->getYDim(),
+			  DeclAttr->getZDim())};
+
+    // Diagnose if any of the argument values of ReqdWorkGroupSizeAttr
+    // are not equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
+    // equals to 0.
+    if (MergeExpr->getResultAsAPSInt() == 0 &&
+        llvm::is_contained(Results, ArgCheck::NotEqualToOne)) {
+      Diag(DeclAttr->getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
+	  << DeclAttr << &A;
+      return nullptr;
+    }
+
+    // Do not diagnose if all argument values of ReqdWorkGroupSizeAttr
+    // are equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
+    // equals to 0.
+    if (llvm::all_of(Results,
+                     [](ArgCheck Ch) { return Ch == ArgCheck::EqualToOne; }))
+      return nullptr;
+  }
+
+  if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
+    const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getValue());
+    ArgCheck Results[] = {
+	AreArgValuesToOne(DeclAttr->getXDim(), DeclAttr->getYDim(),
+                          DeclAttr->getZDim())};
+
+    // Diagnose if any of the argument values of SYCLIntelMaxWorkGroupSizeAttr
+    // are not equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
+    // equals to 0.
+    if (MergeExpr->getResultAsAPSInt() == 0 &&
+        llvm::is_contained(Results, ArgCheck::NotEqualToOne)) {
+      Diag(DeclAttr->getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
+          << DeclAttr << &A;
+      return nullptr;
+    }
+
+    // Do not diagnose if all argument values of SYCLIntelMaxWorkGroupSizeAttr
+    // are equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
+    // equals to 0.
+    if (llvm::all_of(Results,
+                     [](ArgCheck Ch) { return Ch == ArgCheck::EqualToOne; }))
+      return nullptr;
+  }
+
   return ::new (Context)
       SYCLIntelMaxGlobalWorkDimAttr(Context, A, A.getValue());
 }
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 72e6d730ca13a..a304fa03fccc8 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
@@ -44,6 +44,11 @@ struct Func {
 [[intel::max_global_work_dim(2)]] void baz();  // expected-note {{previous attribute is here}}
 [[intel::max_global_work_dim(1)]] void baz();  // expected-warning {{attribute 'max_global_work_dim' is already applied with different arguments}}
 
+struct TRIFuncObj {
+  [[intel::max_global_work_dim(0)]] void operator()() const; // expected-note {{previous attribute is here}}
+};
+[[intel::max_global_work_dim(1)]] void TRIFuncObj::operator()() const {} // expected-warning {{attribute 'max_global_work_dim' is already applied with different arguments}}
+
 // Checks correctness of mutual usage of different work_group_size attributes:
 // reqd_work_group_size, max_work_group_size, and max_global_work_dim.
 // In case the value of 'max_global_work_dim' attribute equals to 0 we shall
@@ -133,10 +138,43 @@ struct TRIFuncObjBad3 {
   operator()() const {}
 };
 
+struct TRIFuncObjBad4 {
+  [[intel::max_global_work_dim(0)]] void
+  operator()() const;
+};
+
+[[cl::reqd_work_group_size(4, 4, 4)]] // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+void TRIFuncObjBad4::operator()() const {}
+
+struct TRIFuncObjBad5 {
+  [[intel::max_global_work_dim(0)]] void
+  operator()() const;
+};
+
+[[intel::reqd_work_group_size(4, 4, 4)]] // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+void TRIFuncObjBad5::operator()() const {}
+
+struct TRIFuncObjBad6 {
+  [[intel::max_global_work_dim(0)]] void
+  operator()() const;
+};
+
+[[intel::max_work_group_size(4, 4, 4)]] // expected-error{{'max_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+void TRIFuncObjBad6::operator()() const {}
+
+/*
+struct TRIFuncObjBad6 {
+  [[intel::max_work_group_size(4, 4, 4)]] void
+  operator()() const;
+};
+
+[[intel::max_global_work_dim(0)]]
+void TRIFuncObjBad6::operator()() const {}
+*/
 // Tests for incorrect argument values for Intel FPGA function attributes:
 // reqd_work_group_size, max_work_group_size and max_global_work_dim.
 
-struct TRIFuncObjBad4 {
+struct TRIFuncObjBad7 {
   // expected-error@+2{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
   // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}}
   [[intel::reqd_work_group_size(-4, 1)]]
@@ -144,25 +182,25 @@ struct TRIFuncObjBad4 {
   operator()() const {}
 };
 
-struct TRIFuncObjBad5 {
+struct TRIFuncObjBad8 {
   [[intel::max_work_group_size(4, 4, 4.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}}
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
 
-struct TRIFuncObjBad6 {
+struct TRIFuncObjBad9 {
   [[cl::reqd_work_group_size(0, 4, 4)]] // expected-error{{'reqd_work_group_size' attribute must be greater than 0}}
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
 
-struct TRIFuncObjBad7 {
+struct TRIFuncObjBad10 {
   [[intel::reqd_work_group_size(4)]]
   [[intel::max_global_work_dim(-2)]] // expected-error{{'max_global_work_dim' attribute requires a non-negative integral compile time constant expression}}
   void operator()() const {}
 };
 
-struct TRIFuncObjBad8 {
+struct TRIFuncObjBad11 {
   [[intel::max_work_group_size(4, 4, 4)]]
   [[intel::max_global_work_dim(4.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}}
   void operator()() const {}
@@ -390,7 +428,13 @@ int main() {
 
     h.single_task<class test_kernel22>(TRIFuncObjBad8());
 
-    h.single_task<class test_kernel23>(
+    h.single_task<class test_kernel23>(TRIFuncObjBad9());
+
+    h.single_task<class test_kernel24>(TRIFuncObjBad10());
+
+    h.single_task<class test_kernel25>(TRIFuncObjBad11());
+
+    h.single_task<class test_kernel26>(
         []() [[intel::max_global_work_dim(4)]]{}); // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}}
 #endif // TRIGGER_ERROR
   });

From a9ad2c4301001993f62cbae9e721028b97a6f713 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Mon, 24 May 2021 09:23:57 -0700
Subject: [PATCH 22/29] Fix clang format issues

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

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 4ac3f6b328cdf..38838b482c183 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3115,8 +3115,8 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
           A->getValue()->getIntegerConstantExpr(Ctx)->getSExtValue();
 
       if (MaxGlobalWorkDim == 0 &&
-	  (XDimVal.getZExtValue() != 1 || YDimVal.getZExtValue() != 1 ||
-	   ZDimVal.getZExtValue() != 1)) {
+          (XDimVal.getZExtValue() != 1 || YDimVal.getZExtValue() != 1 ||
+           ZDimVal.getZExtValue() != 1)) {
         S.Diag(AL.getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
             << AL << A;
         return;
@@ -3548,7 +3548,7 @@ static void handleSYCLIntelSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D,
 // the argument values are invalid.
 enum class ArgCheck { Unknown, EqualToOne, NotEqualToOne };
 static ArgCheck AreArgValuesToOne(const Expr *E, const Expr *E1,
-		                  const Expr *E2) {
+                                  const Expr *E2) {
   // If any of the expression is value dependent, test is invalid.
   const auto *XDimExpr = dyn_cast<ConstantExpr>(E);
   const auto *YDimExpr = dyn_cast<ConstantExpr>(E1);
@@ -3560,8 +3560,8 @@ static ArgCheck AreArgValuesToOne(const Expr *E, const Expr *E1,
   return (XDimExpr->getResultAsAPSInt() == 1 ||
           YDimExpr->getResultAsAPSInt() == 1 ||
           ZDimExpr->getResultAsAPSInt() == 1)
-          ? ArgCheck::EqualToOne
-          : ArgCheck::NotEqualToOne;
+             ? ArgCheck::EqualToOne
+             : ArgCheck::NotEqualToOne;
 }
 
 void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
@@ -3669,9 +3669,8 @@ SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr(
   // equals to 0.
   if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
     const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getValue());
-    ArgCheck Results[] = {
-        AreArgValuesToOne(DeclAttr->getXDim(), DeclAttr->getYDim(),
-			  DeclAttr->getZDim())};
+    ArgCheck Results[] = {AreArgValuesToOne(
+        DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim())};
 
     // Diagnose if any of the argument values of ReqdWorkGroupSizeAttr
     // are not equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
@@ -3679,7 +3678,7 @@ SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr(
     if (MergeExpr->getResultAsAPSInt() == 0 &&
         llvm::is_contained(Results, ArgCheck::NotEqualToOne)) {
       Diag(DeclAttr->getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
-	  << DeclAttr << &A;
+          << DeclAttr << &A;
       return nullptr;
     }
 
@@ -3693,9 +3692,8 @@ SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr(
 
   if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
     const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getValue());
-    ArgCheck Results[] = {
-	AreArgValuesToOne(DeclAttr->getXDim(), DeclAttr->getYDim(),
-                          DeclAttr->getZDim())};
+    ArgCheck Results[] = {AreArgValuesToOne(
+        DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim())};
 
     // Diagnose if any of the argument values of SYCLIntelMaxWorkGroupSizeAttr
     // are not equal to one when value of SYCLIntelMaxGlobalWorkDimAttr

From ce62808d31847dd6d509bbe70e0ed9b4e80da72d Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Tue, 25 May 2021 06:33:26 -0700
Subject: [PATCH 23/29] address review comments

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 clang/lib/Sema/SemaDeclAttr.cpp               | 156 ++++++++----------
 .../intel-max-global-work-dim-device.cpp      | 144 ++++++++++++----
 ...ice-intel-max-global-work-dim-template.cpp |   6 +-
 3 files changed, 186 insertions(+), 120 deletions(-)

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 38838b482c183..7aa625d33063a 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -2998,9 +2998,6 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) {
 
   ASTContext &Ctx = S.getASTContext();
 
-  if (AL.getKind() == ParsedAttr::AT_SYCLIntelMaxWorkGroupSize)
-    S.CheckDeprecatedSYCLAttributeSpelling(AL);
-
   if (const auto *A = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
     if (!((getExprValue(AL.getArgAsExpr(0), Ctx) <=
            getExprValue(A->getXDim(), Ctx)) &&
@@ -3110,16 +3107,24 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
       }
     }
 
-    if (const auto *A = D->getAttr<SYCLIntelMaxGlobalWorkDimAttr>()) {
-      int64_t MaxGlobalWorkDim =
-          A->getValue()->getIntegerConstantExpr(Ctx)->getSExtValue();
-
-      if (MaxGlobalWorkDim == 0 &&
-          (XDimVal.getZExtValue() != 1 || YDimVal.getZExtValue() != 1 ||
-           ZDimVal.getZExtValue() != 1)) {
-        S.Diag(AL.getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
-            << AL << A;
-        return;
+    // If the declaration has a SYCLIntelMaxWorkGroupSizeAttr or
+    // ReqdWorkGroupSizeAttr, check to see if they hold equal values
+    // (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr
+    // equals to 0.
+    if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxGlobalWorkDimAttr>()) {
+      if (const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue())) {
+	// If the value is dependent, we can not test anything.
+        if (!DeclExpr)
+	  return;
+
+	// Test the attribute value.
+        if (DeclExpr->getResultAsAPSInt() == 0 &&
+            (XDimVal.getZExtValue() != 1 || YDimVal.getZExtValue() != 1 ||
+             ZDimVal.getZExtValue() != 1)) {
+          S.Diag(AL.getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
+                << AL << DeclAttr;
+          return;
+        }
       }
     }
 
@@ -3543,23 +3548,20 @@ static void handleSYCLIntelSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D,
 }
 
 // Handles max_global_work_dim.
-// Returns a ArgCheck value; EqualToOne means all argument values are one,
-// NotEqualToOne means all argument values are not one, and Unknown means that
-// the argument values are invalid.
+// Returns a ArgCheck value; EqualToOne means all argument values are
+// equal to one, NotEqualToOne means all argument values are not
+// equal to one, and Unknown means that the argument values are invalid.
 enum class ArgCheck { Unknown, EqualToOne, NotEqualToOne };
-static ArgCheck AreArgValuesToOne(const Expr *E, const Expr *E1,
-                                  const Expr *E2) {
+static ArgCheck AreArgValuesToOne(const Expr *E, const Expr *E1) {
   // If any of the expression is value dependent, test is invalid.
-  const auto *XDimExpr = dyn_cast<ConstantExpr>(E);
-  const auto *YDimExpr = dyn_cast<ConstantExpr>(E1);
-  const auto *ZDimExpr = dyn_cast<ConstantExpr>(E2);
-  if (!XDimExpr || !YDimExpr || !ZDimExpr)
+  const auto *MaxGExpr = dyn_cast<ConstantExpr>(E);
+  const auto *WorkGExpr = dyn_cast<ConstantExpr>(E1);
+  if (!MaxGExpr || !WorkGExpr)
     return ArgCheck::Unknown;
 
   // Check the argument values.
-  return (XDimExpr->getResultAsAPSInt() == 1 ||
-          YDimExpr->getResultAsAPSInt() == 1 ||
-          ZDimExpr->getResultAsAPSInt() == 1)
+  return (MaxGExpr->getResultAsAPSInt() == 0 &&
+	  WorkGExpr->getResultAsAPSInt() == 1)
              ? ArgCheck::EqualToOne
              : ArgCheck::NotEqualToOne;
 }
@@ -3577,18 +3579,13 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
       return;
     E = Res.get();
 
-    // This attribute requires a non-negative value.
-    if (ArgVal < 0) {
-      Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer)
-          << CI << /*non-negative*/ 1;
-      return;
-    }
     // This attribute must be in the range [0, 3].
-    if (ArgVal > 3) {
+    if (ArgVal < 0 || ArgVal > 3) {
       Diag(E->getBeginLoc(), diag::err_attribute_argument_out_of_range)
           << CI << 0 << 3 << E->getSourceRange();
       return;
     }
+
     // Check to see if there's a duplicate attribute with different values
     // already applied to the declaration.
     if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxGlobalWorkDimAttr>()) {
@@ -3604,42 +3601,45 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
         return;
       }
     }
+
     // If the declaration has a SYCLIntelMaxWorkGroupSizeAttr or
     // ReqdWorkGroupSizeAttr, check to see if they hold equal values
     // (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr
     // equals to 0.
-    auto Check = [this, &CI](const auto *A) -> bool {
-      Expr *XDimExpr = A->getXDim();
-      Expr *YDimExpr = A->getYDim();
-      Expr *ZDimExpr = A->getZDim();
-
-      if (!XDimExpr->isValueDependent() && !YDimExpr->isValueDependent() &&
-          !ZDimExpr->isValueDependent()) {
-        llvm::APSInt XDimVal, YDimVal, ZDimVal;
-        ExprResult XDim = VerifyIntegerConstantExpression(XDimExpr, &XDimVal);
-        ExprResult YDim = VerifyIntegerConstantExpression(YDimExpr, &YDimVal);
-        ExprResult ZDim = VerifyIntegerConstantExpression(ZDimExpr, &ZDimVal);
-
-        if (XDim.isInvalid() || YDim.isInvalid() || ZDim.isInvalid())
-          return false;
-
-        if (XDimVal != 1 || YDimVal != 1 || ZDimVal != 1) {
-          Diag(A->getLocation(), diag::err_sycl_x_y_z_arguments_must_be_one)
-              << A << CI;
-          return false;
-        }
-      }
-      return true;
-    };
-
-    if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
-      if (ArgVal == 0 && !Check(DeclAttr))
+    if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
+      ArgCheck Results[] = {
+        AreArgValuesToOne(E, DeclAttr->getXDim()),
+        AreArgValuesToOne(E, DeclAttr->getYDim()),
+        AreArgValuesToOne(E, DeclAttr->getZDim())};
+
+      // Diagnose if any of the argument values of ReqdWorkGroupSizeAttr
+      // are not equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
+      // equals to 0.
+      if (ArgVal == 0 &&
+	  llvm::any_of(Results,
+                     [](ArgCheck Ch) { return Ch == ArgCheck::NotEqualToOne; })) {
+        Diag(DeclAttr->getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
+            << DeclAttr << CI;
         return;
+      }
     }
 
-    if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
-      if (ArgVal == 0 && !Check(DeclAttr))
+    if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
+      ArgCheck Results[] = {
+        AreArgValuesToOne(E, DeclAttr->getXDim()),
+        AreArgValuesToOne(E, DeclAttr->getYDim()),
+        AreArgValuesToOne(E, DeclAttr->getZDim())};
+
+      // Diagnose if any of the argument values of SYCLIntelMaxWorkGroupSizeAttr
+      // are not equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
+      // equals to 0.
+      if (ArgVal == 0 &&
+	  llvm::any_of(Results,
+                     [](ArgCheck Ch) { return Ch == ArgCheck::NotEqualToOne; })) {
+        Diag(DeclAttr->getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
+            << DeclAttr << CI;
         return;
+      }
     }
   }
 
@@ -3668,49 +3668,37 @@ SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr(
   // (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr
   // equals to 0.
   if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
-    const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getValue());
-    ArgCheck Results[] = {AreArgValuesToOne(
-        DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim())};
+    ArgCheck Results[] = {
+        AreArgValuesToOne(DeclAttr->getXDim(), A.getValue()),
+        AreArgValuesToOne(DeclAttr->getYDim(), A.getValue()),
+        AreArgValuesToOne(DeclAttr->getZDim(), A.getValue())};
 
     // Diagnose if any of the argument values of ReqdWorkGroupSizeAttr
     // are not equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
     // equals to 0.
-    if (MergeExpr->getResultAsAPSInt() == 0 &&
-        llvm::is_contained(Results, ArgCheck::NotEqualToOne)) {
+    if (llvm::any_of(Results,
+                     [](ArgCheck Ch) { return Ch == ArgCheck::NotEqualToOne; })) {
       Diag(DeclAttr->getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
           << DeclAttr << &A;
       return nullptr;
     }
-
-    // Do not diagnose if all argument values of ReqdWorkGroupSizeAttr
-    // are equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
-    // equals to 0.
-    if (llvm::all_of(Results,
-                     [](ArgCheck Ch) { return Ch == ArgCheck::EqualToOne; }))
-      return nullptr;
   }
 
   if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
-    const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getValue());
-    ArgCheck Results[] = {AreArgValuesToOne(
-        DeclAttr->getXDim(), DeclAttr->getYDim(), DeclAttr->getZDim())};
+    ArgCheck Results[] = {
+        AreArgValuesToOne(DeclAttr->getXDim(), A.getValue()),
+        AreArgValuesToOne(DeclAttr->getYDim(), A.getValue()),
+        AreArgValuesToOne(DeclAttr->getZDim(), A.getValue())};
 
     // Diagnose if any of the argument values of SYCLIntelMaxWorkGroupSizeAttr
     // are not equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
     // equals to 0.
-    if (MergeExpr->getResultAsAPSInt() == 0 &&
-        llvm::is_contained(Results, ArgCheck::NotEqualToOne)) {
+    if (llvm::any_of(Results,
+                     [](ArgCheck Ch) { return Ch == ArgCheck::NotEqualToOne; })) {
       Diag(DeclAttr->getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
           << DeclAttr << &A;
       return nullptr;
     }
-
-    // Do not diagnose if all argument values of SYCLIntelMaxWorkGroupSizeAttr
-    // are equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
-    // equals to 0.
-    if (llvm::all_of(Results,
-                     [](ArgCheck Ch) { return Ch == ArgCheck::EqualToOne; }))
-      return nullptr;
   }
 
   return ::new (Context)
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 f48da277bde0b..aa1e387fd0aa0 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
@@ -76,7 +76,7 @@ struct TRIFuncObjGood3 {
 };
 
 struct TRIFuncObjGood4 {
-  [[cl::reqd_work_group_size(1, 1, 1)]]
+  [[sycl::reqd_work_group_size(1, 1, 1)]]
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
@@ -94,7 +94,7 @@ struct TRIFuncObjGood6 {
 };
 
 struct TRIFuncObjGood7 {
-  [[cl::reqd_work_group_size(4, 1, 1)]]
+  [[sycl::reqd_work_group_size(4, 1, 1)]]
   [[intel::max_global_work_dim(3)]] void
   operator()() const {}
 };
@@ -105,6 +105,42 @@ struct TRIFuncObjGood8 {
   operator()() const {}
 };
 
+// FIXME: We do have support yet for MergeExpr of
+// reqd_work_group_size and max_work_group_size
+// attribute, So the test compiles fine now without
+// any diagnostic
+struct TRIFuncObjGood9 {
+  [[intel::max_work_group_size(4, 4, 4)]] void
+  operator()() const;
+};
+
+[[intel::max_global_work_dim(0)]]
+void TRIFuncObjGood9::operator()() const {}
+
+// FIXME: We do have support yet for MergeExpr of
+// reqd_work_group_size and max_work_group_size
+// attribute, So the test compiles fine now without
+// any diagnostic
+struct TRIFuncObjGood10 {
+  [[intel::reqd_work_group_size(4, 4, 4)]] void
+  operator()() const;
+};
+
+[[intel::max_global_work_dim(0)]]
+void TRIFuncObjGood10::operator()() const {}
+
+// FIXME: We do have support yet for MergeExpr of
+// reqd_work_group_size and max_work_group_size
+// attribute, So the test compiles fine now without
+// any diagnostic
+struct TRIFuncObjGood11 {
+  [[sycl::reqd_work_group_size(4, 4, 4)]] void
+  operator()() const;
+};
+
+[[intel::max_global_work_dim(0)]]
+void TRIFuncObjGood11::operator()() const {}
+
 #ifdef TRIGGER_ERROR
 // Checks correctness of mutual usage of different work_group_size attributes:
 // reqd_work_group_size, max_work_group_size and max_global_work_dim.
@@ -127,7 +163,7 @@ struct TRIFuncObjBad1 {
 };
 
 struct TRIFuncObjBad2 {
-  [[cl::reqd_work_group_size(4, 4, 4)]]   // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+  [[sycl::reqd_work_group_size(4, 4, 4)]]   // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
@@ -143,7 +179,7 @@ struct TRIFuncObjBad4 {
   operator()() const;
 };
 
-[[cl::reqd_work_group_size(4, 4, 4)]] // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+[[sycl::reqd_work_group_size(4, 4, 4)]] // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
 void TRIFuncObjBad4::operator()() const {}
 
 struct TRIFuncObjBad5 {
@@ -162,15 +198,6 @@ struct TRIFuncObjBad6 {
 [[intel::max_work_group_size(4, 4, 4)]] // expected-error{{'max_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
 void TRIFuncObjBad6::operator()() const {}
 
-/*
-struct TRIFuncObjBad6 {
-  [[intel::max_work_group_size(4, 4, 4)]] void
-  operator()() const;
-};
-
-[[intel::max_global_work_dim(0)]]
-void TRIFuncObjBad6::operator()() const {}
-*/
 // Tests for incorrect argument values for Intel FPGA function attributes:
 // reqd_work_group_size, max_work_group_size and max_global_work_dim.
 
@@ -189,14 +216,14 @@ struct TRIFuncObjBad8 {
 };
 
 struct TRIFuncObjBad9 {
-  [[cl::reqd_work_group_size(0, 4, 4)]] // expected-error{{'reqd_work_group_size' attribute must be greater than 0}}
+  [[sycl::reqd_work_group_size(0, 4, 4)]] // expected-error{{'reqd_work_group_size' attribute must be greater than 0}}
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
 
 struct TRIFuncObjBad10 {
   [[intel::reqd_work_group_size(4)]]
-  [[intel::max_global_work_dim(-2)]] // expected-error{{'max_global_work_dim' attribute requires a non-negative integral compile time constant expression}}
+  [[intel::max_global_work_dim(-2)]] // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}}
   void operator()() const {}
 };
 
@@ -390,9 +417,60 @@ int main() {
     // CHECK-NEXT:  value: Int 3
     // CHECK-NEXT:  IntegerLiteral{{.*}}3{{$}}
 
+    h.single_task<class test_kernel11>(TRIFuncObjGood9());
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel11
+    // CHECK:       SYCLIntelMaxWorkGroupSizeAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+    // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 0
+    // CHECK-NEXT:  IntegerLiteral{{.*}}0{{$}}
+
+    h.single_task<class test_kernel12>(TRIFuncObjGood10());
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel12
+    // CHECK:       ReqdWorkGroupSizeAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+    // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 0
+    // CHECK-NEXT:  IntegerLiteral{{.*}}0{{$}}
+
+    h.single_task<class test_kernel13>(TRIFuncObjGood11());
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel13
+    // CHECK:       ReqdWorkGroupSizeAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 4
+    // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+    // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
+    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
+    // CHECK-NEXT:  value: Int 0
+    // CHECK-NEXT:  IntegerLiteral{{.*}}0{{$}}
+
     // Ignore duplicate attribute with same argument value.
-    h.single_task<class test_kernell1>(
-    // CHECK-LABEL: FunctionDecl {{.*}}test_kernell1 'void ()'
+    h.single_task<class test_kernell4>(
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernell4 'void ()'
     // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 3
@@ -403,38 +481,38 @@ int main() {
 #ifdef TRIGGER_ERROR
     [[intel::max_global_work_dim(1)]] int Var = 0; // expected-error{{'max_global_work_dim' attribute only applies to functions}}
 
-    h.single_task<class test_kernel12>(
-        []() [[intel::max_global_work_dim(-8)]]{}); // expected-error{{'max_global_work_dim' attribute requires a non-negative integral compile time constant expression}}
+    h.single_task<class test_kernel15>(
+        []() [[intel::max_global_work_dim(-8)]]{}); // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}}
 
-    h.single_task<class test_kernell3>(
+    h.single_task<class test_kernell6>(
         []() [[intel::max_global_work_dim(3),      // expected-note {{previous attribute is here}}
                intel::max_global_work_dim(2)]]{}); // expected-warning{{attribute 'max_global_work_dim' is already applied with different arguments}}
 
-    h.single_task<class test_kernel14>(TRIFuncObjBad());
+    h.single_task<class test_kernel17>(TRIFuncObjBad());
 
-    h.single_task<class test_kernel15>(TRIFuncObjBad1());
+    h.single_task<class test_kernel18>(TRIFuncObjBad1());
 
-    h.single_task<class test_kernel16>(TRIFuncObjBad2());
+    h.single_task<class test_kernel19>(TRIFuncObjBad2());
 
-    h.single_task<class test_kernel17>(TRIFuncObjBad3());
+    h.single_task<class test_kernel20>(TRIFuncObjBad3());
 
-    h.single_task<class test_kernel18>(TRIFuncObjBad4());
+    h.single_task<class test_kernel21>(TRIFuncObjBad4());
 
-    h.single_task<class test_kernel19>(TRIFuncObjBad5());
+    h.single_task<class test_kernel22>(TRIFuncObjBad5());
 
-    h.single_task<class test_kernel20>(TRIFuncObjBad6());
+    h.single_task<class test_kernel23>(TRIFuncObjBad6());
 
-    h.single_task<class test_kernel21>(TRIFuncObjBad7());
+    h.single_task<class test_kernel24>(TRIFuncObjBad7());
 
-    h.single_task<class test_kernel22>(TRIFuncObjBad8());
+    h.single_task<class test_kernel25>(TRIFuncObjBad8());
 
-    h.single_task<class test_kernel23>(TRIFuncObjBad9());
+    h.single_task<class test_kernel26>(TRIFuncObjBad9());
 
-    h.single_task<class test_kernel24>(TRIFuncObjBad10());
+    h.single_task<class test_kernel27>(TRIFuncObjBad10());
 
-    h.single_task<class test_kernel25>(TRIFuncObjBad11());
+    h.single_task<class test_kernel28>(TRIFuncObjBad11());
 
-    h.single_task<class test_kernel26>(
+    h.single_task<class test_kernel29>(
         []() [[intel::max_global_work_dim(4)]]{}); // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}}
 #endif // TRIGGER_ERROR
   });
diff --git a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
index 4b909b9b6d024..56eaf48cecc81 100644
--- a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
+++ b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
@@ -34,7 +34,7 @@ constexpr int bar() { return 0; }
 template <int SIZE>
 class KernelFunctor {
 public:
-  // expected-error@+1{{'max_global_work_dim' attribute requires a non-negative integral compile time constant expression}}
+  // expected-error@+1{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}}
   [[intel::max_global_work_dim(SIZE)]] void operator()() {}
 };
 
@@ -58,7 +58,7 @@ int main() {
 
 // Test that checks template parameter support on function.
 template <int N>
-// expected-error@+1{{'max_global_work_dim' attribute requires a non-negative integral compile time constant expression}}
+// expected-error@+1{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}}
 [[intel::max_global_work_dim(N)]] void func3() {}
 
 // Test that checks template instantiations for different argument values.
@@ -109,7 +109,7 @@ template <int N>
 [[intel::max_global_work_dim(0)]] void func11();
 
 template <int N>
-[[cl::reqd_work_group_size(N, N, N)]] void func12();
+[[sycl::reqd_work_group_size(N, N, N)]] void func12();
 template <int N>
 [[intel::max_global_work_dim(0)]] void func12();
 

From f6c7c576ab05836f27d40fb132273cd39df3e1e7 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Tue, 25 May 2021 06:49:20 -0700
Subject: [PATCH 24/29] fix format issues

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

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 62b2e2d18dd87..f69044fe035e9 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3113,16 +3113,16 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
     // equals to 0.
     if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxGlobalWorkDimAttr>()) {
       if (const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue())) {
-	// If the value is dependent, we can not test anything.
+        // If the value is dependent, we can not test anything.
         if (!DeclExpr)
-	  return;
+          return;
 
-	// Test the attribute value.
+        // Test the attribute value.
         if (DeclExpr->getResultAsAPSInt() == 0 &&
             (XDimVal.getZExtValue() != 1 || YDimVal.getZExtValue() != 1 ||
              ZDimVal.getZExtValue() != 1)) {
           S.Diag(AL.getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
-                << AL << DeclAttr;
+              << AL << DeclAttr;
           return;
         }
       }
@@ -3643,7 +3643,7 @@ static ArgCheck AreArgValuesToOne(const Expr *E, const Expr *E1) {
 
   // Check the argument values.
   return (MaxGExpr->getResultAsAPSInt() == 0 &&
-	  WorkGExpr->getResultAsAPSInt() == 1)
+          WorkGExpr->getResultAsAPSInt() == 1)
              ? ArgCheck::EqualToOne
              : ArgCheck::NotEqualToOne;
 }
@@ -3689,17 +3689,16 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
     // (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr
     // equals to 0.
     if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
-      ArgCheck Results[] = {
-        AreArgValuesToOne(E, DeclAttr->getXDim()),
-        AreArgValuesToOne(E, DeclAttr->getYDim()),
-        AreArgValuesToOne(E, DeclAttr->getZDim())};
+      ArgCheck Results[] = {AreArgValuesToOne(E, DeclAttr->getXDim()),
+                            AreArgValuesToOne(E, DeclAttr->getYDim()),
+                            AreArgValuesToOne(E, DeclAttr->getZDim())};
 
       // Diagnose if any of the argument values of ReqdWorkGroupSizeAttr
       // are not equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
       // equals to 0.
-      if (ArgVal == 0 &&
-	  llvm::any_of(Results,
-                     [](ArgCheck Ch) { return Ch == ArgCheck::NotEqualToOne; })) {
+      if (ArgVal == 0 && llvm::any_of(Results, [](ArgCheck Ch) {
+            return Ch == ArgCheck::NotEqualToOne;
+          })) {
         Diag(DeclAttr->getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
             << DeclAttr << CI;
         return;
@@ -3707,17 +3706,16 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
     }
 
     if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
-      ArgCheck Results[] = {
-        AreArgValuesToOne(E, DeclAttr->getXDim()),
-        AreArgValuesToOne(E, DeclAttr->getYDim()),
-        AreArgValuesToOne(E, DeclAttr->getZDim())};
+      ArgCheck Results[] = {AreArgValuesToOne(E, DeclAttr->getXDim()),
+                            AreArgValuesToOne(E, DeclAttr->getYDim()),
+                            AreArgValuesToOne(E, DeclAttr->getZDim())};
 
       // Diagnose if any of the argument values of SYCLIntelMaxWorkGroupSizeAttr
       // are not equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
       // equals to 0.
-      if (ArgVal == 0 &&
-	  llvm::any_of(Results,
-                     [](ArgCheck Ch) { return Ch == ArgCheck::NotEqualToOne; })) {
+      if (ArgVal == 0 && llvm::any_of(Results, [](ArgCheck Ch) {
+            return Ch == ArgCheck::NotEqualToOne;
+          })) {
         Diag(DeclAttr->getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
             << DeclAttr << CI;
         return;
@@ -3750,16 +3748,16 @@ SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr(
   // (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr
   // equals to 0.
   if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
-    ArgCheck Results[] = {
-        AreArgValuesToOne(DeclAttr->getXDim(), A.getValue()),
-        AreArgValuesToOne(DeclAttr->getYDim(), A.getValue()),
-        AreArgValuesToOne(DeclAttr->getZDim(), A.getValue())};
+    ArgCheck Results[] = {AreArgValuesToOne(DeclAttr->getXDim(), A.getValue()),
+                          AreArgValuesToOne(DeclAttr->getYDim(), A.getValue()),
+                          AreArgValuesToOne(DeclAttr->getZDim(), A.getValue())};
 
     // Diagnose if any of the argument values of ReqdWorkGroupSizeAttr
     // are not equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
     // equals to 0.
-    if (llvm::any_of(Results,
-                     [](ArgCheck Ch) { return Ch == ArgCheck::NotEqualToOne; })) {
+    if (llvm::any_of(Results, [](ArgCheck Ch) {
+          return Ch == ArgCheck::NotEqualToOne;
+        })) {
       Diag(DeclAttr->getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
           << DeclAttr << &A;
       return nullptr;
@@ -3767,16 +3765,16 @@ SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr(
   }
 
   if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
-    ArgCheck Results[] = {
-        AreArgValuesToOne(DeclAttr->getXDim(), A.getValue()),
-        AreArgValuesToOne(DeclAttr->getYDim(), A.getValue()),
-        AreArgValuesToOne(DeclAttr->getZDim(), A.getValue())};
+    ArgCheck Results[] = {AreArgValuesToOne(DeclAttr->getXDim(), A.getValue()),
+                          AreArgValuesToOne(DeclAttr->getYDim(), A.getValue()),
+                          AreArgValuesToOne(DeclAttr->getZDim(), A.getValue())};
 
     // Diagnose if any of the argument values of SYCLIntelMaxWorkGroupSizeAttr
     // are not equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
     // equals to 0.
-    if (llvm::any_of(Results,
-                     [](ArgCheck Ch) { return Ch == ArgCheck::NotEqualToOne; })) {
+    if (llvm::any_of(Results, [](ArgCheck Ch) {
+          return Ch == ArgCheck::NotEqualToOne;
+        })) {
       Diag(DeclAttr->getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
           << DeclAttr << &A;
       return nullptr;

From a1139cd31b85fd1c905ccbebac5125fa7e48d8c8 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Wed, 26 May 2021 08:02:34 -0700
Subject: [PATCH 25/29] address review comments

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

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index f69044fe035e9..7424ee18dc908 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3630,22 +3630,37 @@ static void handleSYCLIntelSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D,
 }
 
 // Handles max_global_work_dim.
-// Returns a ArgCheck value; EqualToOne means all argument values are
+// Returns a OneArgResult value; EqualToOne means all argument values are
 // equal to one, NotEqualToOne means all argument values are not
 // equal to one, and Unknown means that the argument values are invalid.
-enum class ArgCheck { Unknown, EqualToOne, NotEqualToOne };
-static ArgCheck AreArgValuesToOne(const Expr *E, const Expr *E1) {
-  // If any of the expression is value dependent, test is invalid.
-  const auto *MaxGExpr = dyn_cast<ConstantExpr>(E);
-  const auto *WorkGExpr = dyn_cast<ConstantExpr>(E1);
-  if (!MaxGExpr || !WorkGExpr)
-    return ArgCheck::Unknown;
+enum class OneArgResult { Unknown, EqualToOne, NotEqualToOne};
+static OneArgResult AreAllArgsOne(const Expr *Args[], size_t Count) {
 
-  // Check the argument values.
-  return (MaxGExpr->getResultAsAPSInt() == 0 &&
-          WorkGExpr->getResultAsAPSInt() == 1)
-             ? ArgCheck::EqualToOne
-             : ArgCheck::NotEqualToOne;
+  for (size_t Idx = 0; Idx < Count; ++Idx) {
+    const auto *CE = dyn_cast<ConstantExpr>(Args[Idx]);
+    if (!CE)
+      return OneArgResult::Unknown;
+    if (CE->getResultAsAPSInt() != 1)
+      return OneArgResult::NotEqualToOne;
+  }
+  return OneArgResult::EqualToOne;
+}
+
+// If the declaration has a SYCLIntelMaxWorkGroupSizeAttr or
+// ReqdWorkGroupSizeAttr, check to see if they hold equal values
+// (1, 1, 1). Returns true if diagnosed.
+template <typename AttrTy>
+static bool checkWorkGroupSizeAttrExpr(Sema &S, Decl *D,
+                                       const AttributeCommonInfo &AL) {
+  if (const auto *A = D->getAttr<AttrTy>()) {
+    const Expr *Args[3] = {A->getXDim(), A->getYDim(), A->getZDim()};
+    if (OneArgResult::NotEqualToOne == AreAllArgsOne(Args, 3)) {
+      S.Diag(A->getLocation(), diag::err_sycl_x_y_z_arguments_must_be_one)
+            << A << AL;
+      return true;
+    }
+  }
+  return false;
 }
 
 void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
@@ -3688,38 +3703,11 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
     // ReqdWorkGroupSizeAttr, check to see if they hold equal values
     // (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr
     // equals to 0.
-    if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
-      ArgCheck Results[] = {AreArgValuesToOne(E, DeclAttr->getXDim()),
-                            AreArgValuesToOne(E, DeclAttr->getYDim()),
-                            AreArgValuesToOne(E, DeclAttr->getZDim())};
-
-      // Diagnose if any of the argument values of ReqdWorkGroupSizeAttr
-      // are not equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
-      // equals to 0.
-      if (ArgVal == 0 && llvm::any_of(Results, [](ArgCheck Ch) {
-            return Ch == ArgCheck::NotEqualToOne;
-          })) {
-        Diag(DeclAttr->getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
-            << DeclAttr << CI;
+    if (ArgVal == 0) {
+      if (checkWorkGroupSizeAttrExpr<SYCLIntelMaxWorkGroupSizeAttr>(*this, D, CI))
         return;
-      }
-    }
-
-    if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
-      ArgCheck Results[] = {AreArgValuesToOne(E, DeclAttr->getXDim()),
-                            AreArgValuesToOne(E, DeclAttr->getYDim()),
-                            AreArgValuesToOne(E, DeclAttr->getZDim())};
-
-      // Diagnose if any of the argument values of SYCLIntelMaxWorkGroupSizeAttr
-      // are not equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
-      // equals to 0.
-      if (ArgVal == 0 && llvm::any_of(Results, [](ArgCheck Ch) {
-            return Ch == ArgCheck::NotEqualToOne;
-          })) {
-        Diag(DeclAttr->getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
-            << DeclAttr << CI;
+      if (checkWorkGroupSizeAttrExpr<ReqdWorkGroupSizeAttr>(*this, D, CI))
         return;
-      }
     }
   }
 
@@ -3747,38 +3735,12 @@ SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr(
   // ReqdWorkGroupSizeAttr, check to see if they hold equal values
   // (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr
   // equals to 0.
-  if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
-    ArgCheck Results[] = {AreArgValuesToOne(DeclAttr->getXDim(), A.getValue()),
-                          AreArgValuesToOne(DeclAttr->getYDim(), A.getValue()),
-                          AreArgValuesToOne(DeclAttr->getZDim(), A.getValue())};
-
-    // Diagnose if any of the argument values of ReqdWorkGroupSizeAttr
-    // are not equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
-    // equals to 0.
-    if (llvm::any_of(Results, [](ArgCheck Ch) {
-          return Ch == ArgCheck::NotEqualToOne;
-        })) {
-      Diag(DeclAttr->getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
-          << DeclAttr << &A;
+  const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getValue());
+  if (MergeExpr->getResultAsAPSInt() == 0) {
+    if (checkWorkGroupSizeAttrExpr<SYCLIntelMaxWorkGroupSizeAttr>(*this, D, A))
       return nullptr;
-    }
-  }
-
-  if (const auto *DeclAttr = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
-    ArgCheck Results[] = {AreArgValuesToOne(DeclAttr->getXDim(), A.getValue()),
-                          AreArgValuesToOne(DeclAttr->getYDim(), A.getValue()),
-                          AreArgValuesToOne(DeclAttr->getZDim(), A.getValue())};
-
-    // Diagnose if any of the argument values of SYCLIntelMaxWorkGroupSizeAttr
-    // are not equal to one when value of SYCLIntelMaxGlobalWorkDimAttr
-    // equals to 0.
-    if (llvm::any_of(Results, [](ArgCheck Ch) {
-          return Ch == ArgCheck::NotEqualToOne;
-        })) {
-      Diag(DeclAttr->getLoc(), diag::err_sycl_x_y_z_arguments_must_be_one)
-          << DeclAttr << &A;
+    if (checkWorkGroupSizeAttrExpr<ReqdWorkGroupSizeAttr>(*this, D, A))
       return nullptr;
-    }
   }
 
   return ::new (Context)

From 45a294b051a6a6412a1fa6f47ba194e2b4836928 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Wed, 26 May 2021 08:07:49 -0700
Subject: [PATCH 26/29] fix format issues

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

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 7424ee18dc908..212f52101fa39 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3633,7 +3633,7 @@ static void handleSYCLIntelSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D,
 // Returns a OneArgResult value; EqualToOne means all argument values are
 // equal to one, NotEqualToOne means all argument values are not
 // equal to one, and Unknown means that the argument values are invalid.
-enum class OneArgResult { Unknown, EqualToOne, NotEqualToOne};
+enum class OneArgResult { Unknown, EqualToOne, NotEqualToOne };
 static OneArgResult AreAllArgsOne(const Expr *Args[], size_t Count) {
 
   for (size_t Idx = 0; Idx < Count; ++Idx) {
@@ -3656,7 +3656,7 @@ static bool checkWorkGroupSizeAttrExpr(Sema &S, Decl *D,
     const Expr *Args[3] = {A->getXDim(), A->getYDim(), A->getZDim()};
     if (OneArgResult::NotEqualToOne == AreAllArgsOne(Args, 3)) {
       S.Diag(A->getLocation(), diag::err_sycl_x_y_z_arguments_must_be_one)
-            << A << AL;
+          << A << AL;
       return true;
     }
   }
@@ -3704,7 +3704,8 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
     // (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr
     // equals to 0.
     if (ArgVal == 0) {
-      if (checkWorkGroupSizeAttrExpr<SYCLIntelMaxWorkGroupSizeAttr>(*this, D, CI))
+      if (checkWorkGroupSizeAttrExpr<SYCLIntelMaxWorkGroupSizeAttr>(*this, D,
+                                                                    CI))
         return;
       if (checkWorkGroupSizeAttrExpr<ReqdWorkGroupSizeAttr>(*this, D, CI))
         return;

From 65da4b8dedb3a7ae375ebf31d337cc7cb2cc7a4c Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Wed, 26 May 2021 10:07:55 -0700
Subject: [PATCH 27/29] address review comments

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 .../clang/Basic/DiagnosticSemaKinds.td        |   2 +-
 clang/lib/Sema/SemaDeclAttr.cpp               |  15 ++-
 .../intel-max-global-work-dim-device.cpp      | 117 +++++++++---------
 ...ice-intel-max-global-work-dim-template.cpp |   8 +-
 4 files changed, 69 insertions(+), 73 deletions(-)

diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 769d170880362..84e5a96bc97b0 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -11346,7 +11346,7 @@ def err_conflicting_sycl_function_attributes : Error<
 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<
-  "%0 X-, Y- and Z- sizes must be 1 when %1 attribute is used with value 0">;
+  "%0 X, Y and Z sizes must be 1 when %1 attribute is used with value 0">;
 def err_sycl_attribute_internal_function
     : Error<"%0 attribute cannot be applied to a "
             "static function or function in an anonymous namespace">;
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 212f52101fa39..674c10d713a67 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3631,8 +3631,9 @@ static void handleSYCLIntelSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D,
 
 // Handles max_global_work_dim.
 // Returns a OneArgResult value; EqualToOne means all argument values are
-// equal to one, NotEqualToOne means all argument values are not
-// equal to one, and Unknown means that the argument values are invalid.
+// equal to one, NotEqualToOne means at least one argument value is not
+// equal to one, and Unknown means that at least one of the argument values
+// could not be determined.
 enum class OneArgResult { Unknown, EqualToOne, NotEqualToOne };
 static OneArgResult AreAllArgsOne(const Expr *Args[], size_t Count) {
 
@@ -3705,9 +3706,8 @@ void Sema::AddSYCLIntelMaxGlobalWorkDimAttr(Decl *D,
     // equals to 0.
     if (ArgVal == 0) {
       if (checkWorkGroupSizeAttrExpr<SYCLIntelMaxWorkGroupSizeAttr>(*this, D,
-                                                                    CI))
-        return;
-      if (checkWorkGroupSizeAttrExpr<ReqdWorkGroupSizeAttr>(*this, D, CI))
+                                                                    CI) ||
+          checkWorkGroupSizeAttrExpr<ReqdWorkGroupSizeAttr>(*this, D, CI))
         return;
     }
   }
@@ -3738,9 +3738,8 @@ SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr(
   // equals to 0.
   const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getValue());
   if (MergeExpr->getResultAsAPSInt() == 0) {
-    if (checkWorkGroupSizeAttrExpr<SYCLIntelMaxWorkGroupSizeAttr>(*this, D, A))
-      return nullptr;
-    if (checkWorkGroupSizeAttrExpr<ReqdWorkGroupSizeAttr>(*this, D, A))
+    if (checkWorkGroupSizeAttrExpr<SYCLIntelMaxWorkGroupSizeAttr>(*this, D, A) ||
+        checkWorkGroupSizeAttrExpr<ReqdWorkGroupSizeAttr>(*this, D, A))
       return nullptr;
   }
 
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 aa1e387fd0aa0..70f89b9226a4f 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
@@ -37,9 +37,6 @@ struct Func {
 [[intel::max_global_work_dim(1)]] void bar();
 [[intel::max_global_work_dim(1)]] void bar() {}
 
-[[intel::max_global_work_dim(2)]] void func_ignore();
-[[intel::max_global_work_dim(2)]] void func_ignore();
-
 // Checking of different argument values.
 [[intel::max_global_work_dim(2)]] void baz();  // expected-note {{previous attribute is here}}
 [[intel::max_global_work_dim(1)]] void baz();  // expected-warning {{attribute 'max_global_work_dim' is already applied with different arguments}}
@@ -105,41 +102,41 @@ struct TRIFuncObjGood8 {
   operator()() const {}
 };
 
-// FIXME: We do have support yet for MergeExpr of
-// reqd_work_group_size and max_work_group_size
-// attribute, So the test compiles fine now without
-// any diagnostic
-struct TRIFuncObjGood9 {
+// FIXME: We do not have support yet for checking
+// max_work_group_size and max_global_work_dim
+// attributes when merging, so the test compiles without
+// any diagnostic when it shouldn't.
+struct TRIFuncObjBad {
   [[intel::max_work_group_size(4, 4, 4)]] void
   operator()() const;
 };
 
 [[intel::max_global_work_dim(0)]]
-void TRIFuncObjGood9::operator()() const {}
+void TRIFuncObjBad::operator()() const {}
 
-// FIXME: We do have support yet for MergeExpr of
-// reqd_work_group_size and max_work_group_size
-// attribute, So the test compiles fine now without
-// any diagnostic
-struct TRIFuncObjGood10 {
+// FIXME: We do not have support yet for checking
+// max_work_group_size and max_global_work_dim
+// attributes when merging, so the test compiles without
+// any diagnostic when it shouldn't.
+struct TRIFuncObjBad1 {
   [[intel::reqd_work_group_size(4, 4, 4)]] void
   operator()() const;
 };
 
 [[intel::max_global_work_dim(0)]]
-void TRIFuncObjGood10::operator()() const {}
+void TRIFuncObjBad1::operator()() const {}
 
-// FIXME: We do have support yet for MergeExpr of
-// reqd_work_group_size and max_work_group_size
-// attribute, So the test compiles fine now without
-// any diagnostic
-struct TRIFuncObjGood11 {
+// FIXME: We do not have support yet for checking
+// max_work_group_size and max_global_work_dim
+// attributes when merging, so the test compiles without
+// any diagnostic when it shouldn't.
+struct TRIFuncObjBad2 {
   [[sycl::reqd_work_group_size(4, 4, 4)]] void
   operator()() const;
 };
 
 [[intel::max_global_work_dim(0)]]
-void TRIFuncObjGood11::operator()() const {}
+void TRIFuncObjBad2::operator()() const {}
 
 #ifdef TRIGGER_ERROR
 // Checks correctness of mutual usage of different work_group_size attributes:
@@ -148,86 +145,86 @@ void TRIFuncObjGood11::operator()() const {}
 // ensure that if max_work_group_size and reqd_work_group_size attributes exist,
 // they hold equal values (1, 1, 1).
 
-struct TRIFuncObjBad {
+struct TRIFuncObjBad3 {
   [[intel::max_global_work_dim(0)]]
-  [[intel::max_work_group_size(8, 8, 8)]] // expected-error{{'max_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
-  [[sycl::reqd_work_group_size(4, 4, 4)]]   // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+  [[intel::max_work_group_size(8, 8, 8)]] // expected-error{{'max_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+  [[sycl::reqd_work_group_size(4, 4, 4)]] // expected-error{{'reqd_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
   void
   operator()() const {}
 };
 
-struct TRIFuncObjBad1 {
-  [[intel::max_work_group_size(8, 8, 8)]] // expected-error{{'max_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+struct TRIFuncObjBad4 {
+  [[intel::max_work_group_size(8, 8, 8)]] // expected-error{{'max_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
 
-struct TRIFuncObjBad2 {
-  [[sycl::reqd_work_group_size(4, 4, 4)]]   // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+struct TRIFuncObjBad5 {
+  [[sycl::reqd_work_group_size(4, 4, 4)]]   // expected-error{{'reqd_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
 
-struct TRIFuncObjBad3 {
-  [[intel::reqd_work_group_size(4)]]   // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+struct TRIFuncObjBad6 {
+  [[intel::reqd_work_group_size(4)]]   // expected-error{{'reqd_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
 
-struct TRIFuncObjBad4 {
+struct TRIFuncObjBad7 {
   [[intel::max_global_work_dim(0)]] void
   operator()() const;
 };
 
-[[sycl::reqd_work_group_size(4, 4, 4)]] // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
-void TRIFuncObjBad4::operator()() const {}
+[[sycl::reqd_work_group_size(4, 4, 4)]] // expected-error{{'reqd_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+void TRIFuncObjBad7::operator()() const {}
 
-struct TRIFuncObjBad5 {
+struct TRIFuncObjBad8 {
   [[intel::max_global_work_dim(0)]] void
   operator()() const;
 };
 
-[[intel::reqd_work_group_size(4, 4, 4)]] // expected-error{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
-void TRIFuncObjBad5::operator()() const {}
+[[intel::reqd_work_group_size(4, 4, 4)]] // expected-error{{'reqd_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+void TRIFuncObjBad8::operator()() const {}
 
-struct TRIFuncObjBad6 {
+struct TRIFuncObjBad9 {
   [[intel::max_global_work_dim(0)]] void
   operator()() const;
 };
 
-[[intel::max_work_group_size(4, 4, 4)]] // expected-error{{'max_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
-void TRIFuncObjBad6::operator()() const {}
+[[intel::max_work_group_size(4, 4, 4)]] // expected-error{{'max_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+void TRIFuncObjBad9::operator()() const {}
 
 // Tests for incorrect argument values for Intel FPGA function attributes:
 // reqd_work_group_size, max_work_group_size and max_global_work_dim.
 
-struct TRIFuncObjBad7 {
-  // expected-error@+2{{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+struct TRIFuncObjBad10 {
+  // expected-error@+2{{'reqd_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
   // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}}
   [[intel::reqd_work_group_size(-4, 1)]]
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
 
-struct TRIFuncObjBad8 {
+struct TRIFuncObjBad11 {
   [[intel::max_work_group_size(4, 4, 4.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}}
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
 
-struct TRIFuncObjBad9 {
+struct TRIFuncObjBad12 {
   [[sycl::reqd_work_group_size(0, 4, 4)]] // expected-error{{'reqd_work_group_size' attribute must be greater than 0}}
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
 
-struct TRIFuncObjBad10 {
+struct TRIFuncObjBad13 {
   [[intel::reqd_work_group_size(4)]]
   [[intel::max_global_work_dim(-2)]] // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}}
   void operator()() const {}
 };
 
-struct TRIFuncObjBad11 {
+struct TRIFuncObjBad14 {
   [[intel::max_work_group_size(4, 4, 4)]]
   [[intel::max_global_work_dim(4.f)]] // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'float'}}
   void operator()() const {}
@@ -417,7 +414,7 @@ int main() {
     // CHECK-NEXT:  value: Int 3
     // CHECK-NEXT:  IntegerLiteral{{.*}}3{{$}}
 
-    h.single_task<class test_kernel11>(TRIFuncObjGood9());
+    h.single_task<class test_kernel11>(TRIFuncObjBad());
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel11
     // CHECK:       SYCLIntelMaxWorkGroupSizeAttr {{.*}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
@@ -434,7 +431,7 @@ int main() {
     // CHECK-NEXT:  value: Int 0
     // CHECK-NEXT:  IntegerLiteral{{.*}}0{{$}}
 
-    h.single_task<class test_kernel12>(TRIFuncObjGood10());
+    h.single_task<class test_kernel12>(TRIFuncObjBad1());
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel12
     // CHECK:       ReqdWorkGroupSizeAttr {{.*}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
@@ -451,7 +448,7 @@ int main() {
     // CHECK-NEXT:  value: Int 0
     // CHECK-NEXT:  IntegerLiteral{{.*}}0{{$}}
 
-    h.single_task<class test_kernel13>(TRIFuncObjGood11());
+    h.single_task<class test_kernel13>(TRIFuncObjBad2());
     // CHECK-LABEL: FunctionDecl {{.*}}test_kernel13
     // CHECK:       ReqdWorkGroupSizeAttr {{.*}}
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
@@ -488,29 +485,29 @@ int main() {
         []() [[intel::max_global_work_dim(3),      // expected-note {{previous attribute is here}}
                intel::max_global_work_dim(2)]]{}); // expected-warning{{attribute 'max_global_work_dim' is already applied with different arguments}}
 
-    h.single_task<class test_kernel17>(TRIFuncObjBad());
+    h.single_task<class test_kernel17>(TRIFuncObjBad3());
 
-    h.single_task<class test_kernel18>(TRIFuncObjBad1());
+    h.single_task<class test_kernel18>(TRIFuncObjBad4());
 
-    h.single_task<class test_kernel19>(TRIFuncObjBad2());
+    h.single_task<class test_kernel19>(TRIFuncObjBad5());
 
-    h.single_task<class test_kernel20>(TRIFuncObjBad3());
+    h.single_task<class test_kernel20>(TRIFuncObjBad6());
 
-    h.single_task<class test_kernel21>(TRIFuncObjBad4());
+    h.single_task<class test_kernel21>(TRIFuncObjBad7());
 
-    h.single_task<class test_kernel22>(TRIFuncObjBad5());
+    h.single_task<class test_kernel22>(TRIFuncObjBad8());
 
-    h.single_task<class test_kernel23>(TRIFuncObjBad6());
+    h.single_task<class test_kernel23>(TRIFuncObjBad9());
 
-    h.single_task<class test_kernel24>(TRIFuncObjBad7());
+    h.single_task<class test_kernel24>(TRIFuncObjBad10());
 
-    h.single_task<class test_kernel25>(TRIFuncObjBad8());
+    h.single_task<class test_kernel25>(TRIFuncObjBad11());
 
-    h.single_task<class test_kernel26>(TRIFuncObjBad9());
+    h.single_task<class test_kernel26>(TRIFuncObjBad12());
 
-    h.single_task<class test_kernel27>(TRIFuncObjBad10());
+    h.single_task<class test_kernel27>(TRIFuncObjBad13());
 
-    h.single_task<class test_kernel28>(TRIFuncObjBad11());
+    h.single_task<class test_kernel28>(TRIFuncObjBad14());
 
     h.single_task<class test_kernel29>(
         []() [[intel::max_global_work_dim(4)]]{}); // expected-error{{'max_global_work_dim' attribute requires integer constant between 0 and 3 inclusive}}
diff --git a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
index 56eaf48cecc81..02813cce51f4d 100644
--- a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
+++ b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
@@ -74,22 +74,22 @@ template <int size>
 // ensure that if max_work_group_size and reqd_work_group_size attributes exist,
 // they hold equal values (1, 1, 1).
 template <int N>
-[[intel::max_work_group_size(N, N, N)]] void func5(); // expected-error {{'max_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+[[intel::max_work_group_size(N, N, N)]] void func5(); // expected-error {{'max_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
 template <int N>
 [[intel::max_global_work_dim(0)]] void func5();
 
 template <int N>
-[[intel::reqd_work_group_size(N)]] void func6(); // expected-error {{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+[[intel::reqd_work_group_size(N)]] void func6(); // expected-error {{'reqd_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
 template <int N>
 [[intel::max_global_work_dim(0)]] void func6();
 
 template <int N>
-[[intel::reqd_work_group_size(N, N)]] void func7(); // expected-error {{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+[[intel::reqd_work_group_size(N, N)]] void func7(); // expected-error {{'reqd_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
 template <int N>
 [[intel::max_global_work_dim(0)]] void func7();
 
 template <int N>
-[[intel::reqd_work_group_size(N, N, N)]] void func8(); // expected-error {{'reqd_work_group_size' X-, Y- and Z- sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+[[intel::reqd_work_group_size(N, N, N)]] void func8(); // expected-error {{'reqd_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
 template <int N>
 [[intel::max_global_work_dim(0)]] void func8();
 

From 5d5f1faebb1a3ceb5c0561645560d7ab6b048ece Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Wed, 26 May 2021 10:11:35 -0700
Subject: [PATCH 28/29] fix format issues

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

diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 674c10d713a67..f310567cb8f08 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3738,7 +3738,8 @@ SYCLIntelMaxGlobalWorkDimAttr *Sema::MergeSYCLIntelMaxGlobalWorkDimAttr(
   // equals to 0.
   const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getValue());
   if (MergeExpr->getResultAsAPSInt() == 0) {
-    if (checkWorkGroupSizeAttrExpr<SYCLIntelMaxWorkGroupSizeAttr>(*this, D, A) ||
+    if (checkWorkGroupSizeAttrExpr<SYCLIntelMaxWorkGroupSizeAttr>(*this, D,
+                                                                  A) ||
         checkWorkGroupSizeAttrExpr<ReqdWorkGroupSizeAttr>(*this, D, A))
       return nullptr;
   }

From 618bb3ea704631fc4d82415fa46494dc008b4787 Mon Sep 17 00:00:00 2001
From: Soumi Manna <soumi.manna@intel.com>
Date: Wed, 26 May 2021 14:28:39 -0700
Subject: [PATCH 29/29] address review comments

Signed-off-by: Soumi Manna <soumi.manna@intel.com>
---
 .../clang/Basic/DiagnosticSemaKinds.td        |  2 +-
 .../intel-max-global-work-dim-device.cpp      | 29 ++++++-------------
 ...ice-intel-max-global-work-dim-template.cpp |  8 ++---
 3 files changed, 14 insertions(+), 25 deletions(-)

diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 84e5a96bc97b0..03e86a36bf475 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -11346,7 +11346,7 @@ def err_conflicting_sycl_function_attributes : Error<
 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<
-  "%0 X, Y and Z sizes must be 1 when %1 attribute is used with value 0">;
+  "all %0 attribute arguments must be '1' when the %1 attribute argument is '0'">;
 def err_sycl_attribute_internal_function
     : Error<"%0 attribute cannot be applied to a "
             "static function or function in an anonymous namespace">;
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 70f89b9226a4f..29cc7f084e2d4 100644
--- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
+++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp
@@ -147,26 +147,26 @@ void TRIFuncObjBad2::operator()() const {}
 
 struct TRIFuncObjBad3 {
   [[intel::max_global_work_dim(0)]]
-  [[intel::max_work_group_size(8, 8, 8)]] // expected-error{{'max_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
-  [[sycl::reqd_work_group_size(4, 4, 4)]] // expected-error{{'reqd_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+  [[intel::max_work_group_size(8, 8, 8)]] // expected-error{{all 'max_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}}
+  [[sycl::reqd_work_group_size(4, 4, 4)]] // expected-error{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}}
   void
   operator()() const {}
 };
 
 struct TRIFuncObjBad4 {
-  [[intel::max_work_group_size(8, 8, 8)]] // expected-error{{'max_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+  [[intel::max_work_group_size(8, 8, 8)]] // expected-error{{all 'max_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}}
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
 
 struct TRIFuncObjBad5 {
-  [[sycl::reqd_work_group_size(4, 4, 4)]]   // expected-error{{'reqd_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+  [[sycl::reqd_work_group_size(4, 4, 4)]]   // expected-error{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}}
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
 
 struct TRIFuncObjBad6 {
-  [[intel::reqd_work_group_size(4)]]   // expected-error{{'reqd_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+  [[intel::reqd_work_group_size(4)]]   // expected-error{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}}
   [[intel::max_global_work_dim(0)]] void
   operator()() const {}
 };
@@ -176,7 +176,7 @@ struct TRIFuncObjBad7 {
   operator()() const;
 };
 
-[[sycl::reqd_work_group_size(4, 4, 4)]] // expected-error{{'reqd_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+[[sycl::reqd_work_group_size(4, 4, 4)]] // expected-error{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}}
 void TRIFuncObjBad7::operator()() const {}
 
 struct TRIFuncObjBad8 {
@@ -184,7 +184,7 @@ struct TRIFuncObjBad8 {
   operator()() const;
 };
 
-[[intel::reqd_work_group_size(4, 4, 4)]] // expected-error{{'reqd_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+[[intel::reqd_work_group_size(4, 4, 4)]] // expected-error{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}}
 void TRIFuncObjBad8::operator()() const {}
 
 struct TRIFuncObjBad9 {
@@ -192,14 +192,14 @@ struct TRIFuncObjBad9 {
   operator()() const;
 };
 
-[[intel::max_work_group_size(4, 4, 4)]] // expected-error{{'max_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+[[intel::max_work_group_size(4, 4, 4)]] // expected-error{{all 'max_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}}
 void TRIFuncObjBad9::operator()() const {}
 
 // Tests for incorrect argument values for Intel FPGA function attributes:
 // reqd_work_group_size, max_work_group_size and max_global_work_dim.
 
 struct TRIFuncObjBad10 {
-  // expected-error@+2{{'reqd_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+  // expected-error@+2{{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}}
   // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}}
   [[intel::reqd_work_group_size(-4, 1)]]
   [[intel::max_global_work_dim(0)]] void
@@ -486,27 +486,16 @@ int main() {
                intel::max_global_work_dim(2)]]{}); // expected-warning{{attribute 'max_global_work_dim' is already applied with different arguments}}
 
     h.single_task<class test_kernel17>(TRIFuncObjBad3());
-
     h.single_task<class test_kernel18>(TRIFuncObjBad4());
-
     h.single_task<class test_kernel19>(TRIFuncObjBad5());
-
     h.single_task<class test_kernel20>(TRIFuncObjBad6());
-
     h.single_task<class test_kernel21>(TRIFuncObjBad7());
-
     h.single_task<class test_kernel22>(TRIFuncObjBad8());
-
     h.single_task<class test_kernel23>(TRIFuncObjBad9());
-
     h.single_task<class test_kernel24>(TRIFuncObjBad10());
-
     h.single_task<class test_kernel25>(TRIFuncObjBad11());
-
     h.single_task<class test_kernel26>(TRIFuncObjBad12());
-
     h.single_task<class test_kernel27>(TRIFuncObjBad13());
-
     h.single_task<class test_kernel28>(TRIFuncObjBad14());
 
     h.single_task<class test_kernel29>(
diff --git a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
index 02813cce51f4d..41bb2ce80dde5 100644
--- a/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
+++ b/clang/test/SemaSYCL/sycl-device-intel-max-global-work-dim-template.cpp
@@ -74,22 +74,22 @@ template <int size>
 // ensure that if max_work_group_size and reqd_work_group_size attributes exist,
 // they hold equal values (1, 1, 1).
 template <int N>
-[[intel::max_work_group_size(N, N, N)]] void func5(); // expected-error {{'max_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+[[intel::max_work_group_size(N, N, N)]] void func5(); // expected-error {{all 'max_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}}
 template <int N>
 [[intel::max_global_work_dim(0)]] void func5();
 
 template <int N>
-[[intel::reqd_work_group_size(N)]] void func6(); // expected-error {{'reqd_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+[[intel::reqd_work_group_size(N)]] void func6(); // expected-error {{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}}
 template <int N>
 [[intel::max_global_work_dim(0)]] void func6();
 
 template <int N>
-[[intel::reqd_work_group_size(N, N)]] void func7(); // expected-error {{'reqd_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+[[intel::reqd_work_group_size(N, N)]] void func7(); // expected-error {{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}}
 template <int N>
 [[intel::max_global_work_dim(0)]] void func7();
 
 template <int N>
-[[intel::reqd_work_group_size(N, N, N)]] void func8(); // expected-error {{'reqd_work_group_size' X, Y and Z sizes must be 1 when 'max_global_work_dim' attribute is used with value 0}}
+[[intel::reqd_work_group_size(N, N, N)]] void func8(); // expected-error {{all 'reqd_work_group_size' attribute arguments must be '1' when the 'max_global_work_dim' attribute argument is '0'}}
 template <int N>
 [[intel::max_global_work_dim(0)]] void func8();