[OpenACC][NFC] Add OpenACC Clause AST Nodes/infrastructure (#87675)

As a first step in adding clause support for OpenACC to Semantic
Analysis, this patch adds the 'base' AST nodes required for clauses.

This patch has no functional effect at the moment, but followup patches
will add the semantic analysis of clauses (plus individual clauses).

clang/include/clang/AST/ASTNodeTraverser.h

@@ -53,6 +53,7 @@ struct {
   void Visit(TypeLoc);
   void Visit(const Decl *D);
   void Visit(const CXXCtorInitializer *Init);
+  void Visit(const OpenACCClause *C);
   void Visit(const OMPClause *C);
   void Visit(const BlockDecl::Capture &C);
   void Visit(const GenericSelectionExpr::ConstAssociation &A);
@@ -239,6 +240,13 @@ public:
     });
   }
 
+  void Visit(const OpenACCClause *C) {
+    getNodeDelegate().AddChild([=] {
+      getNodeDelegate().Visit(C);
+      // TODO OpenACC: Switch on clauses that have children, and add them.
+    });
+  }
+
   void Visit(const OMPClause *C) {
     getNodeDelegate().AddChild([=] {
       getNodeDelegate().Visit(C);
@@ -799,6 +807,11 @@ public:
       Visit(C);
   }
 
+  void VisitOpenACCConstructStmt(const OpenACCConstructStmt *Node) {
+    for (const auto *C : Node->clauses())
+      Visit(C);
+  }
+
   void VisitInitListExpr(const InitListExpr *ILE) {
     if (auto *Filler = ILE->getArrayFiller()) {
       Visit(Filler, "array_filler");

clang/include/clang/AST/JSONNodeDumper.h

@@ -203,6 +203,7 @@ public:
   void Visit(const TemplateArgument &TA, SourceRange R = {},
              const Decl *From = nullptr, StringRef Label = {});
   void Visit(const CXXCtorInitializer *Init);
+  void Visit(const OpenACCClause *C);
   void Visit(const OMPClause *C);
   void Visit(const BlockDecl::Capture &C);
   void Visit(const GenericSelectionExpr::ConstAssociation &A);

clang/include/clang/AST/OpenACCClause.h

@@ -0,0 +1,135 @@
+//===- OpenACCClause.h - Classes for OpenACC clauses ------------*- 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 clauses.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_CLANG_AST_OPENACCCLAUSE_H
+#define LLVM_CLANG_AST_OPENACCCLAUSE_H
+#include "clang/AST/ASTContext.h"
+#include "clang/Basic/OpenACCKinds.h"
+
+namespace clang {
+/// This is the base type for all OpenACC Clauses.
+class OpenACCClause {
+  OpenACCClauseKind Kind;
+  SourceRange Location;
+
+protected:
+  OpenACCClause(OpenACCClauseKind K, SourceLocation BeginLoc,
+                SourceLocation EndLoc)
+      : Kind(K), Location(BeginLoc, EndLoc) {}
+
+public:
+  OpenACCClauseKind getClauseKind() const { return Kind; }
+  SourceLocation getBeginLoc() const { return Location.getBegin(); }
+  SourceLocation getEndLoc() const { return Location.getEnd(); }
+
+  static bool classof(const OpenACCClause *) { return true; }
+
+  virtual ~OpenACCClause() = default;
+};
+
+/// Represents a clause that has a list of parameters.
+class OpenACCClauseWithParams : public OpenACCClause {
+  /// Location of the '('.
+  SourceLocation LParenLoc;
+
+protected:
+  OpenACCClauseWithParams(OpenACCClauseKind K, SourceLocation BeginLoc,
+                          SourceLocation LParenLoc, SourceLocation EndLoc)
+      : OpenACCClause(K, BeginLoc, EndLoc), LParenLoc(LParenLoc) {}
+
+public:
+  SourceLocation getLParenLoc() const { return LParenLoc; }
+};
+
+template <class Impl> class OpenACCClauseVisitor {
+  Impl &getDerived() { return static_cast<Impl &>(*this); }
+
+public:
+  void VisitClauseList(ArrayRef<const OpenACCClause *> List) {
+    for (const OpenACCClause *Clause : List)
+      Visit(Clause);
+  }
+
+  void Visit(const OpenACCClause *C) {
+    if (!C)
+      return;
+
+    switch (C->getClauseKind()) {
+    case OpenACCClauseKind::Default:
+    case OpenACCClauseKind::Finalize:
+    case OpenACCClauseKind::IfPresent:
+    case OpenACCClauseKind::Seq:
+    case OpenACCClauseKind::Independent:
+    case OpenACCClauseKind::Auto:
+    case OpenACCClauseKind::Worker:
+    case OpenACCClauseKind::Vector:
+    case OpenACCClauseKind::NoHost:
+    case OpenACCClauseKind::If:
+    case OpenACCClauseKind::Self:
+    case OpenACCClauseKind::Copy:
+    case OpenACCClauseKind::UseDevice:
+    case OpenACCClauseKind::Attach:
+    case OpenACCClauseKind::Delete:
+    case OpenACCClauseKind::Detach:
+    case OpenACCClauseKind::Device:
+    case OpenACCClauseKind::DevicePtr:
+    case OpenACCClauseKind::DeviceResident:
+    case OpenACCClauseKind::FirstPrivate:
+    case OpenACCClauseKind::Host:
+    case OpenACCClauseKind::Link:
+    case OpenACCClauseKind::NoCreate:
+    case OpenACCClauseKind::Present:
+    case OpenACCClauseKind::Private:
+    case OpenACCClauseKind::CopyOut:
+    case OpenACCClauseKind::CopyIn:
+    case OpenACCClauseKind::Create:
+    case OpenACCClauseKind::Reduction:
+    case OpenACCClauseKind::Collapse:
+    case OpenACCClauseKind::Bind:
+    case OpenACCClauseKind::VectorLength:
+    case OpenACCClauseKind::NumGangs:
+    case OpenACCClauseKind::NumWorkers:
+    case OpenACCClauseKind::DeviceNum:
+    case OpenACCClauseKind::DefaultAsync:
+    case OpenACCClauseKind::DeviceType:
+    case OpenACCClauseKind::DType:
+    case OpenACCClauseKind::Async:
+    case OpenACCClauseKind::Tile:
+    case OpenACCClauseKind::Gang:
+    case OpenACCClauseKind::Wait:
+    case OpenACCClauseKind::Invalid:
+      llvm_unreachable("Clause visitor not yet implemented");
+    }
+    llvm_unreachable("Invalid Clause kind");
+  }
+};
+
+class OpenACCClausePrinter final
+    : public OpenACCClauseVisitor<OpenACCClausePrinter> {
+  raw_ostream &OS;
+
+public:
+  void VisitClauseList(ArrayRef<const OpenACCClause *> List) {
+    for (const OpenACCClause *Clause : List) {
+      Visit(Clause);
+
+      if (Clause != List.back())
+        OS << ' ';
+    }
+  }
+  OpenACCClausePrinter(raw_ostream &OS) : OS(OS) {}
+};
+
+} // namespace clang
+
+#endif // LLVM_CLANG_AST_OPENACCCLAUSE_H

clang/include/clang/AST/RecursiveASTVisitor.h

@@ -509,6 +509,7 @@ private:
   bool TraverseOpenACCConstructStmt(OpenACCConstructStmt *S);
   bool
   TraverseOpenACCAssociatedStmtConstruct(OpenACCAssociatedStmtConstruct *S);
+  bool VisitOpenACCClauseList(ArrayRef<const OpenACCClause *>);
 };
 
 template <typename Derived>
@@ -3936,8 +3937,8 @@ bool RecursiveASTVisitor<Derived>::VisitOMPXBareClause(OMPXBareClause *C) {
 
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::TraverseOpenACCConstructStmt(
-    OpenACCConstructStmt *) {
-  // TODO OpenACC: When we implement clauses, ensure we traverse them here.
+    OpenACCConstructStmt *C) {
+  TRY_TO(VisitOpenACCClauseList(C->clauses()));
   return true;
 }
 
@@ -3949,6 +3950,14 @@ bool RecursiveASTVisitor<Derived>::TraverseOpenACCAssociatedStmtConstruct(
   return true;
 }
 
+template <typename Derived>
+bool RecursiveASTVisitor<Derived>::VisitOpenACCClauseList(
+    ArrayRef<const OpenACCClause *>) {
+  // TODO OpenACC: When we have Clauses with expressions, we should visit them
+  // here.
+  return true;
+}
+
 DEF_TRAVERSE_STMT(OpenACCComputeConstruct,
                   { TRY_TO(TraverseOpenACCAssociatedStmtConstruct(S)); })
 

clang/include/clang/AST/StmtOpenACC.h

@@ -13,9 +13,11 @@
 #ifndef LLVM_CLANG_AST_STMTOPENACC_H
 #define LLVM_CLANG_AST_STMTOPENACC_H
 
+#include "clang/AST/OpenACCClause.h"
 #include "clang/AST/Stmt.h"
 #include "clang/Basic/OpenACCKinds.h"
 #include "clang/Basic/SourceLocation.h"
+#include <memory>
 
 namespace clang {
 /// This is the base class for an OpenACC statement-level construct, other
@@ -30,13 +32,23 @@ class OpenACCConstructStmt : public Stmt {
   /// the directive.
   SourceRange Range;
 
-  // TODO OPENACC: Clauses should probably be collected in this class.
+  /// The list of clauses.  This is stored here as an ArrayRef, as this is the
+  /// most convienient place to access the list, however the list itself should
+  /// be stored in leaf nodes, likely in trailing-storage.
+  MutableArrayRef<const OpenACCClause *> Clauses;
 
 protected:
   OpenACCConstructStmt(StmtClass SC, OpenACCDirectiveKind K,
                        SourceLocation Start, SourceLocation End)
       : Stmt(SC), Kind(K), Range(Start, End) {}
 
+  // Used only for initialization, the leaf class can initialize this to
+  // trailing storage.
+  void setClauseList(MutableArrayRef<const OpenACCClause *> NewClauses) {
+    assert(Clauses.empty() && "Cannot change clause list");
+    Clauses = NewClauses;
+  }
+
 public:
   OpenACCDirectiveKind getDirectiveKind() const { return Kind; }
 
@@ -47,6 +59,7 @@ public:
 
   SourceLocation getBeginLoc() const { return Range.getBegin(); }
   SourceLocation getEndLoc() const { return Range.getEnd(); }
+  ArrayRef<const OpenACCClause *> clauses() const { return Clauses; }
 
   child_range children() {
     return child_range(child_iterator(), child_iterator());
@@ -101,17 +114,32 @@ public:
 /// 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 {
+class OpenACCComputeConstruct final
+    : public OpenACCAssociatedStmtConstruct,
+      public llvm::TrailingObjects<OpenACCComputeConstruct,
+                                   const OpenACCClause *> {
   friend class ASTStmtWriter;
   friend class ASTStmtReader;
   friend class ASTContext;
-  OpenACCComputeConstruct()
-      : OpenACCAssociatedStmtConstruct(
-            OpenACCComputeConstructClass, OpenACCDirectiveKind::Invalid,
-            SourceLocation{}, SourceLocation{}, /*AssociatedStmt=*/nullptr) {}
+  OpenACCComputeConstruct(unsigned NumClauses)
+      : OpenACCAssociatedStmtConstruct(OpenACCComputeConstructClass,
+                                       OpenACCDirectiveKind::Invalid,
+                                       SourceLocation{}, SourceLocation{},
+                                       /*AssociatedStmt=*/nullptr) {
+    // We cannot send the TrailingObjects storage to the base class (which holds
+    // a reference to the data) until it is constructed, so we have to set it
+    // separately here.
+    std::uninitialized_value_construct(
+        getTrailingObjects<const OpenACCClause *>(),
+        getTrailingObjects<const OpenACCClause *>() + NumClauses);
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  NumClauses));
+  }
 
   OpenACCComputeConstruct(OpenACCDirectiveKind K, SourceLocation Start,
-                          SourceLocation End, Stmt *StructuredBlock)
+                          SourceLocation End,
+                          ArrayRef<const OpenACCClause *> Clauses,
+                          Stmt *StructuredBlock)
       : OpenACCAssociatedStmtConstruct(OpenACCComputeConstructClass, K, Start,
                                        End, StructuredBlock) {
     assert((K == OpenACCDirectiveKind::Parallel ||
@@ -119,6 +147,13 @@ class OpenACCComputeConstruct : public OpenACCAssociatedStmtConstruct {
             K == OpenACCDirectiveKind::Kernels) &&
            "Only parallel, serial, and kernels constructs should be "
            "represented by this type");
+
+    // Initialize the trailing storage.
+    std::uninitialized_copy(Clauses.begin(), Clauses.end(),
+                            getTrailingObjects<const OpenACCClause *>());
+
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  Clauses.size()));
   }
 
   void setStructuredBlock(Stmt *S) { setAssociatedStmt(S); }
@@ -128,10 +163,12 @@ public:
     return T->getStmtClass() == OpenACCComputeConstructClass;
   }
 
-  static OpenACCComputeConstruct *CreateEmpty(const ASTContext &C, EmptyShell);
+  static OpenACCComputeConstruct *CreateEmpty(const ASTContext &C,
+                                              unsigned NumClauses);
   static OpenACCComputeConstruct *
   Create(const ASTContext &C, OpenACCDirectiveKind K, SourceLocation BeginLoc,
-         SourceLocation EndLoc, Stmt *StructuredBlock);
+         SourceLocation EndLoc, ArrayRef<const OpenACCClause *> Clauses,
+         Stmt *StructuredBlock);
 
   Stmt *getStructuredBlock() { return getAssociatedStmt(); }
   const Stmt *getStructuredBlock() const {

clang/include/clang/AST/TextNodeDumper.h

@@ -189,6 +189,8 @@ public:
 
   void Visit(const OMPClause *C);
 
+  void Visit(const OpenACCClause *C);
+
   void Visit(const BlockDecl::Capture &C);
 
   void Visit(const GenericSelectionExpr::ConstAssociation &A);

clang/include/clang/Serialization/ASTRecordReader.h

@@ -24,6 +24,7 @@
 #include "llvm/ADT/APSInt.h"
 
 namespace clang {
+class OpenACCClause;
 class OMPTraitInfo;
 class OMPChildren;
 
@@ -278,6 +279,12 @@ public:
   /// Read an OpenMP children, advancing Idx.
   void readOMPChildren(OMPChildren *Data);
 
+  /// Read an OpenACC clause, advancing Idx.
+  OpenACCClause *readOpenACCClause();
+
+  /// Read a list of OpenACC clauses into the passed SmallVector.
+  void readOpenACCClauseList(MutableArrayRef<const OpenACCClause *> Clauses);
+
   /// Read a source location, advancing Idx.
   SourceLocation readSourceLocation(LocSeq *Seq = nullptr) {
     return Reader->ReadSourceLocation(*F, Record, Idx, Seq);

clang/include/clang/Serialization/ASTRecordWriter.h

@@ -21,6 +21,7 @@
 
 namespace clang {
 
+class OpenACCClause;
 class TypeLoc;
 
 /// An object for streaming information to a record.
@@ -292,6 +293,12 @@ public:
   /// Writes data related to the OpenMP directives.
   void writeOMPChildren(OMPChildren *Data);
 
+  /// Writes out a single OpenACC Clause.
+  void writeOpenACCClause(const OpenACCClause *C);
+
+  /// Writes out a list of OpenACC clauses.
+  void writeOpenACCClauseList(ArrayRef<const OpenACCClause *> Clauses);
+
   /// Emit a string.
   void AddString(StringRef Str) {
     return Writer->AddString(Str, *Record);

clang/lib/AST/CMakeLists.txt

@@ -98,6 +98,7 @@ add_clang_library(clangAST
   NSAPI.cpp
   ODRDiagsEmitter.cpp
   ODRHash.cpp
+  OpenACCClause.cpp
   OpenMPClause.cpp
   OSLog.cpp
   ParentMap.cpp

clang/lib/AST/JSONNodeDumper.cpp

@@ -187,6 +187,8 @@ void JSONNodeDumper::Visit(const CXXCtorInitializer *Init) {
     llvm_unreachable("Unknown initializer type");
 }
 
+void JSONNodeDumper::Visit(const OpenACCClause *C) {}
+
 void JSONNodeDumper::Visit(const OMPClause *C) {}
 
 void JSONNodeDumper::Visit(const BlockDecl::Capture &C) {

clang/lib/AST/OpenACCClause.cpp

@@ -0,0 +1,17 @@
+//===---- OpenACCClause.cpp - Classes for OpenACC Clauses  ----------------===//
+//
+// 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 subclasses of the OpenACCClause class declared in
+// OpenACCClause.h
+//
+//===----------------------------------------------------------------------===//
+
+#include "clang/AST/OpenACCClause.h"
+#include "clang/AST/ASTContext.h"
+
+using namespace clang;

clang/lib/AST/StmtOpenACC.cpp

@@ -15,20 +15,23 @@
 using namespace clang;
 
 OpenACCComputeConstruct *
-OpenACCComputeConstruct::CreateEmpty(const ASTContext &C, EmptyShell) {
-  void *Mem = C.Allocate(sizeof(OpenACCComputeConstruct),
-                         alignof(OpenACCComputeConstruct));
-  auto *Inst = new (Mem) OpenACCComputeConstruct;
+OpenACCComputeConstruct::CreateEmpty(const ASTContext &C, unsigned NumClauses) {
+  void *Mem = C.Allocate(
+      OpenACCComputeConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          NumClauses));
+  auto *Inst = new (Mem) OpenACCComputeConstruct(NumClauses);
   return Inst;
 }
 
 OpenACCComputeConstruct *
 OpenACCComputeConstruct::Create(const ASTContext &C, OpenACCDirectiveKind K,
                                 SourceLocation BeginLoc, SourceLocation EndLoc,
+                                ArrayRef<const OpenACCClause *> Clauses,
                                 Stmt *StructuredBlock) {
-  void *Mem = C.Allocate(sizeof(OpenACCComputeConstruct),
-                         alignof(OpenACCComputeConstruct));
-  auto *Inst =
-      new (Mem) OpenACCComputeConstruct(K, BeginLoc, EndLoc, StructuredBlock);
+  void *Mem = C.Allocate(
+      OpenACCComputeConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          Clauses.size()));
+  auto *Inst = new (Mem)
+      OpenACCComputeConstruct(K, BeginLoc, EndLoc, Clauses, StructuredBlock);
   return Inst;
 }

clang/lib/AST/StmtPrinter.cpp

@@ -1142,7 +1142,13 @@ void StmtPrinter::VisitOMPTargetParallelGenericLoopDirective(
 //===----------------------------------------------------------------------===//
 void StmtPrinter::VisitOpenACCComputeConstruct(OpenACCComputeConstruct *S) {
   Indent() << "#pragma acc " << S->getDirectiveKind();
-  // TODO OpenACC: Print Clauses.
+
+  if (!S->clauses().empty()) {
+    OS << ' ';
+    OpenACCClausePrinter Printer(OS);
+    Printer.VisitClauseList(S->clauses());
+  }
+
   PrintStmt(S->getStructuredBlock());
 }
 

clang/lib/AST/StmtProfile.cpp

@@ -2441,11 +2441,30 @@ void StmtProfiler::VisitTemplateArgument(const TemplateArgument &Arg) {
   }
 }
 
+namespace {
+class OpenACCClauseProfiler
+    : public OpenACCClauseVisitor<OpenACCClauseProfiler> {
+
+public:
+  OpenACCClauseProfiler() = default;
+
+  void VisitOpenACCClauseList(ArrayRef<const OpenACCClause *> Clauses) {
+    for (const OpenACCClause *Clause : Clauses) {
+      // TODO OpenACC: When we have clauses with expressions, we should
+      // profile them too.
+      Visit(Clause);
+    }
+  }
+};
+} // namespace
+
 void StmtProfiler::VisitOpenACCComputeConstruct(
     const OpenACCComputeConstruct *S) {
   // VisitStmt handles children, so the AssociatedStmt is handled.
   VisitStmt(S);
-  // TODO OpenACC: Visit Clauses.
+
+  OpenACCClauseProfiler P;
+  P.VisitOpenACCClauseList(S->clauses());
 }
 
 void Stmt::Profile(llvm::FoldingSetNodeID &ID, const ASTContext &Context,

clang/lib/AST/TextNodeDumper.cpp

@@ -381,6 +381,28 @@ void TextNodeDumper::Visit(const OMPClause *C) {
     OS << " <implicit>";
 }
 
+void TextNodeDumper::Visit(const OpenACCClause *C) {
+  if (!C) {
+    ColorScope Color(OS, ShowColors, NullColor);
+    OS << "<<<NULL>>> OpenACCClause";
+    return;
+  }
+  {
+    ColorScope Color(OS, ShowColors, AttrColor);
+    OS << C->getClauseKind();
+
+    // Handle clauses with parens for types that have no children, likely
+    // because there is no sub expression.
+    switch (C->getClauseKind()) {
+    default:
+      // Nothing to do here.
+      break;
+    }
+  }
+  dumpPointer(C);
+  dumpSourceRange(SourceRange(C->getBeginLoc(), C->getEndLoc()));
+}
+
 void TextNodeDumper::Visit(const GenericSelectionExpr::ConstAssociation &A) {
   const TypeSourceInfo *TSI = A.getTypeSourceInfo();
   if (TSI) {
@@ -2684,5 +2706,4 @@ void TextNodeDumper::VisitHLSLBufferDecl(const HLSLBufferDecl *D) {
 
 void TextNodeDumper::VisitOpenACCConstructStmt(const OpenACCConstructStmt *S) {
   OS << " " << S->getDirectiveKind();
-  // TODO OpenACC: Dump clauses as well.
 }

clang/lib/Sema/SemaOpenACC.cpp

@@ -94,8 +94,10 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective(OpenACCDirectiveKind K,
   case OpenACCDirectiveKind::Parallel:
   case OpenACCDirectiveKind::Serial:
   case OpenACCDirectiveKind::Kernels:
+    // TODO OpenACC: Add clauses to the construct here.
     return OpenACCComputeConstruct::Create(
         getASTContext(), K, StartLoc, EndLoc,
+        /*Clauses=*/std::nullopt,
         AssocStmt.isUsable() ? AssocStmt.get() : nullptr);
   }
   llvm_unreachable("Unhandled case in directive handling?");

clang/lib/Serialization/ASTReader.cpp

@@ -32,6 +32,7 @@
 #include "clang/AST/NestedNameSpecifier.h"
 #include "clang/AST/ODRDiagsEmitter.h"
 #include "clang/AST/ODRHash.h"
+#include "clang/AST/OpenACCClause.h"
 #include "clang/AST/OpenMPClause.h"
 #include "clang/AST/RawCommentList.h"
 #include "clang/AST/TemplateBase.h"
@@ -53,6 +54,7 @@
 #include "clang/Basic/LangOptions.h"
 #include "clang/Basic/Module.h"
 #include "clang/Basic/ObjCRuntime.h"
+#include "clang/Basic/OpenACCKinds.h"
 #include "clang/Basic/OpenMPKinds.h"
 #include "clang/Basic/OperatorKinds.h"
 #include "clang/Basic/PragmaKinds.h"
@@ -11751,3 +11753,66 @@ void ASTRecordReader::readOMPChildren(OMPChildren *Data) {
   for (unsigned I = 0, E = Data->getNumChildren(); I < E; ++I)
     Data->getChildren()[I] = readStmt();
 }
+
+OpenACCClause *ASTRecordReader::readOpenACCClause() {
+  OpenACCClauseKind ClauseKind = readEnum<OpenACCClauseKind>();
+  // TODO OpenACC: We don't have these used anywhere, but eventually we should
+  // be constructing the Clauses with them, so these attributes can go away at
+  // that point.
+  [[maybe_unused]] SourceLocation BeginLoc = readSourceLocation();
+  [[maybe_unused]] SourceLocation EndLoc = readSourceLocation();
+
+  switch (ClauseKind) {
+  case OpenACCClauseKind::Default:
+  case OpenACCClauseKind::Finalize:
+  case OpenACCClauseKind::IfPresent:
+  case OpenACCClauseKind::Seq:
+  case OpenACCClauseKind::Independent:
+  case OpenACCClauseKind::Auto:
+  case OpenACCClauseKind::Worker:
+  case OpenACCClauseKind::Vector:
+  case OpenACCClauseKind::NoHost:
+  case OpenACCClauseKind::If:
+  case OpenACCClauseKind::Self:
+  case OpenACCClauseKind::Copy:
+  case OpenACCClauseKind::UseDevice:
+  case OpenACCClauseKind::Attach:
+  case OpenACCClauseKind::Delete:
+  case OpenACCClauseKind::Detach:
+  case OpenACCClauseKind::Device:
+  case OpenACCClauseKind::DevicePtr:
+  case OpenACCClauseKind::DeviceResident:
+  case OpenACCClauseKind::FirstPrivate:
+  case OpenACCClauseKind::Host:
+  case OpenACCClauseKind::Link:
+  case OpenACCClauseKind::NoCreate:
+  case OpenACCClauseKind::Present:
+  case OpenACCClauseKind::Private:
+  case OpenACCClauseKind::CopyOut:
+  case OpenACCClauseKind::CopyIn:
+  case OpenACCClauseKind::Create:
+  case OpenACCClauseKind::Reduction:
+  case OpenACCClauseKind::Collapse:
+  case OpenACCClauseKind::Bind:
+  case OpenACCClauseKind::VectorLength:
+  case OpenACCClauseKind::NumGangs:
+  case OpenACCClauseKind::NumWorkers:
+  case OpenACCClauseKind::DeviceNum:
+  case OpenACCClauseKind::DefaultAsync:
+  case OpenACCClauseKind::DeviceType:
+  case OpenACCClauseKind::DType:
+  case OpenACCClauseKind::Async:
+  case OpenACCClauseKind::Tile:
+  case OpenACCClauseKind::Gang:
+  case OpenACCClauseKind::Wait:
+  case OpenACCClauseKind::Invalid:
+    llvm_unreachable("Clause serialization not yet implemented");
+  }
+  llvm_unreachable("Invalid Clause Kind");
+}
+
+void ASTRecordReader::readOpenACCClauseList(
+    MutableArrayRef<const OpenACCClause *> Clauses) {
+  for (unsigned I = 0; I < Clauses.size(); ++I)
+    Clauses[I] = readOpenACCClause();
+}

clang/lib/Serialization/ASTReaderStmt.cpp

@@ -2784,9 +2784,10 @@ void ASTStmtReader::VisitOMPTargetParallelGenericLoopDirective(
 // OpenACC Constructs/Directives.
 //===----------------------------------------------------------------------===//
 void ASTStmtReader::VisitOpenACCConstructStmt(OpenACCConstructStmt *S) {
+  (void)Record.readInt();
   S->Kind = Record.readEnum<OpenACCDirectiveKind>();
   S->Range = Record.readSourceRange();
-  // TODO OpenACC: Deserialize Clauses.
+  Record.readOpenACCClauseList(S->Clauses);
 }
 
 void ASTStmtReader::VisitOpenACCAssociatedStmtConstruct(
@@ -4218,10 +4219,11 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
       S = new (Context) ConceptSpecializationExpr(Empty);
       break;
     }
-    case STMT_OPENACC_COMPUTE_CONSTRUCT:
-      S = OpenACCComputeConstruct::CreateEmpty(Context, Empty);
+    case STMT_OPENACC_COMPUTE_CONSTRUCT: {
+      unsigned NumClauses = Record[ASTStmtReader::NumStmtFields];
+      S = OpenACCComputeConstruct::CreateEmpty(Context, NumClauses);
       break;
-
+    }
     case EXPR_REQUIRES:
       unsigned numLocalParameters = Record[ASTStmtReader::NumExprFields];
       unsigned numRequirement = Record[ASTStmtReader::NumExprFields + 1];

clang/lib/Serialization/ASTWriter.cpp

@@ -29,6 +29,7 @@
 #include "clang/AST/ExprCXX.h"
 #include "clang/AST/LambdaCapture.h"
 #include "clang/AST/NestedNameSpecifier.h"
+#include "clang/AST/OpenACCClause.h"
 #include "clang/AST/OpenMPClause.h"
 #include "clang/AST/RawCommentList.h"
 #include "clang/AST/TemplateName.h"
@@ -44,6 +45,7 @@
 #include "clang/Basic/LangOptions.h"
 #include "clang/Basic/Module.h"
 #include "clang/Basic/ObjCRuntime.h"
+#include "clang/Basic/OpenACCKinds.h"
 #include "clang/Basic/OpenCLOptions.h"
 #include "clang/Basic/SourceLocation.h"
 #include "clang/Basic/SourceManager.h"
@@ -7397,3 +7399,63 @@ void ASTRecordWriter::writeOMPChildren(OMPChildren *Data) {
   for (unsigned I = 0, E = Data->getNumChildren(); I < E; ++I)
     AddStmt(Data->getChildren()[I]);
 }
+
+void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
+  writeEnum(C->getClauseKind());
+  writeSourceLocation(C->getBeginLoc());
+  writeSourceLocation(C->getEndLoc());
+
+  switch (C->getClauseKind()) {
+  case OpenACCClauseKind::Default:
+  case OpenACCClauseKind::Finalize:
+  case OpenACCClauseKind::IfPresent:
+  case OpenACCClauseKind::Seq:
+  case OpenACCClauseKind::Independent:
+  case OpenACCClauseKind::Auto:
+  case OpenACCClauseKind::Worker:
+  case OpenACCClauseKind::Vector:
+  case OpenACCClauseKind::NoHost:
+  case OpenACCClauseKind::If:
+  case OpenACCClauseKind::Self:
+  case OpenACCClauseKind::Copy:
+  case OpenACCClauseKind::UseDevice:
+  case OpenACCClauseKind::Attach:
+  case OpenACCClauseKind::Delete:
+  case OpenACCClauseKind::Detach:
+  case OpenACCClauseKind::Device:
+  case OpenACCClauseKind::DevicePtr:
+  case OpenACCClauseKind::DeviceResident:
+  case OpenACCClauseKind::FirstPrivate:
+  case OpenACCClauseKind::Host:
+  case OpenACCClauseKind::Link:
+  case OpenACCClauseKind::NoCreate:
+  case OpenACCClauseKind::Present:
+  case OpenACCClauseKind::Private:
+  case OpenACCClauseKind::CopyOut:
+  case OpenACCClauseKind::CopyIn:
+  case OpenACCClauseKind::Create:
+  case OpenACCClauseKind::Reduction:
+  case OpenACCClauseKind::Collapse:
+  case OpenACCClauseKind::Bind:
+  case OpenACCClauseKind::VectorLength:
+  case OpenACCClauseKind::NumGangs:
+  case OpenACCClauseKind::NumWorkers:
+  case OpenACCClauseKind::DeviceNum:
+  case OpenACCClauseKind::DefaultAsync:
+  case OpenACCClauseKind::DeviceType:
+  case OpenACCClauseKind::DType:
+  case OpenACCClauseKind::Async:
+  case OpenACCClauseKind::Tile:
+  case OpenACCClauseKind::Gang:
+  case OpenACCClauseKind::Wait:
+  case OpenACCClauseKind::Invalid:
+    llvm_unreachable("Clause serialization not yet implemented");
+  }
+  llvm_unreachable("Invalid Clause Kind");
+}
+
+void ASTRecordWriter::writeOpenACCClauseList(
+    ArrayRef<const OpenACCClause *> Clauses) {
+  for (const OpenACCClause *Clause : Clauses)
+    writeOpenACCClause(Clause);
+}

clang/lib/Serialization/ASTWriterStmt.cpp

@@ -2839,9 +2839,10 @@ void ASTStmtWriter::VisitOMPTargetParallelGenericLoopDirective(
 // OpenACC Constructs/Directives.
 //===----------------------------------------------------------------------===//
 void ASTStmtWriter::VisitOpenACCConstructStmt(OpenACCConstructStmt *S) {
+  Record.push_back(S->clauses().size());
   Record.writeEnum(S->Kind);
   Record.AddSourceRange(S->Range);
-  // TODO OpenACC: Serialize Clauses.
+  Record.writeOpenACCClauseList(S->clauses());
 }
 
 void ASTStmtWriter::VisitOpenACCAssociatedStmtConstruct(