[OpenACC] Finish implementing 'routine' AST/Sema.

This is the last item of the OpenACC 3.3 spec. It includes the
implicit-name version of 'routine', plus significant refactorings to
make the two work together.  The implicit name version is represented as
an attribute on the function call. This patch also implements the
clauses for the implicit-name version, as well as the A.3.4 warning.

Author: erichkeane

clang/include/clang/AST/Attr.h

@@ -38,6 +38,7 @@ class ASTContext;
 class AttributeCommonInfo;
 class FunctionDecl;
 class OMPTraitInfo;
+class OpenACCClause;
 
 /// Attr - This represents one attribute.
 class Attr : public AttributeCommonInfo {

clang/include/clang/AST/DeclOpenACC.h

@@ -102,6 +102,8 @@ class OpenACCDeclareDecl final
   static bool classofKind(Kind K) { return K == OpenACCDeclare; }
 };
 
+// Reprents a 'routine' directive with a name. When this has no name, it is
+// represented as an attribute.
 class OpenACCRoutineDecl final
     : public OpenACCConstructDecl,
       private llvm::TrailingObjects<OpenACCRoutineDecl, const OpenACCClause *> {
@@ -129,6 +131,8 @@ class OpenACCRoutineDecl final
       : OpenACCConstructDecl(OpenACCRoutine, DC, OpenACCDirectiveKind::Routine,
                              StartLoc, DirLoc, EndLoc),
         FuncRef(FuncRef), ParensLoc(LParenLoc, RParenLoc) {
+    assert(LParenLoc.isValid() &&
+           "Cannot represent implicit name with this declaration");
     // Initialize the trailing storage.
     std::uninitialized_copy(Clauses.begin(), Clauses.end(),
                             getTrailingObjects<const OpenACCClause *>());
@@ -148,13 +152,10 @@ class OpenACCRoutineDecl final
   static bool classofKind(Kind K) { return K == OpenACCRoutine; }
 
   const Expr *getFunctionReference() const { return FuncRef; }
-
   Expr *getFunctionReference() { return FuncRef; }
 
   SourceLocation getLParenLoc() const { return ParensLoc.getBegin(); }
   SourceLocation getRParenLoc() const { return ParensLoc.getEnd(); }
-
-  bool hasNameSpecified() const { return !ParensLoc.getBegin().isInvalid(); }
 };
 } // namespace clang
 

clang/include/clang/AST/OpenACCClause.h

@@ -251,6 +251,12 @@ class OpenACCBindClause final : public OpenACCClauseWithParams {
   }
 };
 
+bool operator==(const OpenACCBindClause &LHS, const OpenACCBindClause &RHS);
+inline bool operator!=(const OpenACCBindClause &LHS,
+                       const OpenACCBindClause &RHS) {
+  return !(LHS == RHS);
+}
+
 using DeviceTypeArgument = std::pair<IdentifierInfo *, SourceLocation>;
 /// A 'device_type' or 'dtype' clause, takes a list of either an 'asterisk' or
 /// an identifier. The 'asterisk' means 'the rest'.

clang/include/clang/AST/TextNodeDumper.h

@@ -426,6 +426,7 @@ class TextNodeDumper
   void VisitOpenACCAsteriskSizeExpr(const OpenACCAsteriskSizeExpr *S);
   void VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D);
   void VisitOpenACCRoutineDecl(const OpenACCRoutineDecl *D);
+  void VisitOpenACCRoutineDeclAttr(const OpenACCRoutineDeclAttr *A);
   void VisitEmbedExpr(const EmbedExpr *S);
   void VisitAtomicExpr(const AtomicExpr *AE);
   void VisitConvertVectorExpr(const ConvertVectorExpr *S);

clang/include/clang/Basic/Attr.td

