[clang] 010d011 - [OpenACC] Create AST nodes for 'data' constructs

via cfe-commits cfe-commits at lists.llvm.org
Thu Dec 12 07:28:35 PST 2024


Author: erichkeane
Date: 2024-12-12T07:28:30-08:00
New Revision: 010d0115fc8e3834fc6f747f0841f3b1e467c4da

URL: https://github.com/llvm/llvm-project/commit/010d0115fc8e3834fc6f747f0841f3b1e467c4da
DIFF: https://github.com/llvm/llvm-project/commit/010d0115fc8e3834fc6f747f0841f3b1e467c4da.diff

LOG: [OpenACC] Create AST nodes for 'data' constructs

These constructs are all very similar and closely related, so this patch
creates the AST nodes for them, serialization, printing/etc.
Additionally the restrictions are all added as tests/todos in the tests,
as those will have to be implemented once we get those clauses implemented.

Added: 
    clang/test/AST/ast-print-openacc-data-construct.cpp
    clang/test/SemaOpenACC/data-construct-ast.cpp
    clang/test/SemaOpenACC/data-construct.cpp

Modified: 
    clang/include/clang-c/Index.h
    clang/include/clang/AST/RecursiveASTVisitor.h
    clang/include/clang/AST/StmtOpenACC.h
    clang/include/clang/AST/TextNodeDumper.h
    clang/include/clang/Basic/StmtNodes.td
    clang/include/clang/Serialization/ASTBitCodes.h
    clang/lib/AST/StmtOpenACC.cpp
    clang/lib/AST/StmtPrinter.cpp
    clang/lib/AST/StmtProfile.cpp
    clang/lib/AST/TextNodeDumper.cpp
    clang/lib/CodeGen/CGStmt.cpp
    clang/lib/CodeGen/CodeGenFunction.h
    clang/lib/Parse/ParseOpenACC.cpp
    clang/lib/Sema/SemaExceptionSpec.cpp
    clang/lib/Sema/SemaOpenACC.cpp
    clang/lib/Sema/TreeTransform.h
    clang/lib/Serialization/ASTReaderStmt.cpp
    clang/lib/Serialization/ASTWriterStmt.cpp
    clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
    clang/test/ParserOpenACC/parse-clauses.c
    clang/test/ParserOpenACC/parse-clauses.cpp
    clang/test/ParserOpenACC/parse-constructs.c
    clang/test/SemaOpenACC/combined-construct-collapse-clause.cpp
    clang/test/SemaOpenACC/combined-construct-default-clause.c
    clang/test/SemaOpenACC/combined-construct-if-clause.c
    clang/test/SemaOpenACC/compute-construct-default-clause.c
    clang/test/SemaOpenACC/compute-construct-device_type-clause.c
    clang/test/SemaOpenACC/compute-construct-if-clause.c
    clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp
    clang/tools/libclang/CIndex.cpp
    clang/tools/libclang/CXCursor.cpp

Removed: 
    


################################################################################
diff  --git a/clang/include/clang-c/Index.h b/clang/include/clang-c/Index.h
index 8fc06328f0bcef..29858f00fad74e 100644
--- a/clang/include/clang-c/Index.h
+++ b/clang/include/clang-c/Index.h
@@ -2166,9 +2166,27 @@ enum CXCursorKind {
    */
   CXCursor_OpenACCLoopConstruct = 321,
 
+  /** OpenACC Combined Constructs.
+   */
   CXCursor_OpenACCCombinedConstruct = 322,
 
-  CXCursor_LastStmt = CXCursor_OpenACCCombinedConstruct,
+  /** OpenACC data Construct.
+   */
+  CXCursor_OpenACCDataConstruct = 323,
+
+  /** OpenACC enter data Construct.
+   */
+  CXCursor_OpenACCEnterDataConstruct = 324,
+
+  /** OpenACC exit data Construct.
+   */
+  CXCursor_OpenACCExitDataConstruct = 325,
+
+  /** OpenACC host_data Construct.
+   */
+  CXCursor_OpenACCHostDataConstruct = 326,
+
+  CXCursor_LastStmt = CXCursor_OpenACCHostDataConstruct,
 
   /**
    * Cursor that represents the translation unit itself.

diff  --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h
index 76b598a5db2382..33363072c716f2 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -4058,6 +4058,12 @@ DEF_TRAVERSE_STMT(OpenACCLoopConstruct,
                   { TRY_TO(TraverseOpenACCAssociatedStmtConstruct(S)); })
 DEF_TRAVERSE_STMT(OpenACCCombinedConstruct,
                   { TRY_TO(TraverseOpenACCAssociatedStmtConstruct(S)); })
+DEF_TRAVERSE_STMT(OpenACCDataConstruct,
+                  { TRY_TO(TraverseOpenACCAssociatedStmtConstruct(S)); })
+DEF_TRAVERSE_STMT(OpenACCEnterDataConstruct, {})
+DEF_TRAVERSE_STMT(OpenACCExitDataConstruct, {})
+DEF_TRAVERSE_STMT(OpenACCHostDataConstruct,
+                  { TRY_TO(TraverseOpenACCAssociatedStmtConstruct(S)); })
 
 // Traverse HLSL: Out argument expression
 DEF_TRAVERSE_STMT(HLSLOutArgExpr, {})

diff  --git a/clang/include/clang/AST/StmtOpenACC.h b/clang/include/clang/AST/StmtOpenACC.h
index fa8793e740822f..df73980822c7be 100644
--- a/clang/include/clang/AST/StmtOpenACC.h
+++ b/clang/include/clang/AST/StmtOpenACC.h
@@ -292,5 +292,175 @@ class OpenACCCombinedConstruct final
     return const_cast<OpenACCCombinedConstruct *>(this)->getLoop();
   }
 };
+
+// This class represents a 'data' construct, which has an associated statement
+// and clauses, but is otherwise pretty simple.
+class OpenACCDataConstruct final
+    : public OpenACCAssociatedStmtConstruct,
+      public llvm::TrailingObjects<OpenACCCombinedConstruct,
+                                   const OpenACCClause *> {
+  OpenACCDataConstruct(unsigned NumClauses)
+      : OpenACCAssociatedStmtConstruct(
+            OpenACCDataConstructClass, OpenACCDirectiveKind::Data,
+            SourceLocation{}, SourceLocation{}, SourceLocation{},
+            /*AssociatedStmt=*/nullptr) {
+    std::uninitialized_value_construct(
+        getTrailingObjects<const OpenACCClause *>(),
+        getTrailingObjects<const OpenACCClause *>() + NumClauses);
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  NumClauses));
+  }
+
+  OpenACCDataConstruct(SourceLocation Start, SourceLocation DirectiveLoc,
+                       SourceLocation End,
+                       ArrayRef<const OpenACCClause *> Clauses,
+                       Stmt *StructuredBlock)
+      : OpenACCAssociatedStmtConstruct(OpenACCDataConstructClass,
+                                       OpenACCDirectiveKind::Data, Start,
+                                       DirectiveLoc, End, StructuredBlock) {
+    std::uninitialized_copy(Clauses.begin(), Clauses.end(),
+                            getTrailingObjects<const OpenACCClause *>());
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  Clauses.size()));
+  }
+  void setStructuredBlock(Stmt *S) { setAssociatedStmt(S); }
+
+public:
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == OpenACCDataConstructClass;
+  }
+
+  static OpenACCDataConstruct *CreateEmpty(const ASTContext &C,
+                                           unsigned NumClauses);
+  static OpenACCDataConstruct *Create(const ASTContext &C, SourceLocation Start,
+                                      SourceLocation DirectiveLoc,
+                                      SourceLocation End,
+                                      ArrayRef<const OpenACCClause *> Clauses,
+                                      Stmt *StructuredBlock);
+  Stmt *getStructuredBlock() { return getAssociatedStmt(); }
+  const Stmt *getStructuredBlock() const {
+    return const_cast<OpenACCDataConstruct *>(this)->getStructuredBlock();
+  }
+};
+// This class represents a 'enter data' construct, which JUST has clauses.
+class OpenACCEnterDataConstruct final
+    : public OpenACCConstructStmt,
+      public llvm::TrailingObjects<OpenACCCombinedConstruct,
+                                   const OpenACCClause *> {
+  OpenACCEnterDataConstruct(unsigned NumClauses)
+      : OpenACCConstructStmt(OpenACCEnterDataConstructClass,
+                             OpenACCDirectiveKind::EnterData, SourceLocation{},
+                             SourceLocation{}, SourceLocation{}) {
+    std::uninitialized_value_construct(
+        getTrailingObjects<const OpenACCClause *>(),
+        getTrailingObjects<const OpenACCClause *>() + NumClauses);
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  NumClauses));
+  }
+  OpenACCEnterDataConstruct(SourceLocation Start, SourceLocation DirectiveLoc,
+                            SourceLocation End,
+                            ArrayRef<const OpenACCClause *> Clauses)
+      : OpenACCConstructStmt(OpenACCEnterDataConstructClass,
+                             OpenACCDirectiveKind::EnterData, Start,
+                             DirectiveLoc, End) {
+    std::uninitialized_copy(Clauses.begin(), Clauses.end(),
+                            getTrailingObjects<const OpenACCClause *>());
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  Clauses.size()));
+  }
+
+public:
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == OpenACCEnterDataConstructClass;
+  }
+  static OpenACCEnterDataConstruct *CreateEmpty(const ASTContext &C,
+                                                unsigned NumClauses);
+  static OpenACCEnterDataConstruct *
+  Create(const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc,
+         SourceLocation End, ArrayRef<const OpenACCClause *> Clauses);
+};
+// This class represents a 'exit data' construct, which JUST has clauses.
+class OpenACCExitDataConstruct final
+    : public OpenACCConstructStmt,
+      public llvm::TrailingObjects<OpenACCCombinedConstruct,
+                                   const OpenACCClause *> {
+  OpenACCExitDataConstruct(unsigned NumClauses)
+      : OpenACCConstructStmt(OpenACCExitDataConstructClass,
+                             OpenACCDirectiveKind::ExitData, SourceLocation{},
+                             SourceLocation{}, SourceLocation{}) {
+    std::uninitialized_value_construct(
+        getTrailingObjects<const OpenACCClause *>(),
+        getTrailingObjects<const OpenACCClause *>() + NumClauses);
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  NumClauses));
+  }
+  OpenACCExitDataConstruct(SourceLocation Start, SourceLocation DirectiveLoc,
+                           SourceLocation End,
+                           ArrayRef<const OpenACCClause *> Clauses)
+      : OpenACCConstructStmt(OpenACCExitDataConstructClass,
+                             OpenACCDirectiveKind::ExitData, Start,
+                             DirectiveLoc, End) {
+    std::uninitialized_copy(Clauses.begin(), Clauses.end(),
+                            getTrailingObjects<const OpenACCClause *>());
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  Clauses.size()));
+  }
+
+public:
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == OpenACCExitDataConstructClass;
+  }
+  static OpenACCExitDataConstruct *CreateEmpty(const ASTContext &C,
+                                               unsigned NumClauses);
+  static OpenACCExitDataConstruct *
+  Create(const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc,
+         SourceLocation End, ArrayRef<const OpenACCClause *> Clauses);
+};
+// This class represents a 'host_data' construct, which has an associated
+// statement and clauses, but is otherwise pretty simple.
+class OpenACCHostDataConstruct final
+    : public OpenACCAssociatedStmtConstruct,
+      public llvm::TrailingObjects<OpenACCCombinedConstruct,
+                                   const OpenACCClause *> {
+  OpenACCHostDataConstruct(unsigned NumClauses)
+      : OpenACCAssociatedStmtConstruct(
+            OpenACCHostDataConstructClass, OpenACCDirectiveKind::HostData,
+            SourceLocation{}, SourceLocation{}, SourceLocation{},
+            /*AssociatedStmt=*/nullptr) {
+    std::uninitialized_value_construct(
+        getTrailingObjects<const OpenACCClause *>(),
+        getTrailingObjects<const OpenACCClause *>() + NumClauses);
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  NumClauses));
+  }
+  OpenACCHostDataConstruct(SourceLocation Start, SourceLocation DirectiveLoc,
+                           SourceLocation End,
+                           ArrayRef<const OpenACCClause *> Clauses,
+                           Stmt *StructuredBlock)
+      : OpenACCAssociatedStmtConstruct(OpenACCHostDataConstructClass,
+                                       OpenACCDirectiveKind::HostData, Start,
+                                       DirectiveLoc, End, StructuredBlock) {
+    std::uninitialized_copy(Clauses.begin(), Clauses.end(),
+                            getTrailingObjects<const OpenACCClause *>());
+    setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                  Clauses.size()));
+  }
+  void setStructuredBlock(Stmt *S) { setAssociatedStmt(S); }
+
+public:
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == OpenACCHostDataConstructClass;
+  }
+  static OpenACCHostDataConstruct *CreateEmpty(const ASTContext &C,
+                                               unsigned NumClauses);
+  static OpenACCHostDataConstruct *
+  Create(const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc,
+         SourceLocation End, ArrayRef<const OpenACCClause *> Clauses,
+         Stmt *StructuredBlock);
+  Stmt *getStructuredBlock() { return getAssociatedStmt(); }
+  const Stmt *getStructuredBlock() const {
+    return const_cast<OpenACCHostDataConstruct *>(this)->getStructuredBlock();
+  }
+};
 } // namespace clang
 #endif // LLVM_CLANG_AST_STMTOPENACC_H

