[SYCL] Allow [[sycl::work_group_size_hint]] to accept constant expr args

clang/include/clang/Basic/Attr.td

@@ -2975,10 +2975,27 @@ def ReqdWorkGroupSize : InheritableAttr {
 def WorkGroupSizeHint :  InheritableAttr {
   let Spellings = [GNU<"work_group_size_hint">,
                    CXX11<"sycl", "work_group_size_hint">];
-  let Args = [UnsignedArgument<"XDim">,
-              UnsignedArgument<"YDim">,
-              UnsignedArgument<"ZDim">];
+  let Args = [ExprArgument<"XDim">,
+              ExprArgument<"YDim">,
+              ExprArgument<"ZDim">];
   let Subjects = SubjectList<[Function], ErrorDiag>;
+  let AdditionalMembers = [{
+    Optional<llvm::APSInt> getXDimVal() const {
+      if (const auto *CE = dyn_cast<ConstantExpr>(getXDim()))
+        return CE->getResultAsAPSInt();
+      return None;
+    }
+    Optional<llvm::APSInt> getYDimVal() const {
+      if (const auto *CE = dyn_cast<ConstantExpr>(getYDim()))
+        return CE->getResultAsAPSInt();
+      return None;
+    }
+    Optional<llvm::APSInt> getZDimVal() const {
+      if (const auto *CE = dyn_cast<ConstantExpr>(getZDim()))
+        return CE->getResultAsAPSInt();
+      return None;
+    }
+  }];
   let Documentation = [WorkGroupSizeHintAttrDocs];
 }
 

clang/include/clang/Sema/Sema.h

@@ -10306,6 +10306,10 @@ class Sema final {
   template <typename AttrType>
   void addIntelTripleArgAttr(Decl *D, const AttributeCommonInfo &CI,
                              Expr *XDimExpr, Expr *YDimExpr, Expr *ZDimExpr);
+  void AddWorkGroupSizeHintAttr(Decl *D, const AttributeCommonInfo &CI,
+                                Expr *XDim, Expr *YDim, Expr *ZDim);
+  WorkGroupSizeHintAttr *
+  MergeWorkGroupSizeHintAttr(Decl *D, const WorkGroupSizeHintAttr &A);
   void AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI,
                                 Expr *E);
   IntelReqdSubGroupSizeAttr *

clang/lib/CodeGen/CodeGenFunction.cpp

@@ -622,9 +622,9 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
 
   if (const WorkGroupSizeHintAttr *A = FD->getAttr<WorkGroupSizeHintAttr>()) {
     llvm::Metadata *AttrMDArgs[] = {
-        llvm::ConstantAsMetadata::get(Builder.getInt32(A->getXDim())),
-        llvm::ConstantAsMetadata::get(Builder.getInt32(A->getYDim())),
-        llvm::ConstantAsMetadata::get(Builder.getInt32(A->getZDim()))};
+        llvm::ConstantAsMetadata::get(Builder.getInt(*A->getXDimVal())),
+        llvm::ConstantAsMetadata::get(Builder.getInt(*A->getYDimVal())),
+        llvm::ConstantAsMetadata::get(Builder.getInt(*A->getZDimVal()))};
     Fn->setMetadata("work_group_size_hint", llvm::MDNode::get(Context, AttrMDArgs));
   }
 

clang/lib/Sema/SemaDecl.cpp

@@ -2641,6 +2641,8 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D,
     NewAttr = S.MergeIntelFPGAForcePow2DepthAttr(D, *A);
   else if (const auto *A = dyn_cast<SYCLIntelFPGAInitiationIntervalAttr>(Attr))
     NewAttr = S.MergeSYCLIntelFPGAInitiationIntervalAttr(D, *A);
+  else if (const auto *A = dyn_cast<WorkGroupSizeHintAttr>(Attr))
+    NewAttr = S.MergeWorkGroupSizeHintAttr(D, *A);
   else if (Attr->shouldInheritEvenIfAlreadyPresent() || !DeclHasAttr(D, Attr))
     NewAttr = cast<InheritableAttr>(Attr->clone(S.Context));
 

clang/lib/Sema/SemaDeclAttr.cpp

@@ -3160,31 +3160,113 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
   S.addIntelTripleArgAttr<WorkGroupAttr>(D, AL, XDimExpr, YDimExpr, ZDimExpr);
 }
 
