[OpenACC] implement AST/Sema for 'routine' construct with argument

The 'routine' construct has two forms, one which takes the name of a
function that it applies to, and another where it implicitly figures it
out based on the next declaration. This patch implements the former with
the required restrictions on the name and the function-static-variables
as specified.

What has not been implemented is any clauses for this, any of the A.3.4
warnings, or the other form.

Author: erichkeane

clang/include/clang/AST/DeclOpenACC.h

@@ -101,6 +101,61 @@ class OpenACCDeclareDecl final
   static bool classof(const Decl *D) { return classofKind(D->getKind()); }
   static bool classofKind(Kind K) { return K == OpenACCDeclare; }
 };
+
+class OpenACCRoutineDecl final
+    : public OpenACCConstructDecl,
+      private llvm::TrailingObjects<OpenACCRoutineDecl, const OpenACCClause *> {
+  friend TrailingObjects;
+  friend class ASTDeclReader;
+  friend class ASTDeclWriter;
+
+  Expr *FuncRef = nullptr;
+  SourceRange ParensLoc;
+
+  OpenACCRoutineDecl(unsigned NumClauses)
+      : OpenACCConstructDecl(OpenACCRoutine) {
+    std::uninitialized_value_construct(
+        getTrailingObjects<const OpenACCClause *>(),
+        getTrailingObjects<const OpenACCClause *>() + NumClauses);
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  NumClauses));
+  }
+
+  OpenACCRoutineDecl(DeclContext *DC, SourceLocation StartLoc,
+                     SourceLocation DirLoc, SourceLocation LParenLoc,
+                     Expr *FuncRef, SourceLocation RParenLoc,
+                     SourceLocation EndLoc,
+                     ArrayRef<const OpenACCClause *> Clauses)
+      : OpenACCConstructDecl(OpenACCRoutine, DC, OpenACCDirectiveKind::Routine,
+                             StartLoc, DirLoc, EndLoc),
+        FuncRef(FuncRef), ParensLoc(LParenLoc, RParenLoc) {
+    // Initialize the trailing storage.
+    std::uninitialized_copy(Clauses.begin(), Clauses.end(),
+                            getTrailingObjects<const OpenACCClause *>());
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  Clauses.size()));
+  }
+
+public:
+  static OpenACCRoutineDecl *
+  Create(ASTContext &Ctx, DeclContext *DC, SourceLocation StartLoc,
+         SourceLocation DirLoc, SourceLocation LParenLoc, Expr *FuncRef,
+         SourceLocation RParenLoc, SourceLocation EndLoc,
+         ArrayRef<const OpenACCClause *> Clauses);
+  static OpenACCRoutineDecl *
+  CreateDeserialized(ASTContext &Ctx, GlobalDeclID ID, unsigned NumClauses);
+  static bool classof(const Decl *D) { return classofKind(D->getKind()); }
+  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
 
 #endif

clang/include/clang/AST/JSONNodeDumper.h

@@ -282,6 +282,7 @@ class JSONNodeDumper
   void VisitBlockDecl(const BlockDecl *D);
 
   void VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D);
+  void VisitOpenACCRoutineDecl(const OpenACCRoutineDecl *D);
 
   void VisitDeclRefExpr(const DeclRefExpr *DRE);
   void VisitSYCLUniqueStableNameExpr(const SYCLUniqueStableNameExpr *E);

clang/include/clang/AST/RecursiveASTVisitor.h