diff  --git a/clang/include/clang/AST/TextNodeDumper.h b/clang/include/clang/AST/TextNodeDumper.h
index 988b142a7672a3..e54e7e527b8a36 100644
--- a/clang/include/clang/AST/TextNodeDumper.h
+++ b/clang/include/clang/AST/TextNodeDumper.h
@@ -411,6 +411,10 @@ class TextNodeDumper
   void VisitOpenACCConstructStmt(const OpenACCConstructStmt *S);
   void VisitOpenACCLoopConstruct(const OpenACCLoopConstruct *S);
   void VisitOpenACCCombinedConstruct(const OpenACCCombinedConstruct *S);
+  void VisitOpenACCDataConstruct(const OpenACCDataConstruct *S);
+  void VisitOpenACCEnterDataConstruct(const OpenACCEnterDataConstruct *S);
+  void VisitOpenACCExitDataConstruct(const OpenACCExitDataConstruct *S);
+  void VisitOpenACCHostDataConstruct(const OpenACCHostDataConstruct *S);
   void VisitOpenACCAsteriskSizeExpr(const OpenACCAsteriskSizeExpr *S);
   void VisitEmbedExpr(const EmbedExpr *S);
   void VisitAtomicExpr(const AtomicExpr *AE);

diff  --git a/clang/include/clang/Basic/StmtNodes.td b/clang/include/clang/Basic/StmtNodes.td
index 89f5a76eb11312..0c3c580c218fd7 100644
--- a/clang/include/clang/Basic/StmtNodes.td
+++ b/clang/include/clang/Basic/StmtNodes.td
@@ -308,6 +308,10 @@ def OpenACCAssociatedStmtConstruct
 def OpenACCComputeConstruct : StmtNode<OpenACCAssociatedStmtConstruct>;
 def OpenACCLoopConstruct : StmtNode<OpenACCAssociatedStmtConstruct>;
 def OpenACCCombinedConstruct : StmtNode<OpenACCAssociatedStmtConstruct>;
+def OpenACCDataConstruct : StmtNode<OpenACCAssociatedStmtConstruct>;
+def OpenACCEnterDataConstruct : StmtNode<OpenACCConstructStmt>;
+def OpenACCExitDataConstruct : StmtNode<OpenACCConstructStmt>;
+def OpenACCHostDataConstruct : StmtNode<OpenACCAssociatedStmtConstruct>;
 
 // OpenACC Additional Expressions.
 def OpenACCAsteriskSizeExpr : StmtNode<Expr>;

diff  --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h
index af0e08d800bf28..2be9ade08cac31 100644
--- a/clang/include/clang/Serialization/ASTBitCodes.h
+++ b/clang/include/clang/Serialization/ASTBitCodes.h
@@ -2019,6 +2019,10 @@ enum StmtCode {
   STMT_OPENACC_LOOP_CONSTRUCT,
   STMT_OPENACC_COMBINED_CONSTRUCT,
   EXPR_OPENACC_ASTERISK_SIZE,
+  STMT_OPENACC_DATA_CONSTRUCT,
+  STMT_OPENACC_ENTER_DATA_CONSTRUCT,
+  STMT_OPENACC_EXIT_DATA_CONSTRUCT,
+  STMT_OPENACC_HOST_DATA_CONSTRUCT,
 
   // HLSL Constructs
   EXPR_HLSL_OUT_ARG,

diff  --git a/clang/lib/AST/StmtOpenACC.cpp b/clang/lib/AST/StmtOpenACC.cpp
index 23dd57d235813a..fb73dfb3fa9dee 100644
--- a/clang/lib/AST/StmtOpenACC.cpp
+++ b/clang/lib/AST/StmtOpenACC.cpp
@@ -110,3 +110,89 @@ OpenACCCombinedConstruct *OpenACCCombinedConstruct::Create(
       OpenACCCombinedConstruct(DK, BeginLoc, DirLoc, EndLoc, Clauses, Loop);
   return Inst;
 }
+
+OpenACCDataConstruct *OpenACCDataConstruct::CreateEmpty(const ASTContext &C,
+                                                        unsigned NumClauses) {
+  void *Mem =
+      C.Allocate(OpenACCDataConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          NumClauses));
+  auto *Inst = new (Mem) OpenACCDataConstruct(NumClauses);
+  return Inst;
+}
+
+OpenACCDataConstruct *
+OpenACCDataConstruct::Create(const ASTContext &C, SourceLocation Start,
+                             SourceLocation DirectiveLoc, SourceLocation End,
+                             ArrayRef<const OpenACCClause *> Clauses,
+                             Stmt *StructuredBlock) {
+  void *Mem =
+      C.Allocate(OpenACCDataConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          Clauses.size()));
+  auto *Inst = new (Mem)
+      OpenACCDataConstruct(Start, DirectiveLoc, End, Clauses, StructuredBlock);
+  return Inst;
+}
+
+OpenACCEnterDataConstruct *
+OpenACCEnterDataConstruct::CreateEmpty(const ASTContext &C,
+                                       unsigned NumClauses) {
+  void *Mem = C.Allocate(
+      OpenACCEnterDataConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          NumClauses));
+  auto *Inst = new (Mem) OpenACCEnterDataConstruct(NumClauses);
+  return Inst;
+}
+
+OpenACCEnterDataConstruct *OpenACCEnterDataConstruct::Create(
+    const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc,
+    SourceLocation End, ArrayRef<const OpenACCClause *> Clauses) {
+  void *Mem = C.Allocate(
+      OpenACCEnterDataConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          Clauses.size()));
+  auto *Inst =
+      new (Mem) OpenACCEnterDataConstruct(Start, DirectiveLoc, End, Clauses);
+  return Inst;
+}
+
+OpenACCExitDataConstruct *
+OpenACCExitDataConstruct::CreateEmpty(const ASTContext &C,
+                                      unsigned NumClauses) {
+  void *Mem = C.Allocate(
+      OpenACCExitDataConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          NumClauses));
+  auto *Inst = new (Mem) OpenACCExitDataConstruct(NumClauses);
+  return Inst;
+}
+
+OpenACCExitDataConstruct *OpenACCExitDataConstruct::Create(
+    const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc,
+    SourceLocation End, ArrayRef<const OpenACCClause *> Clauses) {
+  void *Mem = C.Allocate(
+      OpenACCExitDataConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          Clauses.size()));
+  auto *Inst =
+      new (Mem) OpenACCExitDataConstruct(Start, DirectiveLoc, End, Clauses);
+  return Inst;
+}
+
+OpenACCHostDataConstruct *
+OpenACCHostDataConstruct::CreateEmpty(const ASTContext &C,
+                                      unsigned NumClauses) {
+  void *Mem = C.Allocate(
+      OpenACCHostDataConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          NumClauses));
+  auto *Inst = new (Mem) OpenACCHostDataConstruct(NumClauses);
+  return Inst;
+}
+
+OpenACCHostDataConstruct *OpenACCHostDataConstruct::Create(
+    const ASTContext &C, SourceLocation Start, SourceLocation DirectiveLoc,
+    SourceLocation End, ArrayRef<const OpenACCClause *> Clauses,
+    Stmt *StructuredBlock) {
+  void *Mem = C.Allocate(
+      OpenACCHostDataConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          Clauses.size()));
+  auto *Inst = new (Mem) OpenACCHostDataConstruct(Start, DirectiveLoc, End,
+                                                  Clauses, StructuredBlock);
+  return Inst;
+}

diff  --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp
index 7507c9d14327a0..488419add5e79e 100644
--- a/clang/lib/AST/StmtPrinter.cpp
+++ b/clang/lib/AST/StmtPrinter.cpp
@@ -1193,6 +1193,51 @@ void StmtPrinter::VisitOpenACCCombinedConstruct(OpenACCCombinedConstruct *S) {
   PrintStmt(S->getLoop());
 }
 