-// Handles work_group_size_hint.
-static void handleWorkGroupSizeHint(Sema &S, Decl *D, const ParsedAttr &AL) {
-  S.CheckDeprecatedSYCLAttributeSpelling(AL);
+// Returns a DupArgResult value; Same means the args have the same value,
+// Different means the args do not have the same value, and Unknown means that
+// the args cannot (yet) be compared.
+enum class DupArgResult { Unknown, Same, Different };
+static DupArgResult AreArgValuesIdentical(const Expr *LHS, const Expr *RHS) {
+  // If either operand is still value dependent, we can't test anything.
+  const auto *LHSCE = dyn_cast<ConstantExpr>(LHS);
+  const auto *RHSCE = dyn_cast<ConstantExpr>(RHS);
+  if (!LHSCE || !RHSCE)
+    return DupArgResult::Unknown;
+
+  // Otherwise, test that the values.
+  return LHSCE->getResultAsAPSInt() == RHSCE->getResultAsAPSInt()
+             ? DupArgResult::Same
+             : DupArgResult::Different;
+}
+
+void Sema::AddWorkGroupSizeHintAttr(Decl *D, const AttributeCommonInfo &CI,
+                                    Expr *XDim, Expr *YDim, Expr *ZDim) {
+  // Returns nullptr if diagnosing, otherwise returns the original expression
+  // or the original expression converted to a constant expression.
+  auto CheckAndConvertArg = [&](Expr *E) -> Expr * {
+    // We can only check if the expression is not value dependent.
+    if (!E->isValueDependent()) {
+      llvm::APSInt ArgVal;
+      ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal);
+      if (Res.isInvalid())
+        return nullptr;
+      E = Res.get();
 
-  uint32_t WGSize[3];
-  for (unsigned i = 0; i < AL.getNumArgs(); ++i) {
-    if (!checkUInt32Argument(S, AL, AL.getArgAsExpr(i), WGSize[i], i,
-                             /*StrictlyUnsigned=*/true))
-      return;
+      // This attribute requires a strictly positive value.
+      if (ArgVal <= 0) {
+        Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer)
+            << CI << /*positive*/ 0;
+        return nullptr;
+      }
+    }
+
+    return E;
+  };
 
-    if (WGSize[i] == 0) {
-      S.Diag(AL.getLoc(), diag::err_attribute_argument_is_zero)
-          << AL << AL.getArgAsExpr(i)->getSourceRange();
+  // Check all three argument values, and if any are bad, bail out. This will
+  // convert the given expressions into constant expressions when possible.
+  XDim = CheckAndConvertArg(XDim);
+  YDim = CheckAndConvertArg(YDim);
+  ZDim = CheckAndConvertArg(ZDim);
+  if (!XDim || !YDim || !ZDim)
+    return;
+
+  // If the attribute was already applied with different arguments, then
+  // diagnose the second attribute as a duplicate and don't add it.
+  if (const auto *Existing = D->getAttr<WorkGroupSizeHintAttr>()) {
+    DupArgResult Results[] = {AreArgValuesIdentical(XDim, Existing->getXDim()),
+                              AreArgValuesIdentical(YDim, Existing->getYDim()),
+                              AreArgValuesIdentical(ZDim, Existing->getZDim())};
+    // If any of the results are known to be different, we can diagnose at this
+    // point and drop the attribute.
+    if (llvm::is_contained(Results, DupArgResult::Different)) {
+      Diag(CI.getLoc(), diag::warn_duplicate_attribute) << CI;
+      Diag(Existing->getLoc(), diag::note_previous_attribute);
       return;
     }
+    // If all of the results are known to be the same, we can silently drop the
+    // attribute. Otherwise, we have to add the attribute and resolve its
+    // differences later.
+    if (llvm::all_of(Results,
+                     [](DupArgResult V) { return V == DupArgResult::Same; }))
+      return;
   }
 
-  WorkGroupSizeHintAttr *Existing = D->getAttr<WorkGroupSizeHintAttr>();
-  if (Existing &&
-      !(Existing->getXDim() == WGSize[0] && Existing->getYDim() == WGSize[1] &&
-        Existing->getZDim() == WGSize[2]))
-    S.Diag(AL.getLoc(), diag::warn_duplicate_attribute) << AL;
+  D->addAttr(::new (Context)
+                 WorkGroupSizeHintAttr(Context, CI, XDim, YDim, ZDim));
+}
+
+WorkGroupSizeHintAttr *
+Sema::MergeWorkGroupSizeHintAttr(Decl *D, const WorkGroupSizeHintAttr &A) {
+  // Check to see if there's a duplicate attribute already applied.
+  if (const auto *DeclAttr = D->getAttr<WorkGroupSizeHintAttr>()) {
+    DupArgResult Results[] = {
+        AreArgValuesIdentical(DeclAttr->getXDim(), A.getXDim()),
+        AreArgValuesIdentical(DeclAttr->getYDim(), A.getYDim()),
+        AreArgValuesIdentical(DeclAttr->getZDim(), A.getZDim())};
+
+    // If any of the results are known to be different, we can diagnose at this
+    // point and drop the attribute.
+    if (llvm::is_contained(Results, DupArgResult::Different)) {
+      Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A;
+      Diag(A.getLoc(), diag::note_previous_attribute);
+      return nullptr;
+    }
+    // If all of the results are known to be the same, we can silently drop the
+    // attribute. Otherwise, we have to add the attribute and resolve its
+    // differences later.
+    if (llvm::all_of(Results,
+                     [](DupArgResult V) { return V == DupArgResult::Same; }))
+      return nullptr;
+  }
+  return ::new (Context)
+      WorkGroupSizeHintAttr(Context, A, A.getXDim(), A.getYDim(), A.getZDim());
+}
+
+// Handles work_group_size_hint.
+static void handleWorkGroupSizeHint(Sema &S, Decl *D, const ParsedAttr &AL) {
+  S.CheckDeprecatedSYCLAttributeSpelling(AL);
 
-  D->addAttr(::new (S.Context) WorkGroupSizeHintAttr(S.Context, AL, WGSize[0],
-                                                     WGSize[1], WGSize[2]));
+  S.AddWorkGroupSizeHintAttr(D, AL, AL.getArgAsExpr(0), AL.getArgAsExpr(1),
+                             AL.getArgAsExpr(2));
 }
 
 void Sema::AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI,

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

@@ -743,6 +743,25 @@ static void instantiateSYCLIntelESimdVectorizeAttr(
     S.AddSYCLIntelESimdVectorizeAttr(New, *A, Result.getAs<Expr>());
 }
 
+static void instantiateWorkGroupSizeHintAttr(
+    Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
+    const WorkGroupSizeHintAttr *A, Decl *New) {
+  EnterExpressionEvaluationContext Unevaluated(
+      S, Sema::ExpressionEvaluationContext::ConstantEvaluated);
+  ExprResult XResult = S.SubstExpr(A->getXDim(), TemplateArgs);
+  if (XResult.isInvalid())
+    return;
+  ExprResult YResult = S.SubstExpr(A->getYDim(), TemplateArgs);
+  if (YResult.isInvalid())
+    return;
+  ExprResult ZResult = S.SubstExpr(A->getZDim(), TemplateArgs);
+  if (ZResult.isInvalid())
+    return;
+
+  S.AddWorkGroupSizeHintAttr(New, *A, XResult.get(), YResult.get(),
+                             ZResult.get());
+}
+
 /// Determine whether the attribute A might be relevent to the declaration D.
 /// If not, we can skip instantiating it. The attribute may or may not have
 /// been instantiated yet.
@@ -986,6 +1005,10 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
                                              SYCLIntelESimdVectorize, New);
       continue;
     }
