[OpenACC] Implement AST for OpenACC Compute Constructs

clang/include/clang-c/Index.h

@@ -2145,7 +2145,11 @@ enum CXCursorKind {
    */
   CXCursor_OMPScopeDirective = 306,
 
-  CXCursor_LastStmt = CXCursor_OMPScopeDirective,
+  /** OpenACC Compute Construct.
+   */
+  CXCursor_OpenACCComputeConstruct = 307,
+
+  CXCursor_LastStmt = CXCursor_OpenACCComputeConstruct,
 
   /**
    * Cursor that represents the translation unit itself.

clang/include/clang/AST/RecursiveASTVisitor.h

@@ -34,6 +34,7 @@
 #include "clang/AST/Stmt.h"
 #include "clang/AST/StmtCXX.h"
 #include "clang/AST/StmtObjC.h"
+#include "clang/AST/StmtOpenACC.h"
 #include "clang/AST/StmtOpenMP.h"
 #include "clang/AST/TemplateBase.h"
 #include "clang/AST/TemplateName.h"
@@ -505,6 +506,9 @@ template <typename Derived> class RecursiveASTVisitor {
   bool VisitOMPClauseWithPostUpdate(OMPClauseWithPostUpdate *Node);
 
   bool PostVisitStmt(Stmt *S);
+  bool TraverseOpenACCConstructStmt(OpenACCConstructStmt *S);
+  bool
+  TraverseOpenACCAssociatedStmtConstruct(OpenACCAssociatedStmtConstruct *S);
 };
 
 template <typename Derived>
@@ -3910,6 +3914,24 @@ bool RecursiveASTVisitor<Derived>::VisitOMPXBareClause(OMPXBareClause *C) {
   return true;
 }
 
+template <typename Derived>
+bool RecursiveASTVisitor<Derived>::TraverseOpenACCConstructStmt(
+    OpenACCConstructStmt *) {
+  // TODO OpenACC: When we implement clauses, ensure we traverse them here.
+  return true;
+}
+
+template <typename Derived>
+bool RecursiveASTVisitor<Derived>::TraverseOpenACCAssociatedStmtConstruct(
+    OpenACCAssociatedStmtConstruct *S) {
+  TRY_TO(TraverseOpenACCConstructStmt(S));
+  TRY_TO(TraverseStmt(S->getAssociatedStmt()));
+  return true;
+}
+
+DEF_TRAVERSE_STMT(OpenACCComputeConstruct,
+                  { TRY_TO(TraverseOpenACCAssociatedStmtConstruct(S)); })
+
 // FIXME: look at the following tricky-seeming exprs to see if we
 // need to recurse on anything.  These are ones that have methods
 // returning decls or qualtypes or nestednamespecifier -- though I'm

clang/include/clang/AST/StmtOpenACC.h

@@ -0,0 +1,142 @@
+//===- StmtOpenACC.h - Classes for OpenACC directives  ----------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+/// \file
+/// This file defines OpenACC AST classes for statement-level contructs.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_CLANG_AST_STMTOPENACC_H
+#define LLVM_CLANG_AST_STMTOPENACC_H
+
+#include "clang/AST/Stmt.h"
+#include "clang/Basic/OpenACCKinds.h"
+#include "clang/Basic/SourceLocation.h"
+
+namespace clang {
+/// This is the base class for an OpenACC statement-level construct, other
+/// construct types are expected to inherit from this.
+class OpenACCConstructStmt : public Stmt {
+  friend class ASTStmtWriter;
+  friend class ASTStmtReader;
+  /// The directive kind. Each implementation of this interface should handle
+  /// specific kinds.
+  OpenACCDirectiveKind Kind = OpenACCDirectiveKind::Invalid;
+  /// The location of the directive statement, from the '#' to the last token of
+  /// the directive.
+  SourceRange Range;
+
+  // TODO OPENACC: Clauses should probably be collected in this class.
+
+protected:
+  OpenACCConstructStmt(StmtClass SC, OpenACCDirectiveKind K,
+                       SourceLocation Start, SourceLocation End)
+      : Stmt(SC), Kind(K), Range(Start, End) {}
+
+public:
+  OpenACCDirectiveKind getDirectiveKind() const { return Kind; }
+
+  static bool classof(const Stmt *S) {
+    return S->getStmtClass() >= firstOpenACCConstructStmtConstant &&
+           S->getStmtClass() <= lastOpenACCConstructStmtConstant;
+  }
+
+  SourceLocation getBeginLoc() const { return Range.getBegin(); }
+  SourceLocation getEndLoc() const { return Range.getEnd(); }
+
+  child_range children() {
+    return child_range(child_iterator(), child_iterator());
+  }
+
+  const_child_range children() const {
+    return const_cast<OpenACCConstructStmt *>(this)->children();
+  }
+};
+
+/// This is a base class for any OpenACC statement-level constructs that have an
+/// associated statement. This class is not intended to be instantiated, but is
+/// a convenient place to hold the associated statement.
+class OpenACCAssociatedStmtConstruct : public OpenACCConstructStmt {
+  friend class ASTStmtWriter;
+  friend class ASTStmtReader;
+  template <typename Derived> friend class RecursiveASTVisitor;
+  Stmt *AssociatedStmt = nullptr;
+
+protected:
+  OpenACCAssociatedStmtConstruct(StmtClass SC, OpenACCDirectiveKind K,
+                                 SourceLocation Start, SourceLocation End)
+      : OpenACCConstructStmt(SC, K, Start, End) {}
+
+  void setAssociatedStmt(Stmt *S) { AssociatedStmt = S; }
+  Stmt *getAssociatedStmt() { return AssociatedStmt; }
+  const Stmt *getAssociatedStmt() const {
+    return const_cast<OpenACCAssociatedStmtConstruct *>(this)
+        ->getAssociatedStmt();
+  }
+
+public:
+  child_range children() {
+    if (getAssociatedStmt())
+      return child_range(&AssociatedStmt, &AssociatedStmt + 1);
+    return child_range(child_iterator(), child_iterator());
+  }
+
+  const_child_range children() const {
+    return const_cast<OpenACCAssociatedStmtConstruct *>(this)->children();
+  }
+};
+/// This class represents a compute construct, representing a 'Kind' of
+/// `parallel', 'serial', or 'kernel'. These constructs are associated with a
+/// 'structured block', defined as:
+///
+///  in C or C++, an executable statement, possibly compound, with a single
+///  entry at the top and a single exit at the bottom
+///
+/// At the moment there is no real motivation to have a different AST node for
+/// those three, as they are semantically identical, and have only minor
+/// differences in the permitted list of clauses, which can be differentiated by
+/// the 'Kind'.
+class OpenACCComputeConstruct : public OpenACCAssociatedStmtConstruct {
+  friend class ASTStmtWriter;
+  friend class ASTStmtReader;
+  friend class ASTContext;
+  OpenACCComputeConstruct()
+      : OpenACCAssociatedStmtConstruct(OpenACCComputeConstructClass,
+                                       OpenACCDirectiveKind::Invalid,
+                                       SourceLocation{}, SourceLocation{}) {}
+
+  OpenACCComputeConstruct(OpenACCDirectiveKind K, SourceLocation Start,
+                          SourceLocation End)
+      : OpenACCAssociatedStmtConstruct(OpenACCComputeConstructClass, K, Start,
+                                       End) {
+    assert((K == OpenACCDirectiveKind::Parallel ||
+            K == OpenACCDirectiveKind::Serial ||
+            K == OpenACCDirectiveKind::Kernels) &&
+           "Only parallel, serial, and kernels constructs should be "
+           "represented by this type");
+  }
+
+  void setStructuredBlock(Stmt *S) { setAssociatedStmt(S); }
+
+public:
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == OpenACCComputeConstructClass;
+  }
+
+  static OpenACCComputeConstruct *CreateEmpty(const ASTContext &C, EmptyShell);
+  static OpenACCComputeConstruct *Create(const ASTContext &C,
+                                         OpenACCDirectiveKind K,
+                                         SourceLocation BeginLoc,
+                                         SourceLocation EndLoc);
+
+  Stmt *getStructuredBlock() { return getAssociatedStmt(); }
+  const Stmt *getStructuredBlock() const {
+    return const_cast<OpenACCComputeConstruct *>(this)->getStructuredBlock();
+  }
+};
+} // namespace clang
+#endif // LLVM_CLANG_AST_STMTOPENACC_H

clang/include/clang/AST/StmtVisitor.h

@@ -13,13 +13,14 @@
 #ifndef LLVM_CLANG_AST_STMTVISITOR_H
 #define LLVM_CLANG_AST_STMTVISITOR_H
 
-#include "clang/AST/ExprConcepts.h"
 #include "clang/AST/ExprCXX.h"
+#include "clang/AST/ExprConcepts.h"
 #include "clang/AST/ExprObjC.h"
 #include "clang/AST/ExprOpenMP.h"
 #include "clang/AST/Stmt.h"
 #include "clang/AST/StmtCXX.h"
 #include "clang/AST/StmtObjC.h"
+#include "clang/AST/StmtOpenACC.h"
 #include "clang/AST/StmtOpenMP.h"
 #include "clang/Basic/LLVM.h"
 #include "llvm/ADT/STLExtras.h"

clang/include/clang/AST/TextNodeDumper.h

@@ -401,6 +401,7 @@ class TextNodeDumper
   void
   VisitLifetimeExtendedTemporaryDecl(const LifetimeExtendedTemporaryDecl *D);
   void VisitHLSLBufferDecl(const HLSLBufferDecl *D);
+  void VisitOpenACCConstructStmt(const OpenACCConstructStmt *S);
 };
 
 } // namespace clang

clang/include/clang/Basic/OpenACCKinds.h

@@ -16,6 +16,7 @@
 
 #include "clang/Basic/Diagnostic.h"
 #include "llvm/Support/ErrorHandling.h"
+#include "llvm/Support/raw_ostream.h"
 
 namespace clang {
 // Represents the Construct/Directive kind of a pragma directive. Note the
@@ -65,8 +66,9 @@ enum class OpenACCDirectiveKind {
   Invalid,
 };
 
-inline const StreamingDiagnostic &operator<<(const StreamingDiagnostic &Out,
-                                             OpenACCDirectiveKind K) {
+template <typename StreamTy>
+inline StreamTy &PrintOpenACCDirectiveKind(StreamTy &Out,
+                                           OpenACCDirectiveKind K) {
   switch (K) {
   case OpenACCDirectiveKind::Parallel:
     return Out << "parallel";
@@ -134,6 +136,16 @@ inline const StreamingDiagnostic &operator<<(const StreamingDiagnostic &Out,
   llvm_unreachable("Uncovered directive kind");
 }
 
+inline const StreamingDiagnostic &operator<<(const StreamingDiagnostic &Out,
+                                             OpenACCDirectiveKind K) {
+  return PrintOpenACCDirectiveKind(Out, K);
+}
+
+inline llvm::raw_ostream &operator<<(llvm::raw_ostream &Out,
+                                     OpenACCDirectiveKind K) {
+  return PrintOpenACCDirectiveKind(Out, K);
+}
+
 enum class OpenACCAtomicKind {
   Read,
   Write,
@@ -253,8 +265,8 @@ enum class OpenACCClauseKind {
   Invalid,
 };
 
-inline const StreamingDiagnostic &operator<<(const StreamingDiagnostic &Out,
-                                             OpenACCClauseKind K) {
+template <typename StreamTy>
+inline StreamTy &PrintOpenACCClauseKind(StreamTy &Out, OpenACCClauseKind K) {
   switch (K) {
   case OpenACCClauseKind::Finalize:
     return Out << "finalize";
@@ -387,6 +399,17 @@ inline const StreamingDiagnostic &operator<<(const StreamingDiagnostic &Out,
   }
   llvm_unreachable("Uncovered clause kind");
 }
+
+inline const StreamingDiagnostic &operator<<(const StreamingDiagnostic &Out,
+                                             OpenACCClauseKind K) {
+  return PrintOpenACCClauseKind(Out, K);
+}
+
+inline llvm::raw_ostream &operator<<(llvm::raw_ostream &Out,
+                                     OpenACCClauseKind K) {
+  return PrintOpenACCClauseKind(Out, K);
+}
+
 enum class OpenACCDefaultClauseKind {
   /// 'none' option.
   None,

clang/include/clang/Basic/StmtNodes.td

@@ -296,3 +296,9 @@ def OMPTargetTeamsGenericLoopDirective : StmtNode<OMPLoopDirective>;
 def OMPParallelGenericLoopDirective : StmtNode<OMPLoopDirective>;
 def OMPTargetParallelGenericLoopDirective : StmtNode<OMPLoopDirective>;
 def OMPErrorDirective : StmtNode<OMPExecutableDirective>;
+
+// OpenACC Constructs.
+def OpenACCConstructStmt : StmtNode<Stmt, /*abstract=*/1>;
+def OpenACCAssociatedStmtConstruct
+    : StmtNode<OpenACCConstructStmt, /*abstract=*/1>;
+def OpenACCComputeConstruct : StmtNode<OpenACCAssociatedStmtConstruct>;

clang/include/clang/Serialization/ASTBitCodes.h

@@ -2018,6 +2018,9 @@ enum StmtCode {
 
   // SYCLUniqueStableNameExpr
   EXPR_SYCL_UNIQUE_STABLE_NAME,
+
+  // OpenACC Constructs
+  STMT_OPENACC_COMPUTE_CONSTRUCT,
 };
 
 /// The kinds of designators that can occur in a

clang/lib/AST/ASTStructuralEquivalence.cpp

@@ -74,6 +74,7 @@
 #include "clang/AST/ExprOpenMP.h"
 #include "clang/AST/NestedNameSpecifier.h"
 #include "clang/AST/StmtObjC.h"
+#include "clang/AST/StmtOpenACC.h"
 #include "clang/AST/StmtOpenMP.h"
 #include "clang/AST/TemplateBase.h"
 #include "clang/AST/TemplateName.h"

clang/lib/AST/CMakeLists.txt

@@ -112,6 +112,7 @@ add_clang_library(clangAST
   StmtCXX.cpp
   StmtIterator.cpp
   StmtObjC.cpp
+  StmtOpenACC.cpp
   StmtOpenMP.cpp
   StmtPrinter.cpp
   StmtProfile.cpp

clang/lib/AST/Stmt.cpp

@@ -23,6 +23,7 @@
 #include "clang/AST/ExprOpenMP.h"
 #include "clang/AST/StmtCXX.h"
 #include "clang/AST/StmtObjC.h"
+#include "clang/AST/StmtOpenACC.h"
 #include "clang/AST/StmtOpenMP.h"
 #include "clang/AST/Type.h"
 #include "clang/Basic/CharInfo.h"

clang/lib/AST/StmtOpenACC.cpp

@@ -0,0 +1,33 @@
+//===--- StmtOpenACC.cpp - Classes for OpenACC Constructs -----------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file implements the subclesses of Stmt class declared in StmtOpenACC.h
+//
+//===----------------------------------------------------------------------===//
+
+#include "clang/AST/StmtOpenACC.h"
+#include "clang/AST/ASTContext.h"
+using namespace clang;
+
+OpenACCComputeConstruct *
+OpenACCComputeConstruct::CreateEmpty(const ASTContext &C, EmptyShell) {
+  void *Mem = C.Allocate(sizeof(OpenACCComputeConstruct),
+                         alignof(OpenACCComputeConstruct));
+  auto *Inst = new (Mem) OpenACCComputeConstruct;
+  return Inst;
+}
+
+OpenACCComputeConstruct *
+OpenACCComputeConstruct::Create(const ASTContext &C, OpenACCDirectiveKind K,
+                                SourceLocation BeginLoc,
+                                SourceLocation EndLoc) {
+  void *Mem = C.Allocate(sizeof(OpenACCComputeConstruct),
+                         alignof(OpenACCComputeConstruct));
+  auto *Inst = new (Mem) OpenACCComputeConstruct(K, BeginLoc, EndLoc);
+  return Inst;
+}

clang/lib/AST/StmtPrinter.cpp

@@ -1137,6 +1137,15 @@ void StmtPrinter::VisitOMPTargetParallelGenericLoopDirective(
   PrintOMPExecutableDirective(Node);
 }
 
+//===----------------------------------------------------------------------===//
+//  OpenACC construct printing methods
+//===----------------------------------------------------------------------===//
+void StmtPrinter::VisitOpenACCComputeConstruct(OpenACCComputeConstruct *S) {
+  Indent() << "#pragma acc " << S->getDirectiveKind();
+  // TODO OpenACC: Print Clauses.
+  PrintStmt(S->getStructuredBlock());
+}
+
 //===----------------------------------------------------------------------===//
 //  Expr printing methods.
 //===----------------------------------------------------------------------===//

clang/lib/AST/StmtProfile.cpp

@@ -2441,6 +2441,13 @@ void StmtProfiler::VisitTemplateArgument(const TemplateArgument &Arg) {
   }
 }
 
+void StmtProfiler::VisitOpenACCComputeConstruct(
+    const OpenACCComputeConstruct *S) {
+  // VisitStmt handles children, so the AssociatedStmt is handled.
+  VisitStmt(S);
+  // TODO OpenACC: Visit Clauses.
+}
+
 void Stmt::Profile(llvm::FoldingSetNodeID &ID, const ASTContext &Context,
                    bool Canonical, bool ProfileLambdaExpr) const {
   StmtProfilerWithPointers Profiler(ID, Context, Canonical, ProfileLambdaExpr);