+void StmtPrinter::VisitOpenACCDataConstruct(OpenACCDataConstruct *S) {
+  Indent() << "#pragma acc data";
+
+  if (!S->clauses().empty()) {
+    OS << ' ';
+    OpenACCClausePrinter Printer(OS, Policy);
+    Printer.VisitClauseList(S->clauses());
+  }
+  OS << '\n';
+
+  PrintStmt(S->getStructuredBlock());
+}
+void StmtPrinter::VisitOpenACCEnterDataConstruct(OpenACCEnterDataConstruct *S) {
+  Indent() << "#pragma acc enter data";
+
+  if (!S->clauses().empty()) {
+    OS << ' ';
+    OpenACCClausePrinter Printer(OS, Policy);
+    Printer.VisitClauseList(S->clauses());
+  }
+  OS << '\n';
+}
+void StmtPrinter::VisitOpenACCExitDataConstruct(OpenACCExitDataConstruct *S) {
+  Indent() << "#pragma acc exit data";
+
+  if (!S->clauses().empty()) {
+    OS << ' ';
+    OpenACCClausePrinter Printer(OS, Policy);
+    Printer.VisitClauseList(S->clauses());
+  }
+  OS << '\n';
+}
+void StmtPrinter::VisitOpenACCHostDataConstruct(OpenACCHostDataConstruct *S) {
+  Indent() << "#pragma acc host_data";
+
+  if (!S->clauses().empty()) {
+    OS << ' ';
+    OpenACCClausePrinter Printer(OS, Policy);
+    Printer.VisitClauseList(S->clauses());
+  }
+  OS << '\n';
+
+  PrintStmt(S->getStructuredBlock());
+}
+
 //===----------------------------------------------------------------------===//
 //  Expr printing methods.
 //===----------------------------------------------------------------------===//

diff  --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 3dfbef1cdb712d..e9ff674097c8fa 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2697,6 +2697,37 @@ void StmtProfiler::VisitOpenACCCombinedConstruct(
   P.VisitOpenACCClauseList(S->clauses());
 }
 
+void StmtProfiler::VisitOpenACCDataConstruct(const OpenACCDataConstruct *S) {
+  VisitStmt(S);
+
+  OpenACCClauseProfiler P{*this};
+  P.VisitOpenACCClauseList(S->clauses());
+}
+
+void StmtProfiler::VisitOpenACCEnterDataConstruct(
+    const OpenACCEnterDataConstruct *S) {
+  VisitStmt(S);
+
+  OpenACCClauseProfiler P{*this};
+  P.VisitOpenACCClauseList(S->clauses());
+}
+
+void StmtProfiler::VisitOpenACCExitDataConstruct(
+    const OpenACCExitDataConstruct *S) {
+  VisitStmt(S);
+
+  OpenACCClauseProfiler P{*this};
+  P.VisitOpenACCClauseList(S->clauses());
+}
+
+void StmtProfiler::VisitOpenACCHostDataConstruct(
+    const OpenACCHostDataConstruct *S) {
+  VisitStmt(S);
+
+  OpenACCClauseProfiler P{*this};
+  P.VisitOpenACCClauseList(S->clauses());
+}
+
 void StmtProfiler::VisitHLSLOutArgExpr(const HLSLOutArgExpr *S) {
   VisitStmt(S);
 }

diff  --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index 2552c11a395320..209ad3a5f10ac4 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -2936,6 +2936,25 @@ void TextNodeDumper::VisitOpenACCCombinedConstruct(
   OS << " " << S->getDirectiveKind();
 }
 
+void TextNodeDumper::VisitOpenACCDataConstruct(const OpenACCDataConstruct *S) {
+  OS << " " << S->getDirectiveKind();
+}
+
+void TextNodeDumper::VisitOpenACCEnterDataConstruct(
+    const OpenACCEnterDataConstruct *S) {
+  OS << " " << S->getDirectiveKind();
+}
+
+void TextNodeDumper::VisitOpenACCExitDataConstruct(
+    const OpenACCExitDataConstruct *S) {
+  OS << " " << S->getDirectiveKind();
+}
+
+void TextNodeDumper::VisitOpenACCHostDataConstruct(
+    const OpenACCHostDataConstruct *S) {
+  OS << " " << S->getDirectiveKind();
+}
+
 void TextNodeDumper::VisitEmbedExpr(const EmbedExpr *S) {
   AddChild("begin", [=] { OS << S->getStartingElementPos(); });
   AddChild("number of elements", [=] { OS << S->getDataElementCount(); });

diff  --git a/clang/lib/CodeGen/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp
index 698baf853507f4..6c7a594fb10c4c 100644
--- a/clang/lib/CodeGen/CGStmt.cpp
+++ b/clang/lib/CodeGen/CGStmt.cpp
@@ -458,6 +458,18 @@ void CodeGenFunction::EmitStmt(const Stmt *S, ArrayRef<const Attr *> Attrs) {
   case Stmt::OpenACCCombinedConstructClass:
     EmitOpenACCCombinedConstruct(cast<OpenACCCombinedConstruct>(*S));
     break;
+  case Stmt::OpenACCDataConstructClass:
+    EmitOpenACCDataConstruct(cast<OpenACCDataConstruct>(*S));
+    break;
+  case Stmt::OpenACCEnterDataConstructClass:
+    EmitOpenACCEnterDataConstruct(cast<OpenACCEnterDataConstruct>(*S));
+    break;
+  case Stmt::OpenACCExitDataConstructClass:
+    EmitOpenACCExitDataConstruct(cast<OpenACCExitDataConstruct>(*S));
+    break;
+  case Stmt::OpenACCHostDataConstructClass:
+    EmitOpenACCHostDataConstruct(cast<OpenACCHostDataConstruct>(*S));
+    break;
   }
 }
 

diff  --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h
index eaea0d8a08ac06..092d55355a0a17 100644
--- a/clang/lib/CodeGen/CodeGenFunction.h
+++ b/clang/lib/CodeGen/CodeGenFunction.h
@@ -4094,6 +4094,30 @@ class CodeGenFunction : public CodeGenTypeCache {
     EmitStmt(S.getLoop());
   }
 
+  void EmitOpenACCDataConstruct(const OpenACCDataConstruct &S) {
+    // TODO OpenACC: Implement this.  It is currently implemented as a 'no-op',
+    // simply emitting its structured block, but in the future we will implement
+    // some sort of IR.
+    EmitStmt(S.getStructuredBlock());
+  }
+
+  void EmitOpenACCEnterDataConstruct(const OpenACCEnterDataConstruct &S) {
+    // TODO OpenACC: Implement this.  It is currently implemented as a 'no-op',
+    // but in the future we will implement some sort of IR.
+  }
+
+  void EmitOpenACCExitDataConstruct(const OpenACCExitDataConstruct &S) {
+    // TODO OpenACC: Implement this.  It is currently implemented as a 'no-op',
+    // but in the future we will implement some sort of IR.
+  }
+
+  void EmitOpenACCHostDataConstruct(const OpenACCHostDataConstruct &S) {
+    // TODO OpenACC: Implement this.  It is currently implemented as a 'no-op',
+    // simply emitting its structured block, but in the future we will implement
+    // some sort of IR.
+    EmitStmt(S.getStructuredBlock());
+  }
+
   //===--------------------------------------------------------------------===//
   //                         LValue Expression Emission
   //===--------------------------------------------------------------------===//