@@ -1825,6 +1825,11 @@ DEF_TRAVERSE_DECL(OMPAllocateDecl, {
 DEF_TRAVERSE_DECL(OpenACCDeclareDecl,
                   { TRY_TO(VisitOpenACCClauseList(D->clauses())); })
 
+DEF_TRAVERSE_DECL(OpenACCRoutineDecl, {
+  TRY_TO(TraverseStmt(D->getFunctionReference()));
+  TRY_TO(VisitOpenACCClauseList(D->clauses()));
+})
+
 // A helper method for TemplateDecl's children.
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::TraverseTemplateParameterListHelper(

clang/include/clang/AST/TextNodeDumper.h

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

clang/include/clang/Basic/Attr.td

@@ -5026,3 +5026,11 @@ def Atomic : StmtAttr {
   let Documentation = [AtomicDocs];
   let StrictEnumParameters = 1;
 }
+
+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.
+  let Spellings = [];
+  let Subjects = SubjectList<[Function]>;
+  let Documentation = [InternalOnly];
+}

clang/include/clang/Basic/DeclNodes.td

@@ -112,3 +112,4 @@ def RequiresExprBody : DeclNode<Decl>, DeclContext;
 def LifetimeExtendedTemporary : DeclNode<Decl>;
 def HLSLBuffer : DeclNode<Named, "HLSLBuffer">, DeclContext;
 def OpenACCDeclare : DeclNode<Decl, "#pragma acc declare">;
+def OpenACCRoutine : DeclNode<Decl, "#pragma acc routine">;

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12821,8 +12821,9 @@ def err_wasm_builtin_arg_must_be_integer_type : Error <
   "%ordinal0 argument must be an integer">;
 
 // OpenACC diagnostics.
-def warn_acc_construct_unimplemented
-    : Warning<"OpenACC construct '%0' not yet implemented, pragma ignored">,
+def warn_acc_routine_unimplemented
+    : Warning<"OpenACC construct 'routine' with implicit function not yet "
+              "implemented, pragma ignored">,
       InGroup<SourceUsesOpenACC>;
 def warn_acc_clause_unimplemented
     : Warning<"OpenACC clause '%0' not yet implemented, clause ignored">,
@@ -13060,6 +13061,13 @@ def err_acc_declare_same_scope
 def err_acc_multiple_references
     : Error<"variable referenced in '%0' clause of OpenACC 'declare' directive "
             "was already referenced">;
+def err_acc_routine_not_func
+    : Error<"OpenACC routine name '%0' does not name a function">;
+def err_acc_routine_overload_set
+    : Error<"OpenACC routine name '%0' names a set of overloads">;
+def err_acc_magic_static_in_routine
+    : Error<"function static variables are not permitted in functions to which "
+            "an OpenACC 'routine' directive applies">;
 
 // AMDGCN builtins diagnostics
 def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size value">;

clang/include/clang/Sema/Sema.h

@@ -13436,6 +13436,10 @@ class Sema final : public SemaBase {
                             bool ForCallExpr = false);
   ExprResult SubstExpr(Expr *E,
                        const MultiLevelTemplateArgumentList &TemplateArgs);
+  /// Substitute an expression as if it is a address-of-operand, which makes it
+  /// act like a CXXIdExpression rather than an attempt to call.
+  ExprResult SubstCXXIdExpr(Expr *E,
+                            const MultiLevelTemplateArgumentList &TemplateArgs);
 
   // A RAII type used by the TemplateDeclInstantiator and TemplateInstantiator
   // to disable constraint evaluation, then restore the state.

clang/include/clang/Sema/SemaOpenACC.h

@@ -167,6 +167,11 @@ class SemaOpenACC : public SemaBase {
   // contexts, we can just store the declaration and location of the reference.
   llvm::DenseMap<const clang::DeclaratorDecl *, SourceLocation>
       DeclareVarReferences;
+  // The 'routine' construct disallows magic-statics in a function referred to
+  // by a 'routine' directive.  So record any of these that we see so we can
+  // check them later.
+  llvm::SmallDenseMap<const clang::FunctionDecl *, SourceLocation>
+      MagicStaticLocs;
 
 public:
   ComputeConstructInfo &getActiveComputeConstructInfo() {
@@ -723,6 +728,7 @@ class SemaOpenACC : public SemaBase {
   /// MiscLoc: First misc location, if necessary (not all constructs).
   /// Exprs: List of expressions on the construct itself, if necessary (not all
   /// constructs).
+  /// FuncRef: used only for Routine, this is the function being referenced.
   /// AK: The atomic kind of the directive, if necessary (atomic only)
   /// RParenLoc: Location of the right paren, if it exists (not on all
   /// constructs).
@@ -735,23 +741,12 @@ class SemaOpenACC : public SemaBase {
       OpenACCAtomicKind AK, SourceLocation RParenLoc, SourceLocation EndLoc,
       ArrayRef<OpenACCClause *> Clauses, StmtResult AssocStmt);
 
-  StmtResult ActOnEndStmtDirective(
-      OpenACCDirectiveKind K, SourceLocation StartLoc, SourceLocation DirLoc,
-      SourceLocation LParenLoc, SourceLocation MiscLoc, ArrayRef<Expr *> Exprs,
-      SourceLocation RParenLoc, SourceLocation EndLoc,
-      ArrayRef<OpenACCClause *> Clauses, StmtResult AssocStmt) {
-    return ActOnEndStmtDirective(K, StartLoc, DirLoc, LParenLoc, MiscLoc, Exprs,
-                                 OpenACCAtomicKind::None, RParenLoc, EndLoc,
-                                 Clauses, AssocStmt);
-  }
-
   /// Called after the directive has been completely parsed, including the
   /// declaration group or associated statement.
-  DeclGroupRef ActOnEndDeclDirective(OpenACCDirectiveKind K,
-                                     SourceLocation StartLoc,
-                                     SourceLocation DirLoc,
-                                     SourceLocation EndLoc,
-                                     ArrayRef<OpenACCClause *> Clauses);
+  DeclGroupRef ActOnEndDeclDirective(
+      OpenACCDirectiveKind K, SourceLocation StartLoc, SourceLocation DirLoc,
+      SourceLocation LParenLoc, Expr *FuncRef, SourceLocation RParenLoc,
+      SourceLocation EndLoc, ArrayRef<OpenACCClause *> Clauses);
 
   /// Called when encountering an 'int-expr' for OpenACC, and manages
   /// conversions and diagnostics to 'int'.
@@ -764,6 +759,9 @@ class SemaOpenACC : public SemaBase {
                       Expr *VarExpr);
   /// 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.
+  void ActOnVariableDeclarator(VarDecl *VD);
 
   // Called after 'ActOnVar' specifically for a 'link' clause, which has to do
   // some minor additional checks.
@@ -773,6 +771,8 @@ class SemaOpenACC : public SemaBase {
   // checked during both phases of template translation.
   bool CheckDeclareClause(SemaOpenACC::OpenACCParsedClause &Clause);
 
+  ExprResult ActOnRoutineName(Expr *RoutineName);
+
   /// Called while semantically analyzing the reduction clause, ensuring the var
   /// is the correct kind of reference.
   ExprResult CheckReductionVar(OpenACCDirectiveKind DirectiveKind,

clang/include/clang/Serialization/ASTBitCodes.h

@@ -1531,7 +1531,10 @@ enum DeclCode {
   // An OpenACCDeclareDecl record.
   DECL_OPENACC_DECLARE,
 
-  DECL_LAST = DECL_OPENACC_DECLARE
+  // An OpenACCRoutineDecl record.
+  DECL_OPENACC_ROUTINE,
+
+  DECL_LAST = DECL_OPENACC_ROUTINE
 };
 
 /// Record codes for each kind of statement or expression.

clang/lib/AST/DeclBase.cpp

@@ -994,6 +994,7 @@ unsigned Decl::getIdentifierNamespaceForKind(Kind DeclKind) {
     case RequiresExprBody:
     case ImplicitConceptSpecialization:
     case OpenACCDeclare:
+    case OpenACCRoutine:
       // Never looked up by name.
       return 0;
   }

clang/lib/AST/DeclOpenACC.cpp

@@ -31,3 +31,22 @@ OpenACCDeclareDecl::CreateDeserialized(ASTContext &Ctx, GlobalDeclID ID,
   return new (Ctx, ID, additionalSizeToAlloc<const OpenACCClause *>(NumClauses))
       OpenACCDeclareDecl(NumClauses);
 }
+
+OpenACCRoutineDecl *
+OpenACCRoutineDecl::Create(ASTContext &Ctx, DeclContext *DC,
+                           SourceLocation StartLoc, SourceLocation DirLoc,
+                           SourceLocation LParenLoc, Expr *FuncRef,
+                           SourceLocation RParenLoc, SourceLocation EndLoc,
+                           ArrayRef<const OpenACCClause *> Clauses) {
+  return new (Ctx, DC,
+              additionalSizeToAlloc<const OpenACCClause *>(Clauses.size()))
+      OpenACCRoutineDecl(DC, StartLoc, DirLoc, LParenLoc, FuncRef, RParenLoc,
+                         EndLoc, Clauses);
+}
+
+OpenACCRoutineDecl *
+OpenACCRoutineDecl::CreateDeserialized(ASTContext &Ctx, GlobalDeclID ID,
+                                       unsigned NumClauses) {
+  return new (Ctx, ID, additionalSizeToAlloc<const OpenACCClause *>(NumClauses))
+      OpenACCRoutineDecl(NumClauses);
+}

clang/lib/AST/DeclPrinter.cpp

@@ -113,6 +113,7 @@ namespace {
     void VisitHLSLBufferDecl(HLSLBufferDecl *D);
 
     void VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D);
+    void VisitOpenACCRoutineDecl(OpenACCRoutineDecl *D);
 
     void printTemplateParameters(const TemplateParameterList *Params,
                                  bool OmitTemplateKW = false);
@@ -497,7 +498,7 @@ void DeclPrinter::VisitDeclContext(DeclContext *DC, bool Indent) {
         isa<OMPDeclareMapperDecl>(*D) || isa<OMPRequiresDecl>(*D) ||
         isa<OMPAllocateDecl>(*D))
       Terminator = nullptr;
-    else if (isa<OpenACCDeclareDecl>(*D))
+    else if (isa<OpenACCDeclareDecl, OpenACCRoutineDecl>(*D))
       Terminator = nullptr;
     else if (isa<ObjCMethodDecl>(*D) && cast<ObjCMethodDecl>(*D)->hasBody())
       Terminator = nullptr;
@@ -1922,3 +1923,25 @@ void DeclPrinter::VisitOpenACCDeclareDecl(OpenACCDeclareDecl *D) {
     Printer.VisitClauseList(D->clauses());
   }
 }
+void DeclPrinter::VisitOpenACCRoutineDecl(OpenACCRoutineDecl *D) {
+  if (!D->isInvalidDecl()) {
+    Out << "#pragma acc routine";
+
+    if (D->hasNameSpecified()) {
+      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>";
+
+      Out << ")";
+    }
+
+    OpenACCClausePrinter Printer(Out, Policy);
+    Printer.VisitClauseList(D->clauses());
+  }
+}

clang/lib/AST/JSONNodeDumper.cpp

@@ -1363,6 +1363,7 @@ void JSONNodeDumper::VisitOpenACCAsteriskSizeExpr(
     const OpenACCAsteriskSizeExpr *E) {}
 
 void JSONNodeDumper::VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D) {}
+void JSONNodeDumper::VisitOpenACCRoutineDecl(const OpenACCRoutineDecl *D) {}
 
 void JSONNodeDumper::VisitPredefinedExpr(const PredefinedExpr *PE) {
   JOS.attribute("name", PredefinedExpr::getIdentKindName(PE->getIdentKind()));

clang/lib/AST/StmtPrinter.cpp

@@ -264,7 +264,8 @@ void StmtPrinter::VisitDeclStmt(DeclStmt *Node) {
   Indent();
   PrintRawDeclStmt(Node);
   // Certain pragma declarations shouldn't have a semi-colon after them.
-  if (!Node->isSingleDecl() || !isa<OpenACCDeclareDecl>(Node->getSingleDecl()))
+  if (!Node->isSingleDecl() ||
+      !isa<OpenACCDeclareDecl, OpenACCRoutineDecl>(Node->getSingleDecl()))
     OS << ";";
   OS << NL;
 }

clang/lib/AST/TextNodeDumper.cpp

@@ -3079,6 +3079,23 @@ void TextNodeDumper::VisitOpenACCDeclareDecl(const OpenACCDeclareDecl *D) {
         AddChild([=] { Visit(S); });
     });
 }
+void TextNodeDumper::VisitOpenACCRoutineDecl(const OpenACCRoutineDecl *D) {
+  OS << " " << D->getDirectiveKind();
+
+  if (D->hasNameSpecified()) {
+    OS << " name_specified";
+    dumpSourceRange(SourceRange{D->getLParenLoc(), D->getRParenLoc()});
+  }
+
+  AddChild([=] { Visit(D->getFunctionReference()); });
+
+  for (const OpenACCClause *C : D->clauses())
+    AddChild([=] {
+      Visit(C);
+      for (const Stmt *S : C->children())
+        AddChild([=] { Visit(S); });
+    });
+}
 
 void TextNodeDumper::VisitEmbedExpr(const EmbedExpr *S) {
   AddChild("begin", [=] { OS << S->getStartingElementPos(); });

clang/lib/CodeGen/CGDecl.cpp

@@ -180,6 +180,8 @@ void CodeGenFunction::EmitDecl(const Decl &D) {
 
   case Decl::OpenACCDeclare:
     return CGM.EmitOpenACCDeclare(cast<OpenACCDeclareDecl>(&D), this);
+  case Decl::OpenACCRoutine:
+    return CGM.EmitOpenACCRoutine(cast<OpenACCRoutineDecl>(&D), this);
 
   case Decl::Typedef:      // typedef int X;
   case Decl::TypeAlias: {  // using X = int; [C++0x]
@@ -2852,6 +2854,11 @@ void CodeGenModule::EmitOpenACCDeclare(const OpenACCDeclareDecl *D,
   // This is a no-op, we cna just ignore these declarations.
 }
 
+void CodeGenModule::EmitOpenACCRoutine(const OpenACCRoutineDecl *D,
+                                       CodeGenFunction *CGF) {
+  // This is a no-op, we cna just ignore these declarations.
+}
+
 void CodeGenModule::EmitOMPRequiresDecl(const OMPRequiresDecl *D) {
   getOpenMPRuntime().processRequiresDirective(D);
 }

clang/lib/CodeGen/CodeGenModule.h

@@ -1575,6 +1575,8 @@ class CodeGenModule : public CodeGenTypeCache {
 
   // Emit code for the OpenACC Declare declaration.
   void EmitOpenACCDeclare(const OpenACCDeclareDecl *D, CodeGenFunction *CGF);
+  // Emit code for the OpenACC Routine declaration.
+  void EmitOpenACCRoutine(const OpenACCRoutineDecl *D, CodeGenFunction *CGF);
 
   /// Emit a code for requires directive.
   /// \param D Requires declaration