[git] Add top-level .gitignore

[hliao]

Summary

[D56411] Temp solution fixing CUDA template issue

• template with overloadable kernel function as the template function need revising CheckCUDACall checking.

[SelectionDAG] Harden the checking of RegClass when adding operand

• If the operand index is out-of-range, expect nullptr is returned.

[AMDGPU] Allow using integral non-type template parameters

• Allow using integral non-type template parameters in the following attributes
  
  attribute((amdgpu_flat_work_group_size(<min>, <max>))) attribute((amdgpu_waves_per_eu(<min>[, <max>])))

[hliao]

by mistake, Ctrl-C is not fast enough to stop it.

.gitignore

+build.*/

clang/include/clang/Basic/Attr.td

 // use this in CUDA, since CUDA does not use the same terminology.
 //
 // FIXME: SubjectList should be for OpenCLKernelFunction, but is not to
 // workaround needing to see kernel attribute before others to know if
 // this should be rejected on non-kernels.
 
 def AMDGPUFlatWorkGroupSize : InheritableAttr {
   let Spellings = [Clang<"amdgpu_flat_work_group_size", 0>];
-  let Args = [UnsignedArgument<"Min">, UnsignedArgument<"Max">];
+  let Args = [ExprArgument<"Min">, ExprArgument<"Max">];
   let Documentation = [AMDGPUFlatWorkGroupSizeDocs];
   let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">;
 }
 
 def AMDGPUWavesPerEU : InheritableAttr {
   let Spellings = [Clang<"amdgpu_waves_per_eu", 0>];
-  let Args = [UnsignedArgument<"Min">, UnsignedArgument<"Max", 1>];
+  let Args = [ExprArgument<"Min">, ExprArgument<"Max", 1>];
   let Documentation = [AMDGPUWavesPerEUDocs];
   let Subjects = SubjectList<[Function], ErrorDiag, "kernel functions">;
 }
 
 def AMDGPUNumSGPR : InheritableAttr {
   let Spellings = [Clang<"amdgpu_num_sgpr", 0>];
   let Args = [UnsignedArgument<"NumSGPR">];
   let Documentation = [AMDGPUNumSGPRNumVGPRDocs];

clang/include/clang/Sema/Sema.h

struct ExpressionEvaluationContextRecord {
     llvm::SmallPtrSet<const Expr *, 8> PossibleDerefs;
 
     /// \brief Describes whether we are in an expression constext which we have
     /// to handle differently.
     enum ExpressionKind {
       EK_Decltype, EK_TemplateArgument, EK_Other
     } ExprContext;
 
+    /// If we are checking arguments of a template, this is the template
+    /// under check.
+    TemplateDecl *Template;
+
     ExpressionEvaluationContextRecord(ExpressionEvaluationContext Context,
                                       unsigned NumCleanupObjects,
                                       CleanupInfo ParentCleanup,
                                       Decl *ManglingContextDecl,
                                       ExpressionKind ExprContext)
         : Context(Context), ParentCleanup(ParentCleanup),
           NumCleanupObjects(NumCleanupObjects), NumTypos(0),
           ManglingContextDecl(ManglingContextDecl), MangleNumbering(),
-          ExprContext(ExprContext) {}
+          ExprContext(ExprContext), Template(nullptr) {}
 
     /// Retrieve the mangling numbering context, used to consistently
     /// number constructs like lambdas for mangling.
     MangleNumberingContext &getMangleNumberingContext(ASTContext &Ctx);
 
     bool isUnevaluated() const {
       return Context == ExpressionEvaluationContext::Unevaluated ||
              Context == ExpressionEvaluationContext::UnevaluatedAbstract ||
bool CheckTemplateArgumentList(TemplateDecl *Template,
                                  bool UpdateArgsWithConversions = true);
 
   bool CheckTemplateTypeArgument(TemplateTypeParmDecl *Param,
                                  TemplateArgumentLoc &Arg,
                            SmallVectorImpl<TemplateArgument> &Converted);
 
   bool CheckTemplateArgument(TemplateTypeParmDecl *Param,
                              TypeSourceInfo *Arg);
-  ExprResult CheckTemplateArgument(NonTypeTemplateParmDecl *Param,
-                                   QualType InstantiatedParamType, Expr *Arg,
-                                   TemplateArgument &Converted,
-                               CheckTemplateArgumentKind CTAK = CTAK_Specified);
+  ExprResult
+  CheckTemplateArgument(NonTypeTemplateParmDecl *Param,
+                        CheckTemplateArgumentKind CTAK = CTAK_Specified,
+                        TemplateDecl *Template = nullptr);
   bool CheckTemplateTemplateArgument(TemplateParameterList *Params,
                                      TemplateArgumentLoc &Arg);
 
   ExprResult
   BuildExpressionFromDeclTemplateArgument(const TemplateArgument &Arg,
                                           QualType ParamType,
                                           SourceLocation Loc);
   ExprResult

 
   void AddParameterABIAttr(SourceRange AttrRange, Decl *D,
                            ParameterABI ABI, unsigned SpellingListIndex);
 
   enum class RetainOwnershipKind {NS, CF, OS};
   void AddXConsumedAttr(Decl *D, SourceRange SR, unsigned SpellingIndex,
                         RetainOwnershipKind K, bool IsTemplateInstantiation);
 
+  /// addAMDGPUFlatWorkGroupSizeAttr - Adds an amdgpu_flat_work_group_size
+  /// attribute to a particular declaration.
+  void addAMDGPUFlatWorkGroupSizeAttr(SourceRange AttrRange, Decl *D, Expr *Min,
+                                      Expr *Max, unsigned SpellingListIndex);
+
+  /// addAMDGPUWavePersEUAttr - Adds an amdgpu_waves_per_eu attribute to a
+  /// particular declaration.
+  void addAMDGPUWavesPerEUAttr(SourceRange AttrRange, Decl *D, Expr *Min,
+                               Expr *Max, unsigned SpellingListIndex);
+
   bool checkNSReturnsRetainedReturnType(SourceLocation loc, QualType type);
 
   //===--------------------------------------------------------------------===//
   // C++ Coroutines TS
   //
   bool ActOnCoroutineBodyStart(Scope *S, SourceLocation KwLoc,
                                StringRef Keyword);
   ExprResult ActOnCoawaitExpr(Scope *S, SourceLocation KwLoc, Expr *E);

clang/lib/CodeGen/TargetInfo.cpp

const auto *ReqdWGS = M.getLangOpts().OpenCL ?
     FD->getAttr<ReqdWorkGroupSizeAttr>() : nullptr;
 
   if (M.getLangOpts().OpenCL && FD->hasAttr<OpenCLKernelAttr>() &&
       (M.getTriple().getOS() == llvm::Triple::AMDHSA))
     F->addFnAttr("amdgpu-implicitarg-num-bytes", "48");
 
   const auto *FlatWGS = FD->getAttr<AMDGPUFlatWorkGroupSizeAttr>();
   if (ReqdWGS || FlatWGS) {
-    unsigned Min = FlatWGS ? FlatWGS->getMin() : 0;
-    unsigned Max = FlatWGS ? FlatWGS->getMax() : 0;
+    unsigned Min = 0;
+    unsigned Max = 0;
+    if (FlatWGS) {
+      Min = FlatWGS->getMin()
+                ->EvaluateKnownConstInt(M.getContext())
+                .getExtValue();
+      Max = FlatWGS->getMax()
+                ->EvaluateKnownConstInt(M.getContext())
+                .getExtValue();
+    }
     if (ReqdWGS && Min == 0 && Max == 0)
       Min = Max = ReqdWGS->getXDim() * ReqdWGS->getYDim() * ReqdWGS->getZDim();
 
     if (Min != 0) {
       assert(Min <= Max && "Min must be less than or equal Max");
 
       std::string AttrVal = llvm::utostr(Min) + "," + llvm::utostr(Max);
       F->addFnAttr("amdgpu-flat-work-group-size", AttrVal);
     } else
       assert(Max == 0 && "Max must be zero");
   }
 
   if (const auto *Attr = FD->getAttr<AMDGPUWavesPerEUAttr>()) {
-    unsigned Min = Attr->getMin();
-    unsigned Max = Attr->getMax();
+    unsigned Min =
+        Attr->getMin()->EvaluateKnownConstInt(M.getContext()).getExtValue();
+    unsigned Max = Attr->getMax() ? Attr->getMax()
+                                        ->EvaluateKnownConstInt(M.getContext())
+                                        .getExtValue()
+                                  : 0;
 
     if (Min != 0) {
       assert((Max == 0 || Min <= Max) && "Min must be less than or equal Max");
 
       std::string AttrVal = llvm::utostr(Min);
       if (Max != 0)
         AttrVal = AttrVal + "," + llvm::utostr(Max);
       F->addFnAttr("amdgpu-waves-per-eu", AttrVal);

clang/lib/Sema/SemaCUDA.cpp

Sema::DeviceDiagBuilder Sema::CUDADiagIfHostCode(SourceLocation Loc,
   }();
   return DeviceDiagBuilder(DiagKind, Loc, DiagID,
                            dyn_cast<FunctionDecl>(CurContext), *this);
 }
 
 bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
   assert(Callee && "Callee may not be null.");
+
+  auto &ExprEvalCtx = ExprEvalContexts.back();
+  if (ExprEvalCtx.isUnevaluated())
+    return true;
+
+  FunctionDecl *Caller = nullptr;
+  if (auto *Template = ExprEvalContexts.back().Template) {
+    if (auto *FD = dyn_cast<FunctionDecl>(Template->getTemplatedDecl()))
+      Caller = FD;
+  } else if (ExprEvalCtx.isConstantEvaluated())
+    return true;
+
   // FIXME: Is bailing out early correct here?  Should we instead assume that
   // the caller is a global initializer?
-  FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext);
+  if (!Caller)
+    Caller = dyn_cast<FunctionDecl>(CurContext);
   if (!Caller)
     return true;
 
   // If the caller is known-emitted, mark the callee as known-emitted.
   // Otherwise, mark the call in our call graph so we can traverse it later.
   bool CallerKnownEmitted = IsKnownEmitted(*this, Caller);
   if (CallerKnownEmitted) {
     // Host-side references to a __global__ function refer to the stub, so the

clang/lib/Sema/SemaDeclAttr.cpp

 static bool checkUInt32Argument(Sema &S, const AttrInfo &AI, const Expr *Expr,
                                 uint32_t &Val, unsigned Idx = UINT_MAX,
                                 bool StrictlyUnsigned = false) {
   llvm::APSInt I(32);
   if (Expr->isTypeDependent() || Expr->isValueDependent() ||
       !Expr->isIntegerConstantExpr(I, S.Context)) {
     if (Idx != UINT_MAX)
       S.Diag(getAttrLoc(AI), diag::err_attribute_argument_n_type)
-          << AI << Idx << AANT_ArgumentIntegerConstant
+          << &AI << Idx << AANT_ArgumentIntegerConstant
           << Expr->getSourceRange();
     else
       S.Diag(getAttrLoc(AI), diag::err_attribute_argument_type)
-          << AI << AANT_ArgumentIntegerConstant << Expr->getSourceRange();
+          << &AI << AANT_ArgumentIntegerConstant << Expr->getSourceRange();
     return false;
   }
 
   if (!I.isIntN(32)) {
     S.Diag(Expr->getExprLoc(), diag::err_ice_too_large)
         << I.toString(10, false) << 32 << /* Unsigned */ 1;
     return false;
   }
 
   if (StrictlyUnsigned && I.isSigned() && I.isNegative()) {
     S.Diag(getAttrLoc(AI), diag::err_attribute_requires_positive_integer)
-        << AI << /*non-negative*/ 1;
+        << &AI << /*non-negative*/ 1;
     return false;
   }
 
   Val = (uint32_t)I.getZExtValue();
   return true;
 }
 
 /// Wrapper around checkUInt32Argument, with an extra check to be sure
case llvm::Triple::riscv64:
     handleRISCVInterruptAttr(S, D, AL);
     break;
   default:
     handleARMInterruptAttr(S, D, AL);
     break;
   }
 }
 
-static void handleAMDGPUFlatWorkGroupSizeAttr(Sema &S, Decl *D,
-                                              const ParsedAttr &AL) {
+static bool
+checkAMDGPUFlatWorkGroupSizeArguments(Sema &S, Expr *MinExpr, Expr *MaxExpr,
+                                      const AMDGPUFlatWorkGroupSizeAttr &Attr) {
+  // Accept template arguments for now as they depend on something else.
+  // We'll get to check them when they eventually get instantiated.
+  if (MinExpr->isValueDependent() || MaxExpr->isValueDependent())
+    return false;
+
   uint32_t Min = 0;
-  Expr *MinExpr = AL.getArgAsExpr(0);
-  if (!checkUInt32Argument(S, AL, MinExpr, Min))
-    return;
+  if (!checkUInt32Argument(S, Attr, MinExpr, Min, 0))
+    return true;
 
   uint32_t Max = 0;
-  Expr *MaxExpr = AL.getArgAsExpr(1);
-  if (!checkUInt32Argument(S, AL, MaxExpr, Max))
-    return;
+  if (!checkUInt32Argument(S, Attr, MaxExpr, Max, 1))
+    return true;
 
   if (Min == 0 && Max != 0) {
-    S.Diag(AL.getLoc(), diag::err_attribute_argument_invalid) << AL << 0;
-    return;
+    S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
+        << &Attr << 0;
+    return true;
   }
   if (Min > Max) {
-    S.Diag(AL.getLoc(), diag::err_attribute_argument_invalid) << AL << 1;
-    return;
+    S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
+        << &Attr << 1;
+    return true;
   }
 
-  D->addAttr(::new (S.Context)
-             AMDGPUFlatWorkGroupSizeAttr(AL.getLoc(), S.Context, Min, Max,
-                                         AL.getAttributeSpellingListIndex()));
+  return false;
 }
 
-static void handleAMDGPUWavesPerEUAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
-  uint32_t Min = 0;
-  Expr *MinExpr = AL.getArgAsExpr(0);
-  if (!checkUInt32Argument(S, AL, MinExpr, Min))
+void Sema::addAMDGPUFlatWorkGroupSizeAttr(SourceRange AttrRange, Decl *D,
+                                          Expr *MinExpr, Expr *MaxExpr,
+                                          unsigned SpellingListIndex) {
+  AMDGPUFlatWorkGroupSizeAttr TmpAttr(AttrRange, Context, MinExpr, MaxExpr,
+                                      SpellingListIndex);
+
+  if (checkAMDGPUFlatWorkGroupSizeArguments(*this, MinExpr, MaxExpr, TmpAttr))
     return;
 
-  uint32_t Max = 0;
-  if (AL.getNumArgs() == 2) {
-    Expr *MaxExpr = AL.getArgAsExpr(1);
-    if (!checkUInt32Argument(S, AL, MaxExpr, Max))
-      return;
-  }
+  D->addAttr(::new (Context) AMDGPUFlatWorkGroupSizeAttr(
+      AttrRange, Context, MinExpr, MaxExpr, SpellingListIndex));
+}
+
+static void handleAMDGPUFlatWorkGroupSizeAttr(Sema &S, Decl *D,
+                                              const ParsedAttr &AL) {
+  Expr *MinExpr = AL.getArgAsExpr(0);
+
+  S.addAMDGPUFlatWorkGroupSizeAttr(AL.getRange(), D, MinExpr, MaxExpr,
+                                   AL.getAttributeSpellingListIndex());
 
+static bool checkAMDGPUWavesPerEUArguments(Sema &S, Expr *MinExpr,
+                                           Expr *MaxExpr,
+                                           const AMDGPUWavesPerEUAttr &Attr) {
+  if (S.DiagnoseUnexpandedParameterPack(MinExpr) ||
+      (MaxExpr && S.DiagnoseUnexpandedParameterPack(MaxExpr)))
+    return true;
+
+  // Accept template arguments for now as they depend on something else.
+  // We'll get to check them when they eventually get instantiated.
+  if (MinExpr->isValueDependent() || (MaxExpr && MaxExpr->isValueDependent()))
+    return false;
+
+  uint32_t Min = 0;
+  if (!checkUInt32Argument(S, Attr, MinExpr, Min, 0))
+    return true;
+
+  uint32_t Max = 0;
+  if (MaxExpr && !checkUInt32Argument(S, Attr, MaxExpr, Max, 1))
+    return true;
+
   if (Min == 0 && Max != 0) {
-    S.Diag(AL.getLoc(), diag::err_attribute_argument_invalid) << AL << 0;
-    return;
+    S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
+        << &Attr << 0;
+    return true;
   }
   if (Max != 0 && Min > Max) {
-    S.Diag(AL.getLoc(), diag::err_attribute_argument_invalid) << AL << 1;
+    S.Diag(Attr.getLocation(), diag::err_attribute_argument_invalid)
+        << &Attr << 1;
+    return true;
+  }
+
+  return false;
+}
+
+void Sema::addAMDGPUWavesPerEUAttr(SourceRange AttrRange, Decl *D,
+                                   Expr *MinExpr, Expr *MaxExpr,
+                                   unsigned SpellingListIndex) {
+  AMDGPUWavesPerEUAttr TmpAttr(AttrRange, Context, MinExpr, MaxExpr,
+                               SpellingListIndex);
+
+  if (checkAMDGPUWavesPerEUArguments(*this, MinExpr, MaxExpr, TmpAttr))
     return;
-  }
+
+  D->addAttr(::new (Context) AMDGPUWavesPerEUAttr(AttrRange, Context, MinExpr,
+                                                  MaxExpr, SpellingListIndex));
 
-  D->addAttr(::new (S.Context)
-             AMDGPUWavesPerEUAttr(AL.getLoc(), S.Context, Min, Max,
-                                  AL.getAttributeSpellingListIndex()));
+static void handleAMDGPUWavesPerEUAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
+  if (!checkAttributeAtLeastNumArgs(S, AL, 1) ||
+      !checkAttributeAtMostNumArgs(S, AL, 2))
+    return;
+
+  Expr *MinExpr = AL.getArgAsExpr(0);
+  Expr *MaxExpr = (AL.getNumArgs() > 1) ? AL.getArgAsExpr(1) : nullptr;
+
+  S.addAMDGPUWavesPerEUAttr(AL.getRange(), D, MinExpr, MaxExpr,
+                            AL.getAttributeSpellingListIndex());
 }
 
 static void handleAMDGPUNumSGPRAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
   uint32_t NumSGPR = 0;
   Expr *NumSGPRExpr = AL.getArgAsExpr(0);
   if (!checkUInt32Argument(S, AL, NumSGPRExpr, NumSGPR))
     return;
 

clang/lib/Sema/SemaTemplate.cpp

SubstDefaultTemplateArgument(Sema &SemaRef,
   // Only substitute for the innermost template argument list.
   MultiLevelTemplateArgumentList TemplateArgLists;
   TemplateArgLists.addOuterTemplateArguments(&TemplateArgs);
   for (unsigned i = 0, e = Param->getDepth(); i != e; ++i)
     TemplateArgLists.addOuterTemplateArguments(None);
 
   EnterExpressionEvaluationContext ConstantEvaluated(
       SemaRef, Sema::ExpressionEvaluationContext::ConstantEvaluated);
+  SemaRef.ExprEvalContexts.back().Template = Template;
   return SemaRef.SubstExpr(Param->getDefaultArgument(), TemplateArgLists);
 }
 
 /// Substitute template arguments into the default template argument for
 /// the given template template parameter.
 ///
 /// \param SemaRef the semantic analysis object for which we are performing
 /// the substitution.
if (NonTypeTemplateParmDecl *NTTP =dyn_cast<NonTypeTemplateParmDecl>(Param)) {
     switch (Arg.getArgument().getKind()) {
     case TemplateArgument::Null:
       llvm_unreachable("Should never see a NULL template argument here");
 
     case TemplateArgument::Expression: {
       TemplateArgument Result;
       unsigned CurSFINAEErrors = NumSFINAEErrors;
       ExprResult Res =
           CheckTemplateArgument(NTTP, NTTPType, Arg.getArgument().getAsExpr(),
-                              Result, CTAK);
+                                Result, CTAK, dyn_cast<TemplateDecl>(Template));
       if (Res.isInvalid())
         return true;
       // If the current template argument causes an error, give up now.
       if (CurSFINAEErrors < NumSFINAEErrors)
         return true;
 
       // If the resulting expression is new, then use it in place of the
       // old expression in the template argument.
static bool CheckTemplateArgumentPointerToMember(Sema &S,
 
   // We found something else, but we don't know specifically what it is.
   S.Diag(Arg->getBeginLoc(), diag::err_template_arg_not_pointer_to_member_form)
       << Arg->getSourceRange();
   S.Diag(DRE->getDecl()->getLocation(), diag::note_template_arg_refers_here);
   return true;
 }
 
+namespace {
+FunctionDecl *GetFunctionDecl(Expr *Arg) {
+  Expr *E = Arg;
+  if (UnaryOperator *UO = dyn_cast<UnaryOperator>(E)) {
+    E = UO ? UO->getSubExpr() : nullptr;
+  }
+  if (DeclRefExpr *DRE = dyn_cast_or_null<DeclRefExpr>(E)) {
+    ValueDecl *Entity = DRE ? DRE->getDecl() : nullptr;
+    if (Entity) {
+      if (auto Callee = dyn_cast<FunctionDecl>(Entity))
+        return Callee;
+    }
+  }
+  return nullptr;
+}
+} // namespace
 /// Check a template argument against its corresponding
 /// non-type template parameter.
 ///
 /// This routine implements the semantics of C++ [temp.arg.nontype].
 /// If an error occurred, it returns ExprError(); otherwise, it
 /// returns the converted template argument. \p ParamType is the
 /// type of the non-type template parameter after it has been instantiated.
 ExprResult Sema::CheckTemplateArgument(NonTypeTemplateParmDecl *Param,
                                        QualType ParamType, Expr *Arg,
                                        TemplateArgument &Converted,
-                                       CheckTemplateArgumentKind CTAK) {
+                                       CheckTemplateArgumentKind CTAK,
+                                       TemplateDecl *Template) {
   SourceLocation StartLoc = Arg->getBeginLoc();
 
   // If the parameter type somehow involves auto, deduce the type now.
   if (getLangOpts().CPlusPlus17 && ParamType->isUndeducedType()) {
     // During template argument deduction, we allow 'decltype(auto)' to
     // match an arbitrary dependent argument.
     // FIXME: The language rules don't say what happens in this case.
     // FIXME: We get an opaque dependent type out of decltype(auto) if the
if (ParamType->isDependentType() || Arg->isTypeDependent()) {
     Converted = TemplateArgument(Arg);
     return Arg;
   }
 
   // The initialization of the parameter from the argument is
   // a constant-evaluated context.
   EnterExpressionEvaluationContext ConstantEvaluated(
       *this, Sema::ExpressionEvaluationContext::ConstantEvaluated);
+  ExprEvalContexts.back().Template = Template;
 
   if (getLangOpts().CPlusPlus17) {
     // C++17 [temp.arg.nontype]p1:
     //   A template-argument for a non-type template parameter shall be
     //   a converted constant expression of the type of the template-parameter.
     APValue Value;
     ExprResult ArgResult = CheckConvertedConstantExpression(
         Arg, ParamType, Value, CCEK_TemplateArg);
if (Arg->getType() == Context.OverloadTy) {
           return ExprError();
 
         Arg = FixOverloadedFunctionReference(Arg, FoundResult, Fn);
         ArgType = Arg->getType();
       } else
         return ExprError();
     }
 
+    if (auto *FD = GetFunctionDecl(Arg))
+      if (getLangOpts().CUDA && !CheckCUDACall(Arg->getBeginLoc(), FD))
+        return ExprError();
+
     if (!ParamType->isMemberPointerType()) {
       if (CheckTemplateArgumentAddressOfObjectOrFunction(*this, Param,
                                                          ParamType,
                                                          Arg, Converted))
         return ExprError();
       return Arg;
     }
 

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

static void instantiateOMPDeclareSimdDeclAttr(
   }
   LinModifiers.append(Attr.modifiers_begin(), Attr.modifiers_end());
   (void)S.ActOnOpenMPDeclareSimdDirective(
       S.ConvertDeclToDeclGroup(New), Attr.getBranchState(), Simdlen.get(),
       Uniforms, Aligneds, Alignments, Linears, LinModifiers, Steps,
       Attr.getRange());
 }
 
+static void instantiateDependentAMDGPUFlatWorkGroupSizeAttr(
+    Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
+    const AMDGPUFlatWorkGroupSizeAttr &Attr, Decl *New) {
+  // Both min and max expression are constant expressions.
+  EnterExpressionEvaluationContext Unevaluated(
+      S, Sema::ExpressionEvaluationContext::ConstantEvaluated);
+
+  ExprResult Result = S.SubstExpr(Attr.getMin(), TemplateArgs);
+  if (Result.isInvalid())
+    return;
+  Expr *MinExpr = Result.getAs<Expr>();
+
+  Result = S.SubstExpr(Attr.getMax(), TemplateArgs);
+  if (Result.isInvalid())
+    return;
+  Expr *MaxExpr = Result.getAs<Expr>();
+
+  S.addAMDGPUFlatWorkGroupSizeAttr(Attr.getLocation(), New, MinExpr, MaxExpr,
+                                   Attr.getSpellingListIndex());
+}
+
+static void instantiateDependentAMDGPUWavesPerEUAttr(
+    Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
+    const AMDGPUWavesPerEUAttr &Attr, Decl *New) {
+  // Both min and max expression are constant expressions.
+  EnterExpressionEvaluationContext Unevaluated(
+      S, Sema::ExpressionEvaluationContext::ConstantEvaluated);
+
+  ExprResult Result = S.SubstExpr(Attr.getMin(), TemplateArgs);
+  if (Result.isInvalid())
+    return;
+  Expr *MinExpr = Result.getAs<Expr>();
+
+  Expr *MaxExpr = nullptr;
+  if (auto Max = Attr.getMax()) {
+    Result = S.SubstExpr(Max, TemplateArgs);
+    if (Result.isInvalid())
+      return;
+    MaxExpr = Result.getAs<Expr>();
+  }
+
+  S.addAMDGPUWavesPerEUAttr(Attr.getLocation(), New, MinExpr, MaxExpr,
+                            Attr.getSpellingListIndex());
+}
+
 void Sema::InstantiateAttrsForDecl(
     const MultiLevelTemplateArgumentList &TemplateArgs, const Decl *Tmpl,
     Decl *New, LateInstantiatedAttrVec *LateAttrs,
     LocalInstantiationScope *OuterMostScope) {
   if (NamedDecl *ND = dyn_cast<NamedDecl>(New)) {
     for (const auto *TmplAttr : Tmpl->attrs()) {
       // FIXME: If any of the special case versions from InstantiateAttrs become
       // applicable to template declaration, we'll need to add them here.
if (const ModeAttr *Mode = dyn_cast<ModeAttr>(TmplAttr)) {
       continue;
     }
 
     if (const auto *OMPAttr = dyn_cast<OMPDeclareSimdDeclAttr>(TmplAttr)) {
       instantiateOMPDeclareSimdDeclAttr(*this, TemplateArgs, *OMPAttr, New);
       continue;
     }
 
+    if (const AMDGPUFlatWorkGroupSizeAttr *AMDGPUFlatWorkGroupSize =
+            dyn_cast<AMDGPUFlatWorkGroupSizeAttr>(TmplAttr)) {
+      instantiateDependentAMDGPUFlatWorkGroupSizeAttr(
+          *this, TemplateArgs, *AMDGPUFlatWorkGroupSize, New);
+    }
+
+    if (const AMDGPUWavesPerEUAttr *AMDGPUFlatWorkGroupSize =
+            dyn_cast<AMDGPUWavesPerEUAttr>(TmplAttr)) {
+      instantiateDependentAMDGPUWavesPerEUAttr(*this, TemplateArgs,
+                                               *AMDGPUFlatWorkGroupSize, New);
+    }
+
     // Existing DLL attribute on the instantiation takes precedence.
     if (TmplAttr->getKind() == attr::DLLExport ||
         TmplAttr->getKind() == attr::DLLImport) {
       if (New->hasAttr<DLLExportAttr>() || New->hasAttr<DLLImportAttr>()) {
         continue;
       }
     }
 

clang/test/SemaCUDA/amdgpu-attrs.cu

-// RUN: %clang_cc1 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s
 #include "Inputs/cuda.h"
 
 
 __attribute__((amdgpu_flat_work_group_size(32, 64)))
 __global__ void flat_work_group_size_32_64() {}
 
 __attribute__((amdgpu_waves_per_eu(2)))
 __global__ void waves_per_eu_2() {}

 
 // expected-error@+2{{attribute 'vec_type_hint' can only be applied to an OpenCL kernel function}}
 __attribute__((vec_type_hint(int)))
 __global__ void vec_type_hint_int() {}
 
 // expected-error@+2{{attribute 'intel_reqd_sub_group_size' can only be applied to an OpenCL kernel function}}
 __attribute__((intel_reqd_sub_group_size(64)))
 __global__ void intel_reqd_sub_group_size_64() {}
+
+// expected-error@+1{{'amdgpu_flat_work_group_size' attribute requires parameter 0 to be an integer constant}}
+__attribute__((amdgpu_flat_work_group_size("32", 64)))
+__global__ void non_int_min_flat_work_group_size_32_64() {}
+// expected-error@+1{{'amdgpu_flat_work_group_size' attribute requires parameter 1 to be an integer constant}}
+__attribute__((amdgpu_flat_work_group_size(32, "64")))
+__global__ void non_int_max_flat_work_group_size_32_64() {}
+
+int nc_min = 32, nc_max = 64;
+// expected-error@+1{{'amdgpu_flat_work_group_size' attribute requires parameter 0 to be an integer constant}}
+__attribute__((amdgpu_flat_work_group_size(nc_min, 64)))
+__global__ void non_cint_min_flat_work_group_size_32_64() {}
+// expected-error@+1{{'amdgpu_flat_work_group_size' attribute requires parameter 1 to be an integer constant}}
+__attribute__((amdgpu_flat_work_group_size(32, nc_max)))
+__global__ void non_cint_max_flat_work_group_size_32_64() {}
+
+const int c_min = 16, c_max = 32;
+__attribute__((amdgpu_flat_work_group_size(c_min * 2, 64)))
+__global__ void cint_min_flat_work_group_size_32_64() {}
+__attribute__((amdgpu_flat_work_group_size(32, c_max * 2)))
+__global__ void cint_max_flat_work_group_size_32_64() {}
+
+// expected-error@+3{{'T' does not refer to a value}}
+// expected-note@+1{{declared here}}
+template<typename T>
+__attribute__((amdgpu_flat_work_group_size(T, 64)))
+__global__ void template_class_min_flat_work_group_size_32_64() {}
+// expected-error@+3{{'T' does not refer to a value}}
+// expected-note@+1{{declared here}}
+template<typename T>
+__attribute__((amdgpu_flat_work_group_size(32, T)))
+__global__ void template_class_max_flat_work_group_size_32_64() {}
+
+template<unsigned a, unsigned b>
+__attribute__((amdgpu_flat_work_group_size(a, b)))
+__global__ void template_flat_work_group_size_32_64() {}
+template __global__ void template_flat_work_group_size_32_64<32, 64>();
+
+template<unsigned a, unsigned b, unsigned c>
+__attribute__((amdgpu_flat_work_group_size(a + b, b + c)))
+__global__ void template_complex_flat_work_group_size_32_64() {}
+template __global__ void template_complex_flat_work_group_size_32_64<16, 16, 48>();
+
+unsigned ipow2(unsigned n) { return n == 0 ? 1 : 2 * ipow2(n - 1); }
+constexpr unsigned ce_ipow2(unsigned n) { return n == 0 ? 1 : 2 * ce_ipow2(n - 1); }
+
+__attribute__((amdgpu_flat_work_group_size(ce_ipow2(5), ce_ipow2(6))))
+__global__ void cexpr_flat_work_group_size_32_64() {}
+// expected-error@+1{{'amdgpu_flat_work_group_size' attribute requires parameter 0 to be an integer constant}}
+__attribute__((amdgpu_flat_work_group_size(ipow2(5), 64)))
+__global__ void non_cexpr_min_flat_work_group_size_32_64() {}
+// expected-error@+1{{'amdgpu_flat_work_group_size' attribute requires parameter 1 to be an integer constant}}
+__attribute__((amdgpu_flat_work_group_size(32, ipow2(6))))
+__global__ void non_cexpr_max_flat_work_group_size_32_64() {}
+
+// expected-error@+1{{'amdgpu_waves_per_eu' attribute requires parameter 0 to be an integer constant}}
+__attribute__((amdgpu_waves_per_eu("2")))
+__global__ void non_int_min_waves_per_eu_2() {}
+// expected-error@+1{{'amdgpu_waves_per_eu' attribute requires parameter 1 to be an integer constant}}
+__attribute__((amdgpu_waves_per_eu(2, "4")))
+__global__ void non_int_max_waves_per_eu_2_4() {}
+
+// expected-error@+1{{'amdgpu_waves_per_eu' attribute requires parameter 0 to be an integer constant}}
+__attribute__((amdgpu_waves_per_eu(nc_min)))
+__global__ void non_cint_min_waves_per_eu_2() {}
+// expected-error@+1{{'amdgpu_waves_per_eu' attribute requires parameter 1 to be an integer constant}}
+__attribute__((amdgpu_waves_per_eu(2, nc_max)))
+__global__ void non_cint_min_waves_per_eu_2_4() {}
+
+__attribute__((amdgpu_waves_per_eu(c_min / 8)))
+__global__ void cint_min_waves_per_eu_2() {}
+__attribute__((amdgpu_waves_per_eu(c_min / 8, c_max / 8)))
+__global__ void cint_min_waves_per_eu_2_4() {}
+
+// expected-error@+3{{'T' does not refer to a value}}
+// expected-note@+1{{declared here}}
+template<typename T>
+__attribute__((amdgpu_waves_per_eu(T)))
+__global__ void cint_min_waves_per_eu_2() {}
+// expected-error@+3{{'T' does not refer to a value}}
+// expected-note@+1{{declared here}}
+template<typename T>
+__attribute__((amdgpu_waves_per_eu(2, T)))
+__global__ void cint_min_waves_per_eu_2_4() {}
+
+template<unsigned a>
+__attribute__((amdgpu_waves_per_eu(a)))
+__global__ void template_waves_per_eu_2() {}
+template __global__ void template_waves_per_eu_2<2>();
+
+template<unsigned a, unsigned b>
+__attribute__((amdgpu_waves_per_eu(a, b)))
+__global__ void template_waves_per_eu_2_4() {}
+template __global__ void template_waves_per_eu_2_4<2, 4>();
+
+template<unsigned a, unsigned b, unsigned c>
+__attribute__((amdgpu_waves_per_eu(a + b, c - b)))
+__global__ void template_complex_waves_per_eu_2_4() {}
+template __global__ void template_complex_waves_per_eu_2_4<1, 1, 5>();
+
+// expected-error@+2{{expression contains unexpanded parameter pack 'Args'}}
+template<unsigned... Args>
+__attribute__((amdgpu_waves_per_eu(Args)))
+__global__ void template_waves_per_eu_2() {}
+template __global__ void template_waves_per_eu_2<2, 4>();
+
+__attribute__((amdgpu_waves_per_eu(ce_ipow2(1))))
+__global__ void cexpr_waves_per_eu_2() {}
+__attribute__((amdgpu_waves_per_eu(ce_ipow2(1), ce_ipow2(2))))
+__global__ void cexpr_waves_per_eu_2_4() {}
+// expected-error@+1{{'amdgpu_waves_per_eu' attribute requires parameter 0 to be an integer constant}}
+__attribute__((amdgpu_waves_per_eu(ipow2(1))))
+__global__ void non_cexpr_waves_per_eu_2() {}
+// expected-error@+1{{'amdgpu_waves_per_eu' attribute requires parameter 1 to be an integer constant}}
+__attribute__((amdgpu_waves_per_eu(2, ipow2(2))))
+__global__ void non_cexpr_waves_per_eu_2_4() {}

clang/test/SemaCUDA/kernel-template-with-func-arg.cu

+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+struct C {
+  __device__ void devfun() {}
+  void hostfun() {}
+  template<class T> __device__ void devtempfun() {}
+  __device__ __host__ void devhostfun() {}
+};
+
+__device__ void devfun() {}
+__host__ void hostfun() {}
+template<class T> __device__ void devtempfun() {}
+__device__ __host__ void devhostfun() {}
+
+template <void (*devF)()> __global__ void kernel() { devF();}
+template <typename T, void(T::*devF)()> __global__ void kernel2(T *p) { (p->*devF)(); }
+
+template<> __global__ void kernel<devfun>();
+template<> __global__ void kernel<hostfun>(); // expected-error {{no function template matches function template specialization 'kernel'}}
+                                              // expected-note@-5 {{candidate template ignored: invalid explicitly-specified argument for template parameter 'devF'}}
+template<> __global__ void kernel<devtempfun<int> >();
+template<> __global__ void kernel<devhostfun>();
+
+template<> __global__ void kernel<&devfun>();
+template<> __global__ void kernel<&hostfun>(); // expected-error {{no function template matches function template specialization 'kernel'}}
+                                               // expected-note@-11 {{candidate template ignored: invalid explicitly-specified argument for template parameter 'devF'}}
+template<> __global__ void kernel<&devtempfun<int> >();
+template<> __global__ void kernel<&devhostfun>();
+
+template<> __global__ void kernel2<C, &C::devfun>(C *p);
+template<> __global__ void kernel2<C, &C::hostfun>(C *p); // expected-error {{no function template matches function template specialization 'kernel2'}}
+                                                          // expected-note@-16 {{candidate template ignored: invalid explicitly-specified argument for template parameter 'devF'}}
+template<> __global__ void kernel2<C, &C::devtempfun<int> >(C *p);
+template<> __global__ void kernel2<C, &C::devhostfun>(C *p);
+
+void fun() {
+  kernel<&devfun><<<1,1>>>();
+  kernel<&hostfun><<<1,1>>>(); // expected-error {{no matching function for call to 'kernel'}}
+                               // expected-note@-24 {{candidate template ignored: invalid explicitly-specified argument for template parameter 'devF'}}
+  kernel<&devtempfun<int> ><<<1,1>>>();
+  kernel<&devhostfun><<<1,1>>>();
+
+  kernel<devfun><<<1,1>>>();
+  kernel<hostfun><<<1,1>>>(); // expected-error {{no matching function for call to 'kernel'}}
+                              // expected-note@-30 {{candidate template ignored: invalid explicitly-specified argument for template parameter 'devF'}}
+  kernel<devtempfun<int> ><<<1,1>>>();
+  kernel<devhostfun><<<1,1>>>();
+
+  C a;
+  kernel2<C, &C::devfun><<<1,1>>>(&a);
+  kernel2<C, &C::hostfun><<<1,1>>>(&a); // expected-error {{no matching function for call to 'kernel2'}}
+                                        // expected-note@-36 {{candidate template ignored: invalid explicitly-specified argument for template parameter 'devF'}}
+  kernel2<C, &C::devtempfun<int> ><<<1,1>>>(&a);
+  kernel2<C, &C::devhostfun><<<1,1>>>(&a);
+}

clang/test/SemaOpenCL/amdgpu-attrs.cl

 } struct_num_vgpr_64;
 
 __attribute__((amdgpu_flat_work_group_size(32, 64))) void func_flat_work_group_size_32_64() {} // expected-error {{'amdgpu_flat_work_group_size' attribute only applies to kernel functions}}
 __attribute__((amdgpu_waves_per_eu(2))) void func_waves_per_eu_2() {} // expected-error {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
 __attribute__((amdgpu_waves_per_eu(2, 4))) void func_waves_per_eu_2_4() {} // expected-error {{'amdgpu_waves_per_eu' attribute only applies to kernel functions}}
 __attribute__((amdgpu_num_sgpr(32))) void func_num_sgpr_32() {} // expected-error {{'amdgpu_num_sgpr' attribute only applies to kernel functions}}
 __attribute__((amdgpu_num_vgpr(64))) void func_num_vgpr_64() {} // expected-error {{'amdgpu_num_vgpr' attribute only applies to kernel functions}}
 
-__attribute__((amdgpu_flat_work_group_size("ABC", "ABC"))) kernel void kernel_flat_work_group_size_ABC_ABC() {} // expected-error {{'amdgpu_flat_work_group_size' attribute requires an integer constant}}
-__attribute__((amdgpu_flat_work_group_size(32, "ABC"))) kernel void kernel_flat_work_group_size_32_ABC() {} // expected-error {{'amdgpu_flat_work_group_size' attribute requires an integer constant}}
-__attribute__((amdgpu_flat_work_group_size("ABC", 64))) kernel void kernel_flat_work_group_size_ABC_64() {} // expected-error {{'amdgpu_flat_work_group_size' attribute requires an integer constant}}
-__attribute__((amdgpu_waves_per_eu("ABC"))) kernel void kernel_waves_per_eu_ABC() {} // expected-error {{'amdgpu_waves_per_eu' attribute requires an integer constant}}
-__attribute__((amdgpu_waves_per_eu(2, "ABC"))) kernel void kernel_waves_per_eu_2_ABC() {} // expected-error {{'amdgpu_waves_per_eu' attribute requires an integer constant}}
-__attribute__((amdgpu_waves_per_eu("ABC", 4))) kernel void kernel_waves_per_eu_ABC_4() {} // expected-error {{'amdgpu_waves_per_eu' attribute requires an integer constant}}
+__attribute__((amdgpu_flat_work_group_size("ABC", "ABC"))) kernel void kernel_flat_work_group_size_ABC_ABC() {} // expected-error {{'amdgpu_flat_work_group_size' attribute requires parameter 0 to be an integer constant}}
+__attribute__((amdgpu_flat_work_group_size(32, "ABC"))) kernel void kernel_flat_work_group_size_32_ABC() {} // expected-error {{'amdgpu_flat_work_group_size' attribute requires parameter 1 to be an integer constant}}
+__attribute__((amdgpu_flat_work_group_size("ABC", 64))) kernel void kernel_flat_work_group_size_ABC_64() {} // expected-error {{'amdgpu_flat_work_group_size' attribute requires parameter 0 to be an integer constant}}
+__attribute__((amdgpu_waves_per_eu("ABC"))) kernel void kernel_waves_per_eu_ABC() {} // expected-error {{'amdgpu_waves_per_eu' attribute requires parameter 0 to be an integer constant}}
+__attribute__((amdgpu_waves_per_eu(2, "ABC"))) kernel void kernel_waves_per_eu_2_ABC() {} // expected-error {{'amdgpu_waves_per_eu' attribute requires parameter 1 to be an integer constant}}
+__attribute__((amdgpu_waves_per_eu("ABC", 4))) kernel void kernel_waves_per_eu_ABC_4() {} // expected-error {{'amdgpu_waves_per_eu' attribute requires parameter 0 to be an integer constant}}
 __attribute__((amdgpu_num_sgpr("ABC"))) kernel void kernel_num_sgpr_ABC() {} // expected-error {{'amdgpu_num_sgpr' attribute requires an integer constant}}
 __attribute__((amdgpu_num_vgpr("ABC"))) kernel void kernel_num_vgpr_ABC() {} // expected-error {{'amdgpu_num_vgpr' attribute requires an integer constant}}
 
 __attribute__((amdgpu_flat_work_group_size(4294967296, 4294967296))) kernel void kernel_flat_work_group_size_L_L() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
 __attribute__((amdgpu_flat_work_group_size(32, 4294967296))) kernel void kernel_flat_work_group_size_32_L() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
 __attribute__((amdgpu_flat_work_group_size(4294967296, 64))) kernel void kernel_flat_work_group_size_L_64() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
 __attribute__((amdgpu_waves_per_eu(4294967296))) kernel void kernel_waves_per_eu_L() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
 __attribute__((amdgpu_waves_per_eu(2, 4294967296))) kernel void kernel_waves_per_eu_2_L() {} // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}

llvm/lib/CodeGen/SelectionDAG/InstrEmitter.cpp

void InstrEmitter::AddOperand(MachineInstrBuilder &MIB,
   } else if (ConstantFPSDNode *F = dyn_cast<ConstantFPSDNode>(Op)) {
     MIB.addFPImm(F->getConstantFPValue());
   } else if (RegisterSDNode *R = dyn_cast<RegisterSDNode>(Op)) {
     unsigned VReg = R->getReg();
     MVT OpVT = Op.getSimpleValueType();
     const TargetRegisterClass *OpRC =
         TLI->isTypeLegal(OpVT) ? TLI->getRegClassFor(OpVT) : nullptr;
     const TargetRegisterClass *IIRC =
-        II ? TRI->getAllocatableClass(TII->getRegClass(*II, IIOpNum, TRI, *MF))
-           : nullptr;
+        II ? TII->getRegClass(*II, IIOpNum, TRI, *MF) : nullptr;
+    assert(!II || IIOpNum < II->getNumOperands() || !IIRC);
+    IIRC = TRI->getAllocatableClass(IIRC);
 
     if (OpRC && IIRC && OpRC != IIRC &&
         TargetRegisterInfo::isVirtualRegister(VReg)) {
       unsigned NewVReg = MRI->createVirtualRegister(IIRC);
       BuildMI(*MBB, InsertPos, Op.getNode()->getDebugLoc(),
                TII->get(TargetOpcode::COPY), NewVReg).addReg(VReg);
       VReg = NewVReg;
     }