diff  --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index bc59de3c1a0ada..8c81936b35296c 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -571,6 +571,8 @@ void SkipUntilEndOfDirective(Parser &P) {
 bool doesDirectiveHaveAssociatedStmt(OpenACCDirectiveKind DirKind) {
   switch (DirKind) {
   default:
+  case OpenACCDirectiveKind::EnterData:
+  case OpenACCDirectiveKind::ExitData:
     return false;
   case OpenACCDirectiveKind::Parallel:
   case OpenACCDirectiveKind::Serial:
@@ -579,6 +581,8 @@ bool doesDirectiveHaveAssociatedStmt(OpenACCDirectiveKind DirKind) {
   case OpenACCDirectiveKind::SerialLoop:
   case OpenACCDirectiveKind::KernelsLoop:
   case OpenACCDirectiveKind::Loop:
+  case OpenACCDirectiveKind::Data:
+  case OpenACCDirectiveKind::HostData:
     return true;
   }
   llvm_unreachable("Unhandled directive->assoc stmt");
@@ -596,6 +600,11 @@ unsigned getOpenACCScopeFlags(OpenACCDirectiveKind DirKind) {
     // so that we can diagnose trying to 'break'/'continue' inside of one.
     return Scope::BreakScope | Scope::ContinueScope |
            Scope::OpenACCComputeConstructScope;
+  case OpenACCDirectiveKind::Data:
+  case OpenACCDirectiveKind::EnterData:
+  case OpenACCDirectiveKind::ExitData:
+  case OpenACCDirectiveKind::HostData:
+    return 0;
   case OpenACCDirectiveKind::Invalid:
     llvm_unreachable("Shouldn't be creating a scope for an invalid construct");
   default:
@@ -1508,10 +1517,10 @@ StmtResult Parser::ParseOpenACCDirectiveStmt() {
     return StmtError();
 
   StmtResult AssocStmt;
-  SemaOpenACC::AssociatedStmtRAII AssocStmtRAII(getActions().OpenACC(),
-                                                DirInfo.DirKind, DirInfo.DirLoc,
-                                                {}, DirInfo.Clauses);
   if (doesDirectiveHaveAssociatedStmt(DirInfo.DirKind)) {
+    SemaOpenACC::AssociatedStmtRAII AssocStmtRAII(
+        getActions().OpenACC(), DirInfo.DirKind, DirInfo.DirLoc, {},
+        DirInfo.Clauses);
     ParsingOpenACCDirectiveRAII DirScope(*this, /*Value=*/false);
     ParseScope ACCScope(this, getOpenACCScopeFlags(DirInfo.DirKind));
 

diff  --git a/clang/lib/Sema/SemaExceptionSpec.cpp b/clang/lib/Sema/SemaExceptionSpec.cpp
index 6a9f43d6f5215e..2be6af293ed543 100644
--- a/clang/lib/Sema/SemaExceptionSpec.cpp
+++ b/clang/lib/Sema/SemaExceptionSpec.cpp
@@ -1396,6 +1396,8 @@ CanThrowResult Sema::canThrow(const Stmt *S) {
   case Expr::ConceptSpecializationExprClass:
   case Expr::RequiresExprClass:
   case Expr::HLSLOutArgExprClass:
+  case Stmt::OpenACCEnterDataConstructClass:
+  case Stmt::OpenACCExitDataConstructClass:
     // These expressions can never throw.
     return CT_Cannot;
 
@@ -1407,6 +1409,8 @@ CanThrowResult Sema::canThrow(const Stmt *S) {
   case Stmt::OpenACCComputeConstructClass:
   case Stmt::OpenACCLoopConstructClass:
   case Stmt::OpenACCCombinedConstructClass:
+  case Stmt::OpenACCDataConstructClass:
+  case Stmt::OpenACCHostDataConstructClass:
   case Stmt::AttributedStmtClass:
   case Stmt::BreakStmtClass:
   case Stmt::CapturedStmtClass:

diff  --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 62c3e778ab178d..5575dd730e5596 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -37,6 +37,10 @@ bool diagnoseConstructAppertainment(SemaOpenACC &S, OpenACCDirectiveKind K,
   case OpenACCDirectiveKind::Serial:
   case OpenACCDirectiveKind::Kernels:
   case OpenACCDirectiveKind::Loop:
+  case OpenACCDirectiveKind::Data:
+  case OpenACCDirectiveKind::EnterData:
+  case OpenACCDirectiveKind::ExitData:
+  case OpenACCDirectiveKind::HostData:
     if (!IsStmt)
       return S.Diag(StartLoc, diag::err_acc_construct_appertainment) << K;
     break;
@@ -1760,6 +1764,31 @@ void CollectActiveReductionClauses(
   }
 }
 
+// Depth needs to be preserved for all associated statements that aren't
+// supposed to modify the compute/combined/loop construct information.
+bool PreserveLoopRAIIDepthInAssociatedStmtRAII(OpenACCDirectiveKind DK) {
+  switch (DK) {
+  case OpenACCDirectiveKind::Parallel:
+  case OpenACCDirectiveKind::ParallelLoop:
+  case OpenACCDirectiveKind::Serial:
+  case OpenACCDirectiveKind::SerialLoop:
+  case OpenACCDirectiveKind::Kernels:
+  case OpenACCDirectiveKind::KernelsLoop:
+  case OpenACCDirectiveKind::Loop:
+    return false;
+  case OpenACCDirectiveKind::Data:
+  case OpenACCDirectiveKind::HostData:
+    return true;
+  case OpenACCDirectiveKind::EnterData:
+  case OpenACCDirectiveKind::ExitData:
+    llvm_unreachable("Doesn't have an associated stmt");
+  default:
+  case OpenACCDirectiveKind::Invalid:
+    llvm_unreachable("Unhandled directive kind?");
+  }
+  llvm_unreachable("Unhandled directive kind?");
+}
+
 } // namespace
 
 SemaOpenACC::SemaOpenACC(Sema &S) : SemaBase(S) {}
@@ -1774,7 +1803,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
       OldLoopVectorClauseLoc(S.LoopVectorClauseLoc),
       OldLoopWithoutSeqInfo(S.LoopWithoutSeqInfo),
       ActiveReductionClauses(S.ActiveReductionClauses),
-      LoopRAII(SemaRef, /*PreserveDepth=*/false) {
+      LoopRAII(SemaRef, PreserveLoopRAIIDepthInAssociatedStmtRAII(DirKind)) {
 
   // Compute constructs end up taking their 'loop'.
   if (DirKind == OpenACCDirectiveKind::Parallel ||
@@ -1950,24 +1979,23 @@ void SemaOpenACC::AssociatedStmtRAII::SetTileInfoBeforeAssociatedStmt(
 }
 
 SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() {
-  SemaRef.ActiveComputeConstructInfo = OldActiveComputeConstructInfo;
-  SemaRef.LoopGangClauseOnKernel = OldLoopGangClauseOnKernel;
-  SemaRef.LoopWorkerClauseLoc = OldLoopWorkerClauseLoc;
-  SemaRef.LoopVectorClauseLoc = OldLoopVectorClauseLoc;
-  SemaRef.LoopWithoutSeqInfo = OldLoopWithoutSeqInfo;
-  SemaRef.ActiveReductionClauses.swap(ActiveReductionClauses);
-
   if (DirKind == OpenACCDirectiveKind::Parallel ||
       DirKind == OpenACCDirectiveKind::Serial ||
       DirKind == OpenACCDirectiveKind::Kernels ||
+      DirKind == OpenACCDirectiveKind::Loop ||
       DirKind == OpenACCDirectiveKind::ParallelLoop ||
       DirKind == OpenACCDirectiveKind::SerialLoop ||
       DirKind == OpenACCDirectiveKind::KernelsLoop) {
-    // Nothing really to do here, the restorations above should be enough for
-    // now.
-  } else if (DirKind == OpenACCDirectiveKind::Loop) {
-    // Nothing really to do here, the LoopInConstruct should handle restorations
-    // correctly.
+    SemaRef.ActiveComputeConstructInfo = OldActiveComputeConstructInfo;
+    SemaRef.LoopGangClauseOnKernel = OldLoopGangClauseOnKernel;
+    SemaRef.LoopWorkerClauseLoc = OldLoopWorkerClauseLoc;
+    SemaRef.LoopVectorClauseLoc = OldLoopVectorClauseLoc;
+    SemaRef.LoopWithoutSeqInfo = OldLoopWithoutSeqInfo;
+    SemaRef.ActiveReductionClauses.swap(ActiveReductionClauses);
+  } else if (DirKind == OpenACCDirectiveKind::Data ||
+             DirKind == OpenACCDirectiveKind::HostData) {
+    // Intentionally doesn't reset the Loop, Compute Construct, or reduction
+    // effects.
   }
 }
 
@@ -2175,6 +2203,10 @@ void SemaOpenACC::ActOnConstruct(OpenACCDirectiveKind K,
   case OpenACCDirectiveKind::SerialLoop:
   case OpenACCDirectiveKind::KernelsLoop:
   case OpenACCDirectiveKind::Loop:
+  case OpenACCDirectiveKind::Data:
+  case OpenACCDirectiveKind::EnterData:
+  case OpenACCDirectiveKind::ExitData:
+  case OpenACCDirectiveKind::HostData:
     // Nothing to do here, there is no real legalization that needs to happen
     // here as these constructs do not take any arguments.
     break;
@@ -3441,6 +3473,24 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective(OpenACCDirectiveKind K,
         getASTContext(), ActiveComputeConstructInfo.Kind, StartLoc, DirLoc,
         EndLoc, Clauses, AssocStmt.isUsable() ? AssocStmt.get() : nullptr);
   }
+  case OpenACCDirectiveKind::Data: {
+    return OpenACCDataConstruct::Create(
+        getASTContext(), StartLoc, DirLoc, EndLoc, Clauses,
+        AssocStmt.isUsable() ? AssocStmt.get() : nullptr);
+  }
+  case OpenACCDirectiveKind::EnterData: {
+    return OpenACCEnterDataConstruct::Create(getASTContext(), StartLoc, DirLoc,
+                                             EndLoc, Clauses);
+  }
+  case OpenACCDirectiveKind::ExitData: {
+    return OpenACCExitDataConstruct::Create(getASTContext(), StartLoc, DirLoc,
+                                            EndLoc, Clauses);
+  }
+  case OpenACCDirectiveKind::HostData: {
+    return OpenACCHostDataConstruct::Create(
+        getASTContext(), StartLoc, DirLoc, EndLoc, Clauses,
+        AssocStmt.isUsable() ? AssocStmt.get() : nullptr);
+  }
   }
   llvm_unreachable("Unhandled case in directive handling?");
 }
@@ -3451,9 +3501,15 @@ StmtResult SemaOpenACC::ActOnAssociatedStmt(
   switch (K) {
   default:
     llvm_unreachable("Unimplemented associated statement application");
+  case OpenACCDirectiveKind::EnterData:
+  case OpenACCDirectiveKind::ExitData:
+    llvm_unreachable(
+        "these don't have associated statements, so shouldn't get here");
   case OpenACCDirectiveKind::Parallel:
   case OpenACCDirectiveKind::Serial:
   case OpenACCDirectiveKind::Kernels:
+  case OpenACCDirectiveKind::Data:
+  case OpenACCDirectiveKind::HostData:
     // There really isn't any checking here that could happen. As long as we
     // have a statement to associate, this should be fine.
     // OpenACC 3.3 Section 6:

diff  --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 02d2fc018e3c35..f2dbf4086a13da 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -4110,6 +4110,42 @@ class TreeTransform {
                                                      EndLoc, Clauses, Loop);
   }
 
+  StmtResult RebuildOpenACCDataConstruct(SourceLocation BeginLoc,
+                                         SourceLocation DirLoc,
+                                         SourceLocation EndLoc,
+                                         ArrayRef<OpenACCClause *> Clauses,
+                                         StmtResult StrBlock) {
+    return getSema().OpenACC().ActOnEndStmtDirective(OpenACCDirectiveKind::Data,
+                                                     BeginLoc, DirLoc, EndLoc,
+                                                     Clauses, StrBlock);
+  }
+
+  StmtResult
+  RebuildOpenACCEnterDataConstruct(SourceLocation BeginLoc,
+                                   SourceLocation DirLoc, SourceLocation EndLoc,
+                                   ArrayRef<OpenACCClause *> Clauses) {
+    return getSema().OpenACC().ActOnEndStmtDirective(
+        OpenACCDirectiveKind::EnterData, BeginLoc, DirLoc, EndLoc, Clauses, {});
+  }
+
+  StmtResult
+  RebuildOpenACCExitDataConstruct(SourceLocation BeginLoc,
+                                  SourceLocation DirLoc, SourceLocation EndLoc,
+                                  ArrayRef<OpenACCClause *> Clauses) {
+    return getSema().OpenACC().ActOnEndStmtDirective(
+        OpenACCDirectiveKind::ExitData, BeginLoc, DirLoc, EndLoc, Clauses, {});
+  }
+
+  StmtResult RebuildOpenACCHostDataConstruct(SourceLocation BeginLoc,
+                                             SourceLocation DirLoc,
+                                             SourceLocation EndLoc,
+                                             ArrayRef<OpenACCClause *> Clauses,
+                                             StmtResult StrBlock) {
+    return getSema().OpenACC().ActOnEndStmtDirective(
+        OpenACCDirectiveKind::HostData, BeginLoc, DirLoc, EndLoc, Clauses,
+        StrBlock);
+  }
+
   ExprResult RebuildOpenACCAsteriskSizeExpr(SourceLocation AsteriskLoc) {
     return getSema().OpenACC().ActOnOpenACCAsteriskSizeExpr(AsteriskLoc);
   }
@@ -12153,6 +12189,88 @@ StmtResult TreeTransform<Derived>::TransformOpenACCCombinedConstruct(
       C->getEndLoc(), TransformedClauses, Loop);
 }
 
+template <typename Derived>
+StmtResult
+TreeTransform<Derived>::TransformOpenACCDataConstruct(OpenACCDataConstruct *C) {
+  getSema().OpenACC().ActOnConstruct(C->getDirectiveKind(), C->getBeginLoc());
+
+  llvm::SmallVector<OpenACCClause *> TransformedClauses =
+      getDerived().TransformOpenACCClauseList(C->getDirectiveKind(),
+                                              C->clauses());
+  if (getSema().OpenACC().ActOnStartStmtDirective(C->getDirectiveKind(),
+                                                  C->getBeginLoc()))
+    return StmtError();
+
+  SemaOpenACC::AssociatedStmtRAII AssocStmtRAII(
+      getSema().OpenACC(), C->getDirectiveKind(), C->getDirectiveLoc(),
+      C->clauses(), TransformedClauses);
+  StmtResult StrBlock = getDerived().TransformStmt(C->getStructuredBlock());
+  StrBlock = getSema().OpenACC().ActOnAssociatedStmt(
+      C->getBeginLoc(), C->getDirectiveKind(), TransformedClauses, StrBlock);
+
+  return getDerived().RebuildOpenACCDataConstruct(
+      C->getBeginLoc(), C->getDirectiveLoc(), C->getEndLoc(),
+      TransformedClauses, StrBlock);
+}
+
+template <typename Derived>
+StmtResult TreeTransform<Derived>::TransformOpenACCEnterDataConstruct(
+    OpenACCEnterDataConstruct *C) {
+  getSema().OpenACC().ActOnConstruct(C->getDirectiveKind(), C->getBeginLoc());
+
+  llvm::SmallVector<OpenACCClause *> TransformedClauses =
+      getDerived().TransformOpenACCClauseList(C->getDirectiveKind(),
+                                              C->clauses());
+  if (getSema().OpenACC().ActOnStartStmtDirective(C->getDirectiveKind(),
+                                                  C->getBeginLoc()))
+    return StmtError();
+
+  return getDerived().RebuildOpenACCEnterDataConstruct(
+      C->getBeginLoc(), C->getDirectiveLoc(), C->getEndLoc(),
+      TransformedClauses);
+}
+
+template <typename Derived>
+StmtResult TreeTransform<Derived>::TransformOpenACCExitDataConstruct(
+    OpenACCExitDataConstruct *C) {
+  getSema().OpenACC().ActOnConstruct(C->getDirectiveKind(), C->getBeginLoc());
+
+  llvm::SmallVector<OpenACCClause *> TransformedClauses =
+      getDerived().TransformOpenACCClauseList(C->getDirectiveKind(),
+                                              C->clauses());
+  if (getSema().OpenACC().ActOnStartStmtDirective(C->getDirectiveKind(),
+                                                  C->getBeginLoc()))
+    return StmtError();
+
+  return getDerived().RebuildOpenACCExitDataConstruct(
+      C->getBeginLoc(), C->getDirectiveLoc(), C->getEndLoc(),
+      TransformedClauses);
+}
+
+template <typename Derived>
+StmtResult TreeTransform<Derived>::TransformOpenACCHostDataConstruct(
+    OpenACCHostDataConstruct *C) {
+  getSema().OpenACC().ActOnConstruct(C->getDirectiveKind(), C->getBeginLoc());
+
+  llvm::SmallVector<OpenACCClause *> TransformedClauses =
+      getDerived().TransformOpenACCClauseList(C->getDirectiveKind(),
+                                              C->clauses());
+  if (getSema().OpenACC().ActOnStartStmtDirective(C->getDirectiveKind(),
+                                                  C->getBeginLoc()))
+    return StmtError();
+
+  SemaOpenACC::AssociatedStmtRAII AssocStmtRAII(
+      getSema().OpenACC(), C->getDirectiveKind(), C->getDirectiveLoc(),
+      C->clauses(), TransformedClauses);
+  StmtResult StrBlock = getDerived().TransformStmt(C->getStructuredBlock());
+  StrBlock = getSema().OpenACC().ActOnAssociatedStmt(
+      C->getBeginLoc(), C->getDirectiveKind(), TransformedClauses, StrBlock);
+
+  return getDerived().RebuildOpenACCHostDataConstruct(
+      C->getBeginLoc(), C->getDirectiveLoc(), C->getEndLoc(),
+      TransformedClauses, StrBlock);
+}
+
 template <typename Derived>
 ExprResult TreeTransform<Derived>::TransformOpenACCAsteriskSizeExpr(
     OpenACCAsteriskSizeExpr *E) {

diff  --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp
index 9f4877b19d8705..21ad6c5a9faa33 100644
--- a/clang/lib/Serialization/ASTReaderStmt.cpp
+++ b/clang/lib/Serialization/ASTReaderStmt.cpp
@@ -2849,6 +2849,27 @@ void ASTStmtReader::VisitOpenACCCombinedConstruct(OpenACCCombinedConstruct *S) {
   VisitOpenACCAssociatedStmtConstruct(S);
 }
 
+void ASTStmtReader::VisitOpenACCDataConstruct(OpenACCDataConstruct *S) {
+  VisitStmt(S);
+  VisitOpenACCAssociatedStmtConstruct(S);
+}
+
+void ASTStmtReader::VisitOpenACCEnterDataConstruct(
+    OpenACCEnterDataConstruct *S) {
+  VisitStmt(S);
+  VisitOpenACCConstructStmt(S);
+}
+
+void ASTStmtReader::VisitOpenACCExitDataConstruct(OpenACCExitDataConstruct *S) {
+  VisitStmt(S);
+  VisitOpenACCConstructStmt(S);
+}
+
+void ASTStmtReader::VisitOpenACCHostDataConstruct(OpenACCHostDataConstruct *S) {
+  VisitStmt(S);
+  VisitOpenACCAssociatedStmtConstruct(S);
+}
+
 //===----------------------------------------------------------------------===//
 // HLSL Constructs/Directives.
 //===----------------------------------------------------------------------===//
@@ -4324,6 +4345,26 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
       S = OpenACCCombinedConstruct::CreateEmpty(Context, NumClauses);
       break;
     }
+    case STMT_OPENACC_DATA_CONSTRUCT: {
+      unsigned NumClauses = Record[ASTStmtReader::NumStmtFields];
+      S = OpenACCDataConstruct::CreateEmpty(Context, NumClauses);
+      break;
+    }
+    case STMT_OPENACC_ENTER_DATA_CONSTRUCT: {
+      unsigned NumClauses = Record[ASTStmtReader::NumStmtFields];
+      S = OpenACCEnterDataConstruct::CreateEmpty(Context, NumClauses);
+      break;
+    }
+    case STMT_OPENACC_EXIT_DATA_CONSTRUCT: {
+      unsigned NumClauses = Record[ASTStmtReader::NumStmtFields];
+      S = OpenACCExitDataConstruct::CreateEmpty(Context, NumClauses);
+      break;
+    }
+    case STMT_OPENACC_HOST_DATA_CONSTRUCT: {
+      unsigned NumClauses = Record[ASTStmtReader::NumStmtFields];
+      S = OpenACCHostDataConstruct::CreateEmpty(Context, NumClauses);
+      break;
+    }
     case EXPR_REQUIRES: {
       unsigned numLocalParameters = Record[ASTStmtReader::NumExprFields];
       unsigned numRequirement = Record[ASTStmtReader::NumExprFields + 1];

diff  --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp
index 603aa5707ce9be..e55cbe1f6ecce6 100644
--- a/clang/lib/Serialization/ASTWriterStmt.cpp
+++ b/clang/lib/Serialization/ASTWriterStmt.cpp
@@ -2926,6 +2926,31 @@ void ASTStmtWriter::VisitOpenACCCombinedConstruct(OpenACCCombinedConstruct *S) {
   Code = serialization::STMT_OPENACC_COMBINED_CONSTRUCT;
 }
 
+void ASTStmtWriter::VisitOpenACCDataConstruct(OpenACCDataConstruct *S) {
+  VisitStmt(S);
+  VisitOpenACCAssociatedStmtConstruct(S);
+  Code = serialization::STMT_OPENACC_DATA_CONSTRUCT;
+}
+
+void ASTStmtWriter::VisitOpenACCEnterDataConstruct(
+    OpenACCEnterDataConstruct *S) {
+  VisitStmt(S);
+  VisitOpenACCConstructStmt(S);
+  Code = serialization::STMT_OPENACC_ENTER_DATA_CONSTRUCT;
+}
+
+void ASTStmtWriter::VisitOpenACCExitDataConstruct(OpenACCExitDataConstruct *S) {
+  VisitStmt(S);
+  VisitOpenACCConstructStmt(S);
+  Code = serialization::STMT_OPENACC_EXIT_DATA_CONSTRUCT;
+}
+
+void ASTStmtWriter::VisitOpenACCHostDataConstruct(OpenACCHostDataConstruct *S) {
+  VisitStmt(S);
+  VisitOpenACCAssociatedStmtConstruct(S);
+  Code = serialization::STMT_OPENACC_HOST_DATA_CONSTRUCT;
+}
+
 //===----------------------------------------------------------------------===//
 // HLSL Constructs/Directives.
 //===----------------------------------------------------------------------===//

diff  --git a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
index b46cd9fe86fc11..ae43c59511bfa7 100644
--- a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
+++ b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
@@ -1825,6 +1825,10 @@ void ExprEngine::Visit(const Stmt *S, ExplodedNode *Pred,
     case Stmt::OpenACCComputeConstructClass:
     case Stmt::OpenACCLoopConstructClass:
     case Stmt::OpenACCCombinedConstructClass:
+    case Stmt::OpenACCDataConstructClass:
+    case Stmt::OpenACCEnterDataConstructClass:
+    case Stmt::OpenACCExitDataConstructClass:
+    case Stmt::OpenACCHostDataConstructClass:
     case Stmt::OMPUnrollDirectiveClass:
     case Stmt::OMPMetaDirectiveClass:
     case Stmt::HLSLOutArgExprClass: {

diff  --git a/clang/test/AST/ast-print-openacc-data-construct.cpp b/clang/test/AST/ast-print-openacc-data-construct.cpp
new file mode 100644
index 00000000000000..fc15add15c6b89
--- /dev/null
+++ b/clang/test/AST/ast-print-openacc-data-construct.cpp
@@ -0,0 +1,25 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-deprecated-clause-alias -Wno-source-uses-openacc -ast-print %s -o - | FileCheck %s
+
+void foo() {
+  int Var;
+  // TODO OpenACC: These are only legal if they have one of a list of clauses on
+  // them, so the 'check' lines should start to include those once we implement
+  // them.  For now, they don't emit those because they are 'not implemented'.
+
+// CHECK: #pragma acc data
+// CHECK-NOT: default(none)
+#pragma acc data default(none)
+  ;
+// CHECK: #pragma acc enter data
+// CHECK-NOT: copyin(Var)
+#pragma acc enter data copyin(Var)
+  ;
+// CHECK: #pragma acc exit data
+// CHECK-NOT: copyout(Var)
+#pragma acc exit data copyout(Var)
+  ;
+// CHECK: #pragma acc host_data
+// CHECK-NOT: use_device(Var)
+#pragma acc host_data use_device(Var)
+  ;
+}

diff  --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 3741ed099cf5c2..e2f0a753dd3780 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -4,37 +4,30 @@
 
 void func() {
 
-  // expected-warning at +2{{OpenACC clause 'finalize' not yet implemented, clause ignored}}
-  // expected-warning at +1{{OpenACC construct 'enter data' not yet implemented, pragma ignored}}
+  // expected-warning at +1{{OpenACC clause 'finalize' not yet implemented, clause ignored}}
 #pragma acc enter data finalize
 
-  // expected-warning at +3{{OpenACC clause 'finalize' not yet implemented, clause ignored}}
   // expected-warning at +2{{OpenACC clause 'finalize' not yet implemented, clause ignored}}
-  // expected-warning at +1{{OpenACC construct 'enter data' not yet implemented, pragma ignored}}
+  // expected-warning at +1{{OpenACC clause 'finalize' not yet implemented, clause ignored}}
 #pragma acc enter data finalize finalize
 
-  // expected-warning at +3{{OpenACC clause 'finalize' not yet implemented, clause ignored}}
-  // expected-error at +2{{invalid OpenACC clause 'invalid'}}
-  // expected-warning at +1{{OpenACC construct 'enter data' not yet implemented, pragma ignored}}
+  // expected-warning at +2{{OpenACC clause 'finalize' not yet implemented, clause ignored}}
+  // expected-error at +1{{invalid OpenACC clause 'invalid'}}
 #pragma acc enter data finalize invalid
 
-  // expected-warning at +3{{OpenACC clause 'finalize' not yet implemented, clause ignored}}
-  // expected-error at +2{{invalid OpenACC clause 'invalid'}}
-  // expected-warning at +1{{OpenACC construct 'enter data' not yet implemented, pragma ignored}}
+  // expected-warning at +2{{OpenACC clause 'finalize' not yet implemented, clause ignored}}
+  // expected-error at +1{{invalid OpenACC clause 'invalid'}}
 #pragma acc enter data finalize invalid invalid finalize
 
-  // expected-warning at +3{{OpenACC clause 'wait' not yet implemented, clause ignored}}
-  // expected-warning at +2{{OpenACC clause 'finalize' not yet implemented, clause ignored}}
-  // expected-warning at +1{{OpenACC construct 'enter data' not yet implemented, pragma ignored}}
+  // expected-warning at +2{{OpenACC clause 'wait' not yet implemented, clause ignored}}
+  // expected-warning at +1{{OpenACC clause 'finalize' not yet implemented, clause ignored}}
 #pragma acc enter data wait finalize
 
-  // expected-warning at +2{{OpenACC clause 'if_present' not yet implemented, clause ignored}}
-  // expected-warning at +1{{OpenACC construct 'host_data' not yet implemented, pragma ignored}}
+  // expected-warning at +1{{OpenACC clause 'if_present' not yet implemented, clause ignored}}
 #pragma acc host_data if_present
 
-  // expected-warning at +3{{OpenACC clause 'if_present' not yet implemented, clause ignored}}
   // expected-warning at +2{{OpenACC clause 'if_present' not yet implemented, clause ignored}}
-  // expected-warning at +1{{OpenACC construct 'host_data' not yet implemented, pragma ignored}}
+  // expected-warning at +1{{OpenACC clause 'if_present' not yet implemented, clause ignored}}
 #pragma acc host_data if_present, if_present
 
   // expected-error at +4{{OpenACC clause 'independent' on 'loop' construct conflicts with previous data dependence clause}}
@@ -528,27 +521,23 @@ void VarListClauses() {
 #pragma acc serial firstprivate(s.array[s.value : 5], s.value), self
   for(int i = 0; i < 5;++i) {}
 
-  // expected-warning at +4{{OpenACC construct 'exit data' not yet implemented}}
   // expected-error at +3{{expected ','}}
   // expected-warning at +2{{OpenACC clause 'delete' not yet implemented, clause ignored}}
   // expected-warning at +1{{OpenACC clause 'async' not yet implemented, clause ignored}}
 #pragma acc exit data delete(s.array[s.value] s.array[s.value :5] ) async
   for(int i = 0; i < 5;++i) {}
 
-  // expected-warning at +3{{OpenACC construct 'exit data' not yet implemented}}
   // expected-warning at +2{{OpenACC clause 'delete' not yet implemented, clause ignored}}
   // expected-warning at +1{{OpenACC clause 'async' not yet implemented, clause ignored}}
 #pragma acc exit data delete(s.array[s.value : 5], s.value),async
   for(int i = 0; i < 5;++i) {}
 
-  // expected-warning at +4{{OpenACC construct 'exit data' not yet implemented}}
   // expected-error at +3{{expected ','}}
   // expected-warning at +2{{OpenACC clause 'use_device' not yet implemented, clause ignored}}
   // expected-warning at +1{{OpenACC clause 'async' not yet implemented, clause ignored}}
 #pragma acc exit data use_device(s.array[s.value] s.array[s.value :5] ),async
   for(int i = 0; i < 5;++i) {}
 
-  // expected-warning at +3{{OpenACC construct 'exit data' not yet implemented}}
   // expected-warning at +2{{OpenACC clause 'use_device' not yet implemented, clause ignored}}
   // expected-warning at +1{{OpenACC clause 'async' not yet implemented, clause ignored}}
 #pragma acc exit data use_device(s.array[s.value : 5], s.value), async

diff  --git a/clang/test/ParserOpenACC/parse-clauses.cpp b/clang/test/ParserOpenACC/parse-clauses.cpp
index 4dc966ea9879f9..1781a279407543 100644
--- a/clang/test/ParserOpenACC/parse-clauses.cpp
+++ b/clang/test/ParserOpenACC/parse-clauses.cpp
@@ -35,7 +35,6 @@ void templ() {
 #pragma acc parallel async
   for(;;){}
 
-  // expected-warning at +2{{OpenACC construct 'exit data' not yet implemented}}
   // expected-warning at +1{{OpenACC clause 'delete' not yet implemented, clause ignored}}
 #pragma acc exit data delete(I)
   ;

diff  --git a/clang/test/ParserOpenACC/parse-constructs.c b/clang/test/ParserOpenACC/parse-constructs.c
index 27b9a6993fd3e8..d3b1ccb48c0349 100644
--- a/clang/test/ParserOpenACC/parse-constructs.c
+++ b/clang/test/ParserOpenACC/parse-constructs.c
@@ -54,16 +54,13 @@ void func() {
   // expected-error at +1{{invalid OpenACC clause 'clause'}}
 #pragma acc kernels clause list
   for(;;){}
-  // expected-error at +2{{invalid OpenACC clause 'clause'}}
-  // expected-warning at +1{{OpenACC construct 'data' not yet implemented, pragma ignored}}
+  // expected-error at +1{{invalid OpenACC clause 'clause'}}
 #pragma acc data clause list
   for(;;){}
-  // expected-error at +2{{invalid OpenACC clause 'clause'}}
-  // expected-warning at +1{{OpenACC construct 'enter data' not yet implemented, pragma ignored}}
+  // expected-error at +1{{invalid OpenACC clause 'clause'}}
 #pragma acc enter data clause list
   for(;;){}
-  // expected-error at +2{{invalid OpenACC clause 'clause'}}
-  // expected-warning at +1{{OpenACC construct 'exit data' not yet implemented, pragma ignored}}
+  // expected-error at +1{{invalid OpenACC clause 'clause'}}
 #pragma acc exit data clause list
   for(;;){}
   // expected-error at +1{{invalid OpenACC directive 'enter invalid'}}
@@ -78,8 +75,7 @@ void func() {
   // expected-error at +1{{expected identifier}}
 #pragma acc exit }
   for(;;){}
-  // expected-error at +2{{invalid OpenACC clause 'clause'}}
-  // expected-warning at +1{{OpenACC construct 'host_data' not yet implemented, pragma ignored}}
+  // expected-error at +1{{invalid OpenACC clause 'clause'}}
 #pragma acc host_data clause list
   for(;;){}
   // expected-error at +1{{invalid OpenACC clause 'clause'}}

diff  --git a/clang/test/SemaOpenACC/combined-construct-collapse-clause.cpp b/clang/test/SemaOpenACC/combined-construct-collapse-clause.cpp
index c7db9669a9879b..31078ea7a0de9e 100644
--- a/clang/test/SemaOpenACC/combined-construct-collapse-clause.cpp
+++ b/clang/test/SemaOpenACC/combined-construct-collapse-clause.cpp
@@ -214,14 +214,15 @@ void no_other_directives() {
 #pragma acc serial loop collapse(2)
   for(unsigned i = 0; i < 5; ++i) {
     for(unsigned j = 0; j < 5; ++j) {
-#pragma acc data // expected-warning{{OpenACC construct 'data' not yet implemented}}
+#pragma acc data
+      ;
     }
   }
   // expected-note at +1{{active 'collapse' clause defined here}}
 #pragma acc kernels loop collapse(2)
   for(unsigned i = 0; i < 5; ++i) {
     // expected-error at +1{{OpenACC 'data' construct cannot appear in intervening code of a 'kernels loop' with a 'collapse' clause}}
-#pragma acc data // expected-warning{{OpenACC construct 'data' not yet implemented}}
+#pragma acc data
     for(unsigned j = 0; j < 5; ++j) {
     }
   }

diff  --git a/clang/test/SemaOpenACC/combined-construct-default-clause.c b/clang/test/SemaOpenACC/combined-construct-default-clause.c
index a9c90240cb1222..7e384ccfc17a07 100644
--- a/clang/test/SemaOpenACC/combined-construct-default-clause.c
+++ b/clang/test/SemaOpenACC/combined-construct-default-clause.c
@@ -28,7 +28,6 @@ void SingleOnly() {
   #pragma acc kernels loop default(none)
   for(int i = 0; i < 5; ++i);
 
-  // expected-warning at +2{{OpenACC construct 'data' not yet implemented}}
   // expected-warning at +1{{OpenACC clause 'default' not yet implemented}}
   #pragma acc data default(none)
   while(0);

diff  --git a/clang/test/SemaOpenACC/combined-construct-if-clause.c b/clang/test/SemaOpenACC/combined-construct-if-clause.c
index 563f1cd25377bd..c0069c4ee9ef44 100644
--- a/clang/test/SemaOpenACC/combined-construct-if-clause.c
+++ b/clang/test/SemaOpenACC/combined-construct-if-clause.c
@@ -43,7 +43,6 @@ void BoolExpr(int *I, float *F) {
 #pragma acc kernels loop if (*I < *F)
   for (unsigned i = 0; i < 5; ++i);
 
-  // expected-warning at +2{{OpenACC construct 'data' not yet implemented}}
   // expected-warning at +1{{OpenACC clause 'if' not yet implemented}}
 #pragma acc data if (*I < *F)
   for (unsigned i = 0; i < 5; ++i);

diff  --git a/clang/test/SemaOpenACC/compute-construct-default-clause.c b/clang/test/SemaOpenACC/compute-construct-default-clause.c
index 70e29f3e8ac051..4aef2cbd7aec4c 100644
--- a/clang/test/SemaOpenACC/compute-construct-default-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-default-clause.c
@@ -28,7 +28,6 @@ void SingleOnly() {
   #pragma acc kernels default(none)
   for(int i = 0; i < 5; ++i);
 
-  // expected-warning at +2{{OpenACC construct 'data' not yet implemented}}
   // expected-warning at +1{{OpenACC clause 'default' not yet implemented}}
   #pragma acc data default(none)
   while(0);

diff  --git a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
index 0ae972d2a99ff4..2f4a037529b500 100644
--- a/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-device_type-clause.c
@@ -34,11 +34,9 @@ void uses() {
 #pragma acc kernels dtype(MACRO)
   while(1);
 
-  // expected-error at +2{{OpenACC 'device_type' clause is not valid on 'enter data' directive}}
-  // expected-warning at +1{{OpenACC construct 'enter data' not yet implemented}}
+  // expected-error at +1{{OpenACC 'device_type' clause is not valid on 'enter data' directive}}
 #pragma acc enter data device_type(I)
-  // expected-error at +2{{OpenACC 'dtype' clause is not valid on 'enter data' directive}}
-  // expected-warning at +1{{OpenACC construct 'enter data' not yet implemented}}
+  // expected-error at +1{{OpenACC 'dtype' clause is not valid on 'enter data' directive}}
 #pragma acc enter data dtype(I)
 
 

diff  --git a/clang/test/SemaOpenACC/compute-construct-if-clause.c b/clang/test/SemaOpenACC/compute-construct-if-clause.c
index 7cdc35275acce0..1336bf2549f3ca 100644
--- a/clang/test/SemaOpenACC/compute-construct-if-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-if-clause.c
@@ -43,7 +43,6 @@ void BoolExpr(int *I, float *F) {
 #pragma acc kernels if (*I < *F)
   while(0);
 
-  // expected-warning at +2{{OpenACC construct 'data' not yet implemented}}
   // expected-warning at +1{{OpenACC clause 'if' not yet implemented}}
 #pragma acc data if (*I < *F)
   while(0);

diff  --git a/clang/test/SemaOpenACC/data-construct-ast.cpp b/clang/test/SemaOpenACC/data-construct-ast.cpp
new file mode 100644
index 00000000000000..9a7fe2cb793a73
--- /dev/null
+++ b/clang/test/SemaOpenACC/data-construct-ast.cpp
@@ -0,0 +1,88 @@
+// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s
+
+// Test this with PCH.
+// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s
+// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s
+
+#ifndef PCH_HELPER
+#define PCH_HELPER
+
+void NormalFunc() {
+  // CHECK-LABEL: NormalFunc
+  // CHECK-NEXT: CompoundStmt
+
+  int Var;
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl
+
+  // TODO OpenACC: these constructs require the clauses to be legal, but we
+  // don't have the clauses implemented yet.  As we implement them, they needed
+  // to be added to the 'check' lines.
+
+#pragma acc data default(none)
+  while (Var);
+  // CHECK-NEXT: OpenACCDataConstruct{{.*}}data
+  // CHECK-NEXT: WhileStmt
+  // CHECK: NullStmt
+#pragma acc enter data copyin(Var)
+  // CHECK-NEXT: OpenACCEnterDataConstruct{{.*}} enter data
+#pragma acc exit data copyout(Var)
+  // CHECK-NEXT: OpenACCExitDataConstruct{{.*}} exit data
+#pragma acc host_data use_device(Var)
+  while (Var);
+  // CHECK-NEXT: OpenACCHostDataConstruct{{.*}} host_data
+  // CHECK-NEXT: WhileStmt
+  // CHECK: NullStmt
+}
+
+template<typename T>
+void TemplFunc() {
+  // CHECK-LABEL: FunctionTemplateDecl {{.*}}TemplFunc
+  // CHECK-NEXT: TemplateTypeParmDecl
+  // CHECK-NEXT: FunctionDecl{{.*}}TemplFunc
+  // CHECK-NEXT: CompoundStmt
+
+  T Var;
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl
+
+#pragma acc data default(none)
+  while (Var);
+  // CHECK-NEXT: OpenACCDataConstruct{{.*}}data
+  // CHECK-NEXT: WhileStmt
+  // CHECK: NullStmt
+#pragma acc enter data copyin(Var)
+  // CHECK-NEXT: OpenACCEnterDataConstruct{{.*}} enter data
+#pragma acc exit data copyout(Var)
+  // CHECK-NEXT: OpenACCExitDataConstruct{{.*}} exit data
+#pragma acc host_data use_device(Var)
+  while (Var);
+  // CHECK-NEXT: OpenACCHostDataConstruct{{.*}} host_data
+  // CHECK-NEXT: WhileStmt
+  // CHECK: NullStmt
+
+  // Instantiation:
+  // CHECK-NEXT: FunctionDecl{{.*}} TemplFunc 'void ()' implicit_instantiation
+  // CHECK-NEXT: TemplateArgument type 'int'
+  // CHECK-NEXT: BuiltinType{{.*}} 'int'
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl
+
+  // CHECK-NEXT: OpenACCDataConstruct{{.*}}data
+  // CHECK-NEXT: WhileStmt
+  // CHECK: NullStmt
+
+  // CHECK-NEXT: OpenACCEnterDataConstruct{{.*}} enter data
+
+  // CHECK-NEXT: OpenACCExitDataConstruct{{.*}} exit data
+
+  // CHECK-NEXT: OpenACCHostDataConstruct{{.*}} host_data
+  // CHECK-NEXT: WhileStmt
+  // CHECK: NullStmt
+}
+void use() {
+  TemplFunc<int>();
+}
+#endif

diff  --git a/clang/test/SemaOpenACC/data-construct.cpp b/clang/test/SemaOpenACC/data-construct.cpp
new file mode 100644
index 00000000000000..0c1959dc427248
--- /dev/null
+++ b/clang/test/SemaOpenACC/data-construct.cpp
@@ -0,0 +1,237 @@
+// RUN: %clang_cc1 %s -fopenacc -verify -Wno-empty-body -Wno-unused-value
+
+void HasStmt() {
+  {
+    // expected-error at +2{{expected statement}}
+#pragma acc data
+  }
+  {
+    // expected-error at +2{{expected statement}}
+#pragma acc host_data
+  }
+  // Don't have statements, so this is fine.
+  {
+#pragma acc enter data
+  }
+  {
+#pragma acc exit data
+  }
+}
+
+void AtLeastOneOf() {
+  int Var;
+// Data
+  // expected-warning at +1{{OpenACC clause 'copy' not yet implemented}}
+#pragma acc data copy(Var)
+  ;
+  // expected-warning at +1{{OpenACC clause 'copyin' not yet implemented}}
+#pragma acc data copyin(Var)
+  ;
+  // expected-warning at +1{{OpenACC clause 'copyout' not yet implemented}}
+#pragma acc data copyout(Var)
+  ;
+  // expected-warning at +1{{OpenACC clause 'create' not yet implemented}}
+#pragma acc data create(Var)
+  ;
+  // expected-warning at +1{{OpenACC clause 'no_create' not yet implemented}}
+#pragma acc data no_create(Var)
+  ;
+  // expected-warning at +1{{OpenACC clause 'present' not yet implemented}}
+#pragma acc data present(Var)
+  ;
+  // expected-warning at +1{{OpenACC clause 'deviceptr' not yet implemented}}
+#pragma acc data deviceptr(Var)
+  ;
+  // expected-warning at +1{{OpenACC clause 'attach' not yet implemented}}
+#pragma acc data attach(Var)
+  ;
+  // expected-warning at +1{{OpenACC clause 'default' not yet implemented}}
+#pragma acc data default(none)
+  ;
+
+  // OpenACC TODO: The following 'data' directives should diagnose, since they
+  // don't have at least one of the above clauses.
+
+  // expected-warning at +1{{OpenACC clause 'if' not yet implemented}}
+#pragma acc data if(Var)
+  ;
+
+  // expected-warning at +1{{OpenACC clause 'async' not yet implemented}}
+#pragma acc data async
+  ;
+
+  // expected-warning at +1{{OpenACC clause 'wait' not yet implemented}}
+#pragma acc data wait
+  ;
+
+  // expected-warning at +1{{OpenACC clause 'device_type' not yet implemented}}
+#pragma acc data device_type(*)
+  ;
+#pragma acc data
+  ;
+
+  // Enter Data
+  // expected-warning at +1{{OpenACC clause 'copyin' not yet implemented}}
+#pragma acc enter data copyin(Var)
+  // expected-warning at +1{{OpenACC clause 'create' not yet implemented}}
+#pragma acc enter data create(Var)
+  // expected-warning at +1{{OpenACC clause 'attach' not yet implemented}}
+#pragma acc enter data attach(Var)
+
+  // OpenACC TODO: The following 'enter data' directives should diagnose, since
+  // they don't have at least one of the above clauses.
+
+  // expected-warning at +1{{OpenACC clause 'if' not yet implemented}}
+#pragma acc enter data if(Var)
+  // expected-warning at +1{{OpenACC clause 'async' not yet implemented}}
+#pragma acc enter data async
+  // expected-warning at +1{{OpenACC clause 'wait' not yet implemented}}
+#pragma acc enter data wait
+#pragma acc enter data
+
+  // Exit Data
+  // expected-warning at +1{{OpenACC clause 'copyout' not yet implemented}}
+#pragma acc exit data copyout(Var)
+  // expected-warning at +1{{OpenACC clause 'delete' not yet implemented}}
+#pragma acc exit data delete(Var)
+  // expected-warning at +1{{OpenACC clause 'detach' not yet implemented}}
+#pragma acc exit data detach(Var)
+
+  // OpenACC TODO: The following 'exit data' directives should diagnose, since
+  // they don't have at least one of the above clauses.
+
+  // expected-warning at +1{{OpenACC clause 'if' not yet implemented}}
+#pragma acc exit data if(Var)
+  // expected-warning at +1{{OpenACC clause 'async' not yet implemented}}
+#pragma acc exit data async
+  // expected-warning at +1{{OpenACC clause 'wait' not yet implemented}}
+#pragma acc exit data wait
+  // expected-warning at +1{{OpenACC clause 'finalize' not yet implemented}}
+#pragma acc exit data finalize
+#pragma acc exit data
+
+  // Host Data
+  // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+#pragma acc host_data use_device(Var)
+  ;
+  // OpenACC TODO: The following 'host_data' directives should diagnose, since
+  // they don't have at least one of the above clauses.
+
+  // expected-warning at +1{{OpenACC clause 'if' not yet implemented}}
+#pragma acc host_data if(Var)
+  ;
+  // expected-warning at +1{{OpenACC clause 'if_present' not yet implemented}}
+#pragma acc host_data if_present
+  ;
+#pragma acc host_data
+  ;
+}
+
+void DataRules() {
+  int Var;
+  // OpenACC TODO: Only 'async' and 'wait' are permitted after a device_type, so
+  // the rest of these should diagnose.
+
+  // expected-warning at +2{{OpenACC clause 'device_type' not yet implemented}}
+  // expected-warning at +1{{OpenACC clause 'copy' not yet implemented}}
+#pragma acc data device_type(*) copy(Var)
+  ;
+  // expected-warning at +2{{OpenACC clause 'device_type' not yet implemented}}
+  // expected-warning at +1{{OpenACC clause 'copyin' not yet implemented}}
+#pragma acc data device_type(*) copyin(Var)
+  ;
+  // expected-warning at +2{{OpenACC clause 'device_type' not yet implemented}}
+  // expected-warning at +1{{OpenACC clause 'copyout' not yet implemented}}
+#pragma acc data device_type(*) copyout(Var)
+  ;
+  // expected-warning at +2{{OpenACC clause 'device_type' not yet implemented}}
+  // expected-warning at +1{{OpenACC clause 'create' not yet implemented}}
+#pragma acc data device_type(*) create(Var)
+  ;
+  // expected-warning at +2{{OpenACC clause 'device_type' not yet implemented}}
+  // expected-warning at +1{{OpenACC clause 'no_create' not yet implemented}}
+#pragma acc data device_type(*) no_create(Var)
+  ;
+  // expected-warning at +2{{OpenACC clause 'device_type' not yet implemented}}
+  // expected-warning at +1{{OpenACC clause 'present' not yet implemented}}
+#pragma acc data device_type(*) present(Var)
+  ;
+  // expected-warning at +2{{OpenACC clause 'device_type' not yet implemented}}
+  // expected-warning at +1{{OpenACC clause 'deviceptr' not yet implemented}}
+#pragma acc data device_type(*) deviceptr(Var)
+  ;
+  // expected-warning at +2{{OpenACC clause 'device_type' not yet implemented}}
+  // expected-warning at +1{{OpenACC clause 'attach' not yet implemented}}
+#pragma acc data device_type(*) attach(Var)
+  ;
+  // expected-warning at +2{{OpenACC clause 'device_type' not yet implemented}}
+  // expected-warning at +1{{OpenACC clause 'default' not yet implemented}}
+#pragma acc data device_type(*) default(none)
+  ;
+  // expected-warning at +2{{OpenACC clause 'device_type' not yet implemented}}
+  // expected-warning at +1{{OpenACC clause 'if' not yet implemented}}
+#pragma acc data device_type(*) if(Var)
+  ;
+  // expected-warning at +2{{OpenACC clause 'device_type' not yet implemented}}
+  // expected-warning at +1{{OpenACC clause 'async' not yet implemented}}
+#pragma acc data device_type(*) async
+  ;
+  // expected-warning at +2{{OpenACC clause 'device_type' not yet implemented}}
+  // expected-warning at +1{{OpenACC clause 'wait' not yet implemented}}
+#pragma acc data device_type(*) wait
+  ;
+}
+
+struct HasMembers {
+  int Member;
+
+  void HostDataError() {
+  // TODO OpenACC: The following 3 should error, as use_device's var only allows
+  // a variable or array, not an array index, or sub expression.
+
+  // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+#pragma acc host_data use_device(this)
+  ;
+  // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+#pragma acc host_data use_device(this->Member)
+  ;
+  // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+#pragma acc host_data use_device(Member)
+  ;
+  }
+};
+
+void HostDataRules() {
+  int Var, Var2;
+  // TODO OpenACC: The following line should diagnose, since only 1 'if' is
+  // allowed per directive on host_data.
+  // expected-warning at +2{{OpenACC clause 'if' not yet implemented}}
+  // expected-warning at +1{{OpenACC clause 'if' not yet implemented}}
+#pragma acc host_data if(Var) if (Var2)
+  ;
+
+  // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+#pragma acc host_data use_device(Var)
+  ;
+
+  int Array[5];
+  // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+#pragma acc host_data use_device(Array)
+  ;
+
+  // TODO OpenACC: The following 3 should error, as use_device's var only allows
+  // a variable or array, not an array index, or sub expression.
+
+  // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+#pragma acc host_data use_device(Array[1:1])
+  ;
+
+  // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+#pragma acc host_data use_device(Array[1])
+  ;
+  HasMembers HM;
+  // expected-warning at +1{{OpenACC clause 'use_device' not yet implemented}}
+#pragma acc host_data use_device(HM.Member)
+  ;
+
+}

diff  --git a/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp b/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp
index dc954e36d765da..b401dd891629aa 100644
--- a/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp
+++ b/clang/test/SemaOpenACC/loop-construct-collapse-clause.cpp
@@ -323,14 +323,15 @@ void no_other_directives() {
 #pragma acc loop collapse(2)
   for(unsigned i = 0; i < 5; ++i) {
     for(unsigned j = 0; j < 5; ++j) {
-#pragma acc data // expected-warning{{OpenACC construct 'data' not yet implemented}}
+#pragma acc data
+      ;
     }
   }
   // expected-note at +1{{active 'collapse' clause defined here}}
 #pragma acc loop collapse(2)
   for(unsigned i = 0; i < 5; ++i) {
     // expected-error at +1{{OpenACC 'data' construct cannot appear in intervening code of a 'loop' with a 'collapse' clause}}
-#pragma acc data // expected-warning{{OpenACC construct 'data' not yet implemented}}
+#pragma acc data
     for(unsigned j = 0; j < 5; ++j) {
     }
   }

diff  --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index 221f419861af10..d0fc69af7c847d 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2185,6 +2185,10 @@ class EnqueueVisitor : public ConstStmtVisitor<EnqueueVisitor, void>,
   void VisitOpenACCComputeConstruct(const OpenACCComputeConstruct *D);
   void VisitOpenACCLoopConstruct(const OpenACCLoopConstruct *D);
   void VisitOpenACCCombinedConstruct(const OpenACCCombinedConstruct *D);
+  void VisitOpenACCDataConstruct(const OpenACCDataConstruct *D);
+  void VisitOpenACCEnterDataConstruct(const OpenACCEnterDataConstruct *D);
+  void VisitOpenACCExitDataConstruct(const OpenACCExitDataConstruct *D);
+  void VisitOpenACCHostDataConstruct(const OpenACCHostDataConstruct *D);
   void VisitOMPExecutableDirective(const OMPExecutableDirective *D);
   void VisitOMPLoopBasedDirective(const OMPLoopBasedDirective *D);
   void VisitOMPLoopDirective(const OMPLoopDirective *D);
@@ -3587,6 +3591,29 @@ void EnqueueVisitor::VisitOpenACCCombinedConstruct(
   for (auto *Clause : C->clauses())
     EnqueueChildren(Clause);
 }
+void EnqueueVisitor::VisitOpenACCDataConstruct(const OpenACCDataConstruct *C) {
+  EnqueueChildren(C);
+  for (auto *Clause : C->clauses())
+    EnqueueChildren(Clause);
+}
+void EnqueueVisitor::VisitOpenACCEnterDataConstruct(
+    const OpenACCEnterDataConstruct *C) {
+  EnqueueChildren(C);
+  for (auto *Clause : C->clauses())
+    EnqueueChildren(Clause);
+}
+void EnqueueVisitor::VisitOpenACCExitDataConstruct(
+    const OpenACCExitDataConstruct *C) {
+  EnqueueChildren(C);
+  for (auto *Clause : C->clauses())
+    EnqueueChildren(Clause);
+}
+void EnqueueVisitor::VisitOpenACCHostDataConstruct(
+    const OpenACCHostDataConstruct *C) {
+  EnqueueChildren(C);
+  for (auto *Clause : C->clauses())
+    EnqueueChildren(Clause);
+}
 
 void EnqueueVisitor::VisitAnnotateAttr(const AnnotateAttr *A) {
   EnqueueChildren(A);
@@ -6342,6 +6369,14 @@ CXString clang_getCursorKindSpelling(enum CXCursorKind Kind) {
     return cxstring::createRef("OpenACCLoopConstruct");
   case CXCursor_OpenACCCombinedConstruct:
     return cxstring::createRef("OpenACCCombinedConstruct");
+  case CXCursor_OpenACCDataConstruct:
+    return cxstring::createRef("OpenACCDataConstruct");
+  case CXCursor_OpenACCEnterDataConstruct:
+    return cxstring::createRef("OpenACCEnterDataConstruct");
+  case CXCursor_OpenACCExitDataConstruct:
+    return cxstring::createRef("OpenACCExitDataConstruct");
+  case CXCursor_OpenACCHostDataConstruct:
+    return cxstring::createRef("OpenACCHostDataConstruct");
   }
 
   llvm_unreachable("Unhandled CXCursorKind");

diff  --git a/clang/tools/libclang/CXCursor.cpp b/clang/tools/libclang/CXCursor.cpp
index c8cf51d8061328..26935c45ce5f83 100644
--- a/clang/tools/libclang/CXCursor.cpp
+++ b/clang/tools/libclang/CXCursor.cpp
@@ -888,6 +888,18 @@ CXCursor cxcursor::MakeCXCursor(const Stmt *S, const Decl *Parent,
   case Stmt::OpenACCCombinedConstructClass:
     K = CXCursor_OpenACCCombinedConstruct;
     break;
+  case Stmt::OpenACCDataConstructClass:
+    K = CXCursor_OpenACCDataConstruct;
+    break;
+  case Stmt::OpenACCEnterDataConstructClass:
+    K = CXCursor_OpenACCEnterDataConstruct;
+    break;
+  case Stmt::OpenACCExitDataConstructClass:
+    K = CXCursor_OpenACCExitDataConstruct;
+    break;
+  case Stmt::OpenACCHostDataConstructClass:
+    K = CXCursor_OpenACCHostDataConstruct;
+    break;
   case Stmt::OMPTargetParallelGenericLoopDirectiveClass:
     K = CXCursor_OMPTargetParallelGenericLoopDirective;
     break;


        


More information about the cfe-commits mailing list