@@ -693,6 +693,10 @@ class Attr {
   // content. Eg) It parses 3 args, but semantically takes 4 args.  Opts out of
   // common attribute error checking.
   bit HasCustomParsing = 0;
+  // Set to true if this attribute requires custom serialization after the
+  // typical attribute serialization. This will cause tablegen to emit a call to
+  // ASTRecordWriter::Add<Name>Attr and ASTRecordReader::read<Name>Attr.
+  bit HasCustomSerialization = 0;
   // Set to true if all of the attribute's arguments should be parsed in an
   // unevaluated context.
   bit ParseArgumentsAsUnevaluated = 0;
@@ -4991,12 +4995,30 @@ def Atomic : StmtAttr {
 }
 
 def OpenACCRoutineAnnot : InheritableAttr {
-  // This attribute is used to mark that a function is targetted by a `routine`
-  // directive, so it dones't have a spelling and is always implicit.
+  // This attribute is used to mark that a function is targeted by a `routine`
+  // directive with a name for the purposes of checking the declaration later.
+  // We don't really need a link back to the declaration, as location is
+  // sufficient.
+  // We abuse the source locations on this a little, since we need two locations
+  // for various diagnostic purposes.  The 'begin' location is the location of
+  // the Routine directive. We are using the 'end' location for any 'bind'
+  // clauses, since this is needed for a diagnostic.
   let Spellings = [];
   let Subjects = SubjectList<[Function]>;
   let Documentation = [InternalOnly];
+}
+
+def OpenACCRoutineDecl :InheritableAttr {
+  // This attribute represents the 'routine' directive when spelled without a
+  // 'name'.
+  let Spellings = [Pragma<"acc", "routine">];
+  let Subjects = SubjectList<[Function]>;
+  let SemaHandler = 0;
+  let HasCustomParsing = 1;
+  let HasCustomSerialization = 1;
+  let Documentation = [InternalOnly];
   let AdditionalMembers = [{
-    SourceLocation BindClause;
+    llvm::SmallVector<const OpenACCClause *> Clauses;
+    void printPrettyPragma(raw_ostream &OS, const PrintingPolicy &Policy) const;
   }];
 }

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -13040,6 +13040,16 @@ def err_acc_magic_static_in_routine
 def err_acc_duplicate_bind
     : Error<"multiple 'routine' directives with 'bind' clauses are not "
             "permitted to refer to the same function">;
+def err_acc_duplicate_unnamed_bind
+    : Error<"OpenACC 'bind' clause on a declaration must bind to the same name "
+            "as previous bind clauses">;
+def warn_acc_confusing_routine_name
+    : Warning<"OpenACC 'routine' directive with a name refers to a function "
+              "with the same name as the function on the following line; this "
+              "may be unintended">,
+      InGroup<DiagGroup<"openacc-confusing-routine-name">>;
+def err_acc_decl_for_routine
+    : Error<"expected function or lambda declaration for 'routine' construct">;
 
 // AMDGCN builtins diagnostics
 def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size value">;

clang/include/clang/Parse/Parser.h

@@ -3697,7 +3697,10 @@ class Parser : public CodeCompletionHandler {
   /// diagnostic. Eventually will be split into a few functions to parse
   /// different situations.
 public:
-  DeclGroupPtrTy ParseOpenACCDirectiveDecl();
+  DeclGroupPtrTy ParseOpenACCDirectiveDecl(AccessSpecifier &AS,
+                                           ParsedAttributes &Attrs,
+                                           DeclSpec::TST TagType,
+                                           Decl *TagDecl);
   StmtResult ParseOpenACCDirectiveStmt();
 
 private:
@@ -3831,6 +3834,11 @@ class Parser : public CodeCompletionHandler {
   OpenACCGangArgRes ParseOpenACCGangArg(SourceLocation GangLoc);
   /// Parses a 'condition' expr, ensuring it results in a
   ExprResult ParseOpenACCConditionExpr();
+  DeclGroupPtrTy
+  ParseOpenACCAfterRoutineDecl(AccessSpecifier &AS, ParsedAttributes &Attrs,
+                               DeclSpec::TST TagType, Decl *TagDecl,
+                               OpenACCDirectiveParseInfo &DirInfo);
+  StmtResult ParseOpenACCAfterRoutineStmt(OpenACCDirectiveParseInfo &DirInfo);
 
 private:
   //===--------------------------------------------------------------------===//

clang/include/clang/Sema/SemaOpenACC.h

@@ -31,6 +31,7 @@
 namespace clang {
 class IdentifierInfo;
 class OpenACCClause;
+class Scope;
 
 class SemaOpenACC : public SemaBase {
 public:
@@ -172,6 +173,9 @@ class SemaOpenACC : public SemaBase {
   // check them later.
   llvm::SmallDenseMap<const clang::FunctionDecl *, SourceLocation>
       MagicStaticLocs;
+  OpenACCRoutineDecl *LastRoutineDecl = nullptr;
+
+  void CheckLastRoutineDeclNameConflict(const NamedDecl *ND);
 
 public:
   ComputeConstructInfo &getActiveComputeConstructInfo() {
@@ -763,10 +767,37 @@ class SemaOpenACC : public SemaBase {
 
   /// Called after the directive has been completely parsed, including the
   /// declaration group or associated statement.
-  DeclGroupRef ActOnEndDeclDirective(
-      OpenACCDirectiveKind K, SourceLocation StartLoc, SourceLocation DirLoc,
-      SourceLocation LParenLoc, Expr *FuncRef, SourceLocation RParenLoc,
-      SourceLocation EndLoc, ArrayRef<OpenACCClause *> Clauses);
+  DeclGroupRef
+  ActOnEndDeclDirective(OpenACCDirectiveKind K, SourceLocation StartLoc,
+                        SourceLocation DirLoc, SourceLocation LParenLoc,
+                        SourceLocation RParenLoc, SourceLocation EndLoc,
+                        ArrayRef<OpenACCClause *> Clauses);
+
+  // Helper functions for ActOnEndRoutine*Directive, which does all the checking
+  // given the proper list of declarations.
+  void CheckRoutineDecl(SourceLocation DirLoc,
+                        ArrayRef<const OpenACCClause *> Clauses,
+                        Decl *NextParsedDecl);
+  OpenACCRoutineDecl *CheckRoutineDecl(SourceLocation StartLoc,
+                                       SourceLocation DirLoc,
+                                       SourceLocation LParenLoc, Expr *FuncRef,
+                                       SourceLocation RParenLoc,
+                                       ArrayRef<const OpenACCClause *> Clauses,
+                                       SourceLocation EndLoc);
+  OpenACCRoutineDeclAttr *
+  mergeRoutineDeclAttr(const OpenACCRoutineDeclAttr &Old);
+  DeclGroupRef
+  ActOnEndRoutineDeclDirective(SourceLocation StartLoc, SourceLocation DirLoc,
+                               SourceLocation LParenLoc, Expr *ReferencedFunc,
+                               SourceLocation RParenLoc,
+                               ArrayRef<const OpenACCClause *> Clauses,
+                               SourceLocation EndLoc, DeclGroupPtrTy NextDecl);
+  StmtResult
+  ActOnEndRoutineStmtDirective(SourceLocation StartLoc, SourceLocation DirLoc,
+                               SourceLocation LParenLoc, Expr *ReferencedFunc,
+                               SourceLocation RParenLoc,
+                               ArrayRef<const OpenACCClause *> Clauses,
+                               SourceLocation EndLoc, Stmt *NextStmt);
 
   /// Called when encountering an 'int-expr' for OpenACC, and manages
   /// conversions and diagnostics to 'int'.
@@ -780,8 +811,14 @@ class SemaOpenACC : public SemaBase {
   /// Helper function called by ActonVar that is used to check a 'cache' var.
   ExprResult ActOnCacheVar(Expr *VarExpr);
   /// Function called when a variable declarator is created, which lets us
-  /// impelment the 'routine' 'function static variables' restriction.
+  /// implement the 'routine' 'function static variables' restriction.
   void ActOnVariableDeclarator(VarDecl *VD);
+  /// Called when a function decl is created, which lets us implement the
+  /// 'routine' 'doesn't match next thing' warning.
+  void ActOnFunctionDeclarator(FunctionDecl *FD);
+  /// Called when a variable is initialized, so we can implement the 'routine
+  /// 'doesn't match the next thing' warning for lambda init.
+  void ActOnVariableInit(VarDecl *VD, QualType InitType);
 
   // Called after 'ActOnVar' specifically for a 'link' clause, which has to do
   // some minor additional checks.

clang/include/clang/Serialization/ASTRecordReader.h

@@ -282,6 +282,8 @@ class ASTRecordReader
   /// statement reading.
   void readOpenACCClauseList(MutableArrayRef<const OpenACCClause *> Clauses);
 
+  void readOpenACCRoutineDeclAttr(OpenACCRoutineDeclAttr *A);
+
   /// Read a source location, advancing Idx.
   SourceLocation readSourceLocation(LocSeq *Seq = nullptr) {
     return Reader->ReadSourceLocation(*F, Record, Idx, Seq);

clang/include/clang/Serialization/ASTRecordWriter.h

@@ -307,6 +307,8 @@ class ASTRecordWriter
   /// Writes out a list of OpenACC clauses.
   void writeOpenACCClauseList(ArrayRef<const OpenACCClause *> Clauses);
 
+  void AddOpenACCRoutineDeclAttr(const OpenACCRoutineDeclAttr *A);
+
   /// Emit a string.
   void AddString(StringRef Str) {
     return Writer->AddString(Str, *Record);

clang/lib/AST/DeclOpenACC.cpp

@@ -12,6 +12,8 @@
 
 #include "clang/AST/DeclOpenACC.h"
 #include "clang/AST/ASTContext.h"
+#include "clang/AST/Attr.h"
+#include "clang/AST/OpenACCClause.h"
 
 using namespace clang;
 
@@ -50,3 +52,12 @@ OpenACCRoutineDecl::CreateDeserialized(ASTContext &Ctx, GlobalDeclID ID,
   return new (Ctx, ID, additionalSizeToAlloc<const OpenACCClause *>(NumClauses))
       OpenACCRoutineDecl(NumClauses);
 }
+
+void OpenACCRoutineDeclAttr::printPrettyPragma(
+    llvm::raw_ostream &OS, const clang::PrintingPolicy &P) const {
+  if (Clauses.size() > 0) {
+    OS << ' ';
+    OpenACCClausePrinter Printer{OS, P};
+    Printer.VisitClauseList(Clauses);
+  }
+}

clang/lib/AST/DeclPrinter.cpp

@@ -49,6 +49,7 @@ namespace {
                              QualType T);
 
     void PrintObjCTypeParams(ObjCTypeParamList *Params);
+    void PrintOpenACCRoutineOnLambda(Decl *D);
 
   public:
     DeclPrinter(raw_ostream &Out, const PrintingPolicy &Policy,
@@ -292,10 +293,34 @@ bool DeclPrinter::prettyPrintAttributes(const Decl *D,
   return hasPrinted;
 }
 
+void DeclPrinter::PrintOpenACCRoutineOnLambda(Decl *D) {
+  CXXRecordDecl *CXXRD = nullptr;
+  if (const auto *VD = dyn_cast<VarDecl>(D)) {
+    if (const auto *Init = VD->getInit())
+      CXXRD = Init->getType().isNull() ? nullptr
+                                       : Init->getType()->getAsCXXRecordDecl();
+  } else if (const auto *FD = dyn_cast<FieldDecl>(D)) {
+    CXXRD =
+        FD->getType().isNull() ? nullptr : FD->getType()->getAsCXXRecordDecl();
+  }
+
+  if (!CXXRD || !CXXRD->isLambda())
+    return;
+
+  if (const auto *Call = CXXRD->getLambdaCallOperator()) {
+    for (auto *A : Call->specific_attrs<OpenACCRoutineDeclAttr>()) {
+      A->printPretty(Out, Policy);
+      Indent();
+    }
+  }
+}
+
 void DeclPrinter::prettyPrintPragmas(Decl *D) {
   if (Policy.PolishForDeclaration)
     return;
 
+  PrintOpenACCRoutineOnLambda(D);
+
   if (D->hasAttrs()) {
     AttrVec &Attrs = D->getAttrs();
     for (auto *A : Attrs) {
@@ -894,6 +919,7 @@ void DeclPrinter::VisitFriendDecl(FriendDecl *D) {
 }
 
 void DeclPrinter::VisitFieldDecl(FieldDecl *D) {
+  prettyPrintPragmas(D);
   // FIXME: add printing of pragma attributes if required.
   if (!Policy.SuppressSpecifiers && D->isMutable())
     Out << "mutable ";
@@ -1930,19 +1956,17 @@ void DeclPrinter::VisitOpenACCRoutineDecl(OpenACCRoutineDecl *D) {
   if (!D->isInvalidDecl()) {
     Out << "#pragma acc routine";
 
-    if (D->hasNameSpecified()) {
-      Out << "(";
+    Out << "(";
 
-      // The referenced function was named here, but this makes us tolerant of
-      // errors.
-      if (D->getFunctionReference())
-        D->getFunctionReference()->printPretty(Out, nullptr, Policy,
-                                               Indentation, "\n", &Context);
-      else
-        Out << "<error>";
+    // The referenced function was named here, but this makes us tolerant of
+    // errors.
+    if (D->getFunctionReference())
+      D->getFunctionReference()->printPretty(Out, nullptr, Policy, Indentation,
+                                             "\n", &Context);
+    else
+      Out << "<error>";
 
-      Out << ")";
-    }
+    Out << ")";
 
     if (!D->clauses().empty()) {
       Out << ' ';

clang/lib/AST/OpenACCClause.cpp

@@ -628,6 +628,18 @@ OpenACCBindClause *OpenACCBindClause::Create(const ASTContext &C,
   return new (Mem) OpenACCBindClause(BeginLoc, LParenLoc, ID, EndLoc);
 }
 
+bool clang::operator==(const OpenACCBindClause &LHS,
+                       const OpenACCBindClause &RHS) {
+  if (LHS.isStringArgument() != RHS.isStringArgument())
+    return false;
+
+  if (LHS.isStringArgument())
+    return LHS.getStringArgument()->getString() ==
+           RHS.getStringArgument()->getString();
+  return LHS.getIdentifierArgument()->getName() ==
+         RHS.getIdentifierArgument()->getName();
+}
+
 //===----------------------------------------------------------------------===//
 //  OpenACC clauses printing methods
 //===----------------------------------------------------------------------===//

clang/lib/AST/TextNodeDumper.cpp

@@ -3090,10 +3090,7 @@ void TextNodeDumper::VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D) {
 void TextNodeDumper::VisitOpenACCRoutineDecl(const OpenACCRoutineDecl *D) {
   OS << " " << D->getDirectiveKind();
 
-  if (D->hasNameSpecified()) {
-    OS << " name_specified";
-    dumpSourceRange(SourceRange{D->getLParenLoc(), D->getRParenLoc()});
-  }
+  dumpSourceRange(SourceRange{D->getLParenLoc(), D->getRParenLoc()});
 
   AddChild([=] { Visit(D->getFunctionReference()); });
 
@@ -3105,6 +3102,16 @@ void TextNodeDumper::VisitOpenACCRoutineDecl(const OpenACCRoutineDecl *D) {
     });
 }
 
+void TextNodeDumper::VisitOpenACCRoutineDeclAttr(
+    const OpenACCRoutineDeclAttr *A) {
+  for (const OpenACCClause *C : A->Clauses)
+    AddChild([=] {
+      Visit(C);
+      for (const Stmt *S : C->children())
+        AddChild([=] { Visit(S); });
+    });
+}
+
 void TextNodeDumper::VisitEmbedExpr(const EmbedExpr *S) {
   AddChild("begin", [=] { OS << S->getStartingElementPos(); });
   AddChild("number of elements", [=] { OS << S->getDataElementCount(); });