+    if (const auto *A = dyn_cast<WorkGroupSizeHintAttr>(TmplAttr)) {
+      instantiateWorkGroupSizeHintAttr(*this, TemplateArgs, A, New);
+      continue;
+    }
     // Existing DLL attribute on the instantiation takes precedence.
     if (TmplAttr->getKind() == attr::DLLExport ||
         TmplAttr->getKind() == attr::DLLImport) {

clang/test/SemaOpenCL/invalid-kernel-attrs.cl

@@ -12,7 +12,7 @@ kernel __attribute__((vec_type_hint(int))) __attribute__((vec_type_hint(float)))
 
 kernel __attribute__((work_group_size_hint(8,16,32,4))) void kernel6() {} //expected-error{{'work_group_size_hint' attribute requires exactly 3 arguments}}
 
-kernel __attribute__((work_group_size_hint(1,2,3))) __attribute__((work_group_size_hint(3,2,1))) void kernel7() {}  //expected-warning{{attribute 'work_group_size_hint' is already applied with different arguments}}
+kernel __attribute__((work_group_size_hint(1,2,3))) __attribute__((work_group_size_hint(3,2,1))) void kernel7() {}  //expected-warning{{attribute 'work_group_size_hint' is already applied with different arguments}} expected-note {{previous attribute is here}}
 
 __attribute__((reqd_work_group_size(8,16,32))) void kernel8(){} // expected-error {{attribute 'reqd_work_group_size' can only be applied to an OpenCL kernel}}
 
@@ -42,7 +42,7 @@ kernel __attribute__((intel_reqd_sub_group_size(-1))) void kernel16() {} // expe
 kernel __attribute__((intel_reqd_sub_group_size(8))) __attribute__((intel_reqd_sub_group_size(16))) void kernel17() {} //expected-warning{{attribute 'intel_reqd_sub_group_size' is already applied with different arguments}} \
                                                                                                                        // expected-note {{previous attribute is here}}
 
-__kernel __attribute__((work_group_size_hint(8,-16,32))) void neg1() {} //expected-error{{'work_group_size_hint' attribute requires a non-negative integral compile time constant expression}}
+__kernel __attribute__((work_group_size_hint(8,-16,32))) void neg1() {} //expected-error{{'work_group_size_hint' attribute requires a positive integral compile time constant expression}}
 __kernel __attribute__((reqd_work_group_size(8, 16, -32))) void neg2() {} //expected-warning{{implicit conversion changes signedness: 'int' to 'unsigned long long'}}
 
 // 4294967294 is a negative integer if treated as signed.

clang/test/SemaSYCL/work-group-size-hint.cpp

@@ -3,13 +3,12 @@
 // Check the basics.
 [[sycl::work_group_size_hint]] void f0(); // expected-error {{'work_group_size_hint' attribute requires exactly 3 arguments}}
 [[sycl::work_group_size_hint(12, 12, 12, 12)]] void f1(); // expected-error {{'work_group_size_hint' attribute requires exactly 3 arguments}}
-[[sycl::work_group_size_hint("derp", 1, 2)]] void f2(); // expected-error {{'work_group_size_hint' attribute requires parameter 0 to be an integer constant}}
+[[sycl::work_group_size_hint("derp", 1, 2)]] void f2(); // expected-error {{integral constant expression must have integral or unscoped enumeration type, not 'const char [5]'}}
 [[sycl::work_group_size_hint(1, 1, 1)]] int i; // expected-error {{'work_group_size_hint' attribute only applies to functions}}
 
-// FIXME: this should produce a conflicting attribute warning but doesn't. It
-// is missing a merge method (and is also missing template instantiation logic).
-[[sycl::work_group_size_hint(4, 1, 1)]] void f3();
-[[sycl::work_group_size_hint(32, 1, 1)]] void f3() {}
+// Produce a conflicting attribute warning when the args are different.
+[[sycl::work_group_size_hint(4, 1, 1)]] void f3(); // expected-note {{previous attribute is here}}
+[[sycl::work_group_size_hint(32, 1, 1)]] void f3() {} // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}}
 
 // FIXME: the attribute is like reqd_work_group_size in that it has a one, two,
 // and three arg form that needs to be supported.
@@ -24,11 +23,43 @@ __attribute__((work_group_size_hint(4, 1, 1))) void f6(); // expected-warning {{
 
 // Catch the easy case where the attributes are all specified at once with
 // different arguments.
-[[sycl::work_group_size_hint(4, 1, 1), sycl::work_group_size_hint(32, 1, 1)]] void f7(); // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}}
+[[sycl::work_group_size_hint(4, 1, 1), sycl::work_group_size_hint(32, 1, 1)]] void f7(); // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}} expected-note {{previous attribute is here}}
 
 // Show that the attribute works on member functions.
 class Functor {
 public:
   [[sycl::work_group_size_hint(16, 1, 1)]] [[sycl::work_group_size_hint(16, 1, 1)]] void operator()() const;
-  [[sycl::work_group_size_hint(16, 1, 1)]] [[sycl::work_group_size_hint(32, 1, 1)]] void operator()(int) const; // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}}
+  [[sycl::work_group_size_hint(16, 1, 1)]] [[sycl::work_group_size_hint(32, 1, 1)]] void operator()(int) const; // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}} expected-note {{previous attribute is here}}
 };
+
+// Ensure that template arguments behave appropriately based on instantiations.
+template <int N>
+[[sycl::work_group_size_hint(N, 1, 1)]] void f8(); // #f8
+
+// Test that template redeclarations also get diagnosed properly.
+template <int X, int Y, int Z>
+[[sycl::work_group_size_hint(1, 1, 1)]] void f9(); // #f9prev
+
+template <int X, int Y, int Z>
+[[sycl::work_group_size_hint(X, Y, Z)]] void f9() {} // #f9
+
+// Test that a template redeclaration where the difference is known up front is
+// diagnosed immediately, even without instantiation.
+template <int X, int Y, int Z>
+[[sycl::work_group_size_hint(X, 1, Z)]] void f10(); // expected-note {{previous attribute is here}}
+template <int X, int Y, int Z>
+[[sycl::work_group_size_hint(X, 2, Z)]] void f10(); // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}}
+
+void instantiate() {
+  f8<1>(); // OK
+  // expected-error@#f8 {{'work_group_size_hint' attribute requires a positive integral compile time constant expression}}
+  f8<-1>(); // expected-note {{in instantiation}}
+  // expected-error@#f8 {{'work_group_size_hint' attribute requires a positive integral compile time constant expression}}
+  f8<0>(); // expected-note {{in instantiation}}
+
+  f9<1, 1, 1>(); // OK, args are the same on the redecl.
+
+  // expected-warning@#f9 {{attribute 'work_group_size_hint' is already applied with different arguments}}
+  // expected-note@#f9prev {{previous attribute is here}}
+  f9<1, 2, 3>(); // expected-note {{in instantiation}}
+}