Index: include/clang-c/Index.h
===================================================================
--- include/clang-c/Index.h
+++ include/clang-c/Index.h
@@ -2274,7 +2274,15 @@
    */
   CXCursor_OMPDistributeDirective        = 260,
 
-  CXCursor_LastStmt                      = CXCursor_OMPDistributeDirective,
+  /** \brief OpenMP target enter data directive.
+   */
+  CXCursor_OMPTargetEnterDataDirective   = 261,
+
+  /** \brief OpenMP target exit data directive.
+   */
+  CXCursor_OMPTargetExitDataDirective    = 262,
+
+  CXCursor_LastStmt                      = CXCursor_OMPTargetExitDataDirective,
 
   /**
    * \brief Cursor that represents the translation unit itself.
Index: include/clang/AST/OpenMPClause.h
===================================================================
--- include/clang/AST/OpenMPClause.h
+++ include/clang/AST/OpenMPClause.h
@@ -2741,6 +2741,8 @@
   OpenMPMapClauseKind MapTypeModifier;
   /// \brief Map type for the 'map' clause.
   OpenMPMapClauseKind MapType;
+  /// \brief Is this an implicit map type or not.
+  bool MapTypeIsImplicit;
   /// \brief Location of the map type.
   SourceLocation MapLoc;
   /// \brief Colon location.
@@ -2777,11 +2779,13 @@
   /// \param N Number of the variables in the clause.
   ///
   explicit OMPMapClause(OpenMPMapClauseKind MapTypeModifier,
-                        OpenMPMapClauseKind MapType, SourceLocation MapLoc,
-                        SourceLocation StartLoc, SourceLocation LParenLoc,
-                        SourceLocation EndLoc, unsigned N)
+                        OpenMPMapClauseKind MapType, bool MapTypeIsImplicit,
+                        SourceLocation MapLoc, SourceLocation StartLoc,
+                        SourceLocation LParenLoc, SourceLocation EndLoc,
+                        unsigned N)
     : OMPVarListClause<OMPMapClause>(OMPC_map, StartLoc, LParenLoc, EndLoc, N),
-      MapTypeModifier(MapTypeModifier), MapType(MapType), MapLoc(MapLoc) {}
+      MapTypeModifier(MapTypeModifier), MapType(MapType),
+      MapTypeIsImplicit(MapTypeIsImplicit), MapLoc(MapLoc) {}
 
   /// \brief Build an empty clause.
   ///
@@ -2790,7 +2794,8 @@
   explicit OMPMapClause(unsigned N)
       : OMPVarListClause<OMPMapClause>(OMPC_map, SourceLocation(),
                                        SourceLocation(), SourceLocation(), N),
-        MapTypeModifier(OMPC_MAP_unknown), MapType(OMPC_MAP_unknown), MapLoc() {}
+        MapTypeModifier(OMPC_MAP_unknown), MapType(OMPC_MAP_unknown),
+        MapTypeIsImplicit(false), MapLoc() {}
 
 public:
   /// \brief Creates clause with a list of variables \a VL.
@@ -2807,7 +2812,8 @@
                               SourceLocation LParenLoc,
                               SourceLocation EndLoc, ArrayRef<Expr *> VL,
                               OpenMPMapClauseKind TypeModifier,
-                              OpenMPMapClauseKind Type, SourceLocation TypeLoc);
+                              OpenMPMapClauseKind Type, bool TypeIsImplicit,
+                              SourceLocation TypeLoc);
   /// \brief Creates an empty clause with the place for \a N variables.
   ///
   /// \param C AST context.
@@ -2818,6 +2824,11 @@
   /// \brief Fetches mapping kind for the clause.
   OpenMPMapClauseKind getMapType() const LLVM_READONLY { return MapType; }
 
+  /// \brief Is this an implicit map type?
+  bool isImplicitMapType() const LLVM_READONLY {
+    return MapTypeIsImplicit;
+  }
+
   /// \brief Fetches the map type modifier for the clause.
   OpenMPMapClauseKind getMapTypeModifier() const LLVM_READONLY {
     return MapTypeModifier;
Index: include/clang/AST/RecursiveASTVisitor.h
===================================================================
--- include/clang/AST/RecursiveASTVisitor.h
+++ include/clang/AST/RecursiveASTVisitor.h
@@ -2429,6 +2429,12 @@
 DEF_TRAVERSE_STMT(OMPTargetDataDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
+DEF_TRAVERSE_STMT(OMPTargetEnterDataDirective,
+                  { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
+DEF_TRAVERSE_STMT(OMPTargetExitDataDirective,
+                  { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
 DEF_TRAVERSE_STMT(OMPTeamsDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
Index: include/clang/AST/StmtOpenMP.h
===================================================================
--- include/clang/AST/StmtOpenMP.h
+++ include/clang/AST/StmtOpenMP.h
@@ -2039,6 +2039,124 @@
   }
 };
 
+/// \brief This represents '#pragma omp target enter data' directive.
+///
+/// \code
+/// #pragma omp target enter data device(0) if(a) map(b[:])
+/// \endcode
+/// In this example directive '#pragma omp target enter data' has clauses
+/// 'device' with the value '0', 'if' with condition 'a' and 'map' with array
+/// section 'b[:]'.
+///
+class OMPTargetEnterDataDirective : public OMPExecutableDirective {
+  friend class ASTStmtReader;
+  /// \brief Build directive with the given start and end location.
+  ///
+  /// \param StartLoc Starting location of the directive kind.
+  /// \param EndLoc Ending Location of the directive.
+  /// \param NumClauses The number of clauses.
+  ///
+  OMPTargetEnterDataDirective(SourceLocation StartLoc, SourceLocation EndLoc,
+                              unsigned NumClauses)
+      : OMPExecutableDirective(this, OMPTargetEnterDataDirectiveClass,
+                               OMPD_target_enter_data, StartLoc, EndLoc,
+                               NumClauses, /*NumChildren=*/0) {}
+
+  /// \brief Build an empty directive.
+  ///
+  /// \param NumClauses Number of clauses.
+  ///
+  explicit OMPTargetEnterDataDirective(unsigned NumClauses)
+      : OMPExecutableDirective(this, OMPTargetEnterDataDirectiveClass,
+                               OMPD_target_enter_data, SourceLocation(),
+                               SourceLocation(), NumClauses,
+                               /*NumChildren=*/0) {}
+
+public:
+  /// \brief Creates directive with a list of \a Clauses.
+  ///
+  /// \param C AST context.
+  /// \param StartLoc Starting location of the directive kind.
+  /// \param EndLoc Ending Location of the directive.
+  /// \param Clauses List of clauses.
+  /// \param AssociatedStmt Statement, associated with the directive.
+  ///
+  static OMPTargetEnterDataDirective *
+  Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+         ArrayRef<OMPClause *> Clauses);
+
+  /// \brief Creates an empty directive with the place for \a N clauses.
+  ///
+  /// \param C AST context.
+  /// \param N The number of clauses.
+  ///
+  static OMPTargetEnterDataDirective *CreateEmpty(const ASTContext &C,
+                                                  unsigned N, EmptyShell);
+
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == OMPTargetEnterDataDirectiveClass;
+  }
+};
+
+/// \brief This represents '#pragma omp target exit data' directive.
+///
+/// \code
+/// #pragma omp target exit data device(0) if(a) map(b[:])
+/// \endcode
+/// In this example directive '#pragma omp target exit data' has clauses
+/// 'device' with the value '0', 'if' with condition 'a' and 'map' with array
+/// section 'b[:]'.
+///
+class OMPTargetExitDataDirective : public OMPExecutableDirective {
+  friend class ASTStmtReader;
+  /// \brief Build directive with the given start and end location.
+  ///
+  /// \param StartLoc Starting location of the directive kind.
+  /// \param EndLoc Ending Location of the directive.
+  /// \param NumClauses The number of clauses.
+  ///
+  OMPTargetExitDataDirective(SourceLocation StartLoc, SourceLocation EndLoc,
+                             unsigned NumClauses)
+      : OMPExecutableDirective(this, OMPTargetExitDataDirectiveClass,
+                               OMPD_target_exit_data, StartLoc, EndLoc,
+                               NumClauses, /*NumChildren=*/0) {}
+
+  /// \brief Build an empty directive.
+  ///
+  /// \param NumClauses Number of clauses.
+  ///
+  explicit OMPTargetExitDataDirective(unsigned NumClauses)
+      : OMPExecutableDirective(this, OMPTargetExitDataDirectiveClass,
+                               OMPD_target_exit_data, SourceLocation(),
+                               SourceLocation(), NumClauses,
+                               /*NumChildren=*/0) {}
+
+public:
+  /// \brief Creates directive with a list of \a Clauses.
+  ///
+  /// \param C AST context.
+  /// \param StartLoc Starting location of the directive kind.
+  /// \param EndLoc Ending Location of the directive.
+  /// \param Clauses List of clauses.
+  /// \param AssociatedStmt Statement, associated with the directive.
+  ///
+  static OMPTargetExitDataDirective *
+  Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+         ArrayRef<OMPClause *> Clauses);
+
+  /// \brief Creates an empty directive with the place for \a N clauses.
+  ///
+  /// \param C AST context.
+  /// \param N The number of clauses.
+  ///
+  static OMPTargetExitDataDirective *CreateEmpty(const ASTContext &C,
+                                                 unsigned N, EmptyShell);
+
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == OMPTargetExitDataDirectiveClass;
+  }
+};
+
 /// \brief This represents '#pragma omp teams' directive.
 ///
 /// \code
Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -7927,6 +7927,10 @@
   "variable already marked as mapped in current construct">;
 def err_omp_not_mappable_type : Error<
   "type %0 is not mappable to target">;
+def err_omp_invalid_map_type_for_directive : Error<
+  "%select{map type '%1' is not allowed|map type must be specified}0 for '#pragma omp %2'">;
+def err_omp_no_map_for_directive : Error<
+  "expected at least one map clause for '#pragma omp %0'">;
 def note_omp_polymorphic_in_target : Note<
   "mappable type cannot be polymorphic">;
 def note_omp_static_member_in_target : Note<
Index: include/clang/Basic/OpenMPKinds.def
===================================================================
--- include/clang/Basic/OpenMPKinds.def
+++ include/clang/Basic/OpenMPKinds.def
@@ -60,6 +60,12 @@
 #ifndef OPENMP_TARGET_DATA_CLAUSE
 #  define OPENMP_TARGET_DATA_CLAUSE(Name)
 #endif
+#ifndef OPENMP_TARGET_ENTER_DATA_CLAUSE
+#define OPENMP_TARGET_ENTER_DATA_CLAUSE(Name)
+#endif
+#ifndef OPENMP_TARGET_EXIT_DATA_CLAUSE
+#define OPENMP_TARGET_EXIT_DATA_CLAUSE(Name)
+#endif
 #ifndef OPENMP_TEAMS_CLAUSE
 #  define OPENMP_TEAMS_CLAUSE(Name)
 #endif
@@ -125,6 +131,8 @@
 OPENMP_DIRECTIVE(teams)
 OPENMP_DIRECTIVE(cancel)
 OPENMP_DIRECTIVE_EXT(target_data, "target data")
+OPENMP_DIRECTIVE_EXT(target_enter_data, "target enter data")
+OPENMP_DIRECTIVE_EXT(target_exit_data, "target exit data")
 OPENMP_DIRECTIVE_EXT(parallel_for, "parallel for")
 OPENMP_DIRECTIVE_EXT(parallel_for_simd, "parallel for simd")
 OPENMP_DIRECTIVE_EXT(parallel_sections, "parallel sections")
@@ -349,6 +357,18 @@
 OPENMP_TARGET_DATA_CLAUSE(device)
 OPENMP_TARGET_DATA_CLAUSE(map)
 
+// Clauses allowed for OpenMP directive 'target enter data'.
+// TODO More clauses for 'target enter data' directive.
+OPENMP_TARGET_ENTER_DATA_CLAUSE(if)
+OPENMP_TARGET_ENTER_DATA_CLAUSE(device)
+OPENMP_TARGET_ENTER_DATA_CLAUSE(map)
+
+// Clauses allowed for OpenMP directive 'target exit data'.
+// TODO More clauses for 'target exit data' directive.
+OPENMP_TARGET_EXIT_DATA_CLAUSE(if)
+OPENMP_TARGET_EXIT_DATA_CLAUSE(device)
+OPENMP_TARGET_EXIT_DATA_CLAUSE(map)
+
 // Clauses allowed for OpenMP directive 'teams'.
 // TODO More clauses for 'teams' directive.
 OPENMP_TEAMS_CLAUSE(default)
@@ -443,6 +463,8 @@
 #undef OPENMP_ATOMIC_CLAUSE
 #undef OPENMP_TARGET_CLAUSE
 #undef OPENMP_TARGET_DATA_CLAUSE
+#undef OPENMP_TARGET_ENTER_DATA_CLAUSE
+#undef OPENMP_TARGET_EXIT_DATA_CLAUSE
 #undef OPENMP_TEAMS_CLAUSE
 #undef OPENMP_SIMD_CLAUSE
 #undef OPENMP_FOR_CLAUSE
Index: include/clang/Basic/StmtNodes.td
===================================================================
--- include/clang/Basic/StmtNodes.td
+++ include/clang/Basic/StmtNodes.td
@@ -216,6 +216,8 @@
 def OMPAtomicDirective : DStmt<OMPExecutableDirective>;
 def OMPTargetDirective : DStmt<OMPExecutableDirective>;
 def OMPTargetDataDirective : DStmt<OMPExecutableDirective>;
+def OMPTargetEnterDataDirective : DStmt<OMPExecutableDirective>;
+def OMPTargetExitDataDirective : DStmt<OMPExecutableDirective>;
 def OMPTeamsDirective : DStmt<OMPExecutableDirective>;
 def OMPCancellationPointDirective : DStmt<OMPExecutableDirective>;
 def OMPCancelDirective : DStmt<OMPExecutableDirective>;
Index: include/clang/Sema/Sema.h
===================================================================
--- include/clang/Sema/Sema.h
+++ include/clang/Sema/Sema.h
@@ -7761,6 +7761,8 @@
   ExprResult
   VerifyPositiveIntegerConstantInClause(Expr *Op, OpenMPClauseKind CKind,
                                         bool StrictlyPositive = true);
+  /// \brief Check for existence of a map clause
+  bool HasMapClause(ArrayRef<OMPClause *> Clauses);
 
 public:
   /// \brief Return true if the provided declaration \a VD should be captured by
@@ -7938,6 +7940,16 @@
   StmtResult ActOnOpenMPTargetDataDirective(ArrayRef<OMPClause *> Clauses,
                                             Stmt *AStmt, SourceLocation StartLoc,
                                             SourceLocation EndLoc);
+  /// \brief Called on well-formed '\#pragma omp target enter data' after
+  /// parsing of the associated statement.
+  StmtResult ActOnOpenMPTargetEnterDataDirective(ArrayRef<OMPClause *> Clauses,
+                                                 SourceLocation StartLoc,
+                                                 SourceLocation EndLoc);
+  /// \brief Called on well-formed '\#pragma omp target exit data' after
+  /// parsing of the associated statement.
+  StmtResult ActOnOpenMPTargetExitDataDirective(ArrayRef<OMPClause *> Clauses,
+                                                SourceLocation StartLoc,
+                                                SourceLocation EndLoc);
   /// \brief Called on well-formed '\#pragma omp teams' after parsing of the
   /// associated statement.
   StmtResult ActOnOpenMPTeamsDirective(ArrayRef<OMPClause *> Clauses,
@@ -8099,7 +8111,8 @@
       CXXScopeSpec &ReductionIdScopeSpec,
       const DeclarationNameInfo &ReductionId, OpenMPDependClauseKind DepKind,
       OpenMPLinearClauseKind LinKind, OpenMPMapClauseKind MapTypeModifier,
-      OpenMPMapClauseKind MapType, SourceLocation DepLinMapLoc);
+      OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
+      SourceLocation DepLinMapLoc);
   /// \brief Called on well-formed 'private' clause.
   OMPClause *ActOnOpenMPPrivateClause(ArrayRef<Expr *> VarList,
                                       SourceLocation StartLoc,
@@ -8168,8 +8181,9 @@
   /// \brief Called on well-formed 'map' clause.
   OMPClause *ActOnOpenMPMapClause(
       OpenMPMapClauseKind MapTypeModifier, OpenMPMapClauseKind MapType,
-      SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
-      SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc);
+      bool IsMapTypeImplicit, SourceLocation MapLoc, SourceLocation ColonLoc,
+      ArrayRef<Expr *> VarList, SourceLocation StartLoc, SourceLocation LParenLoc,
+      SourceLocation EndLoc);
   /// \brief Called on well-formed 'num_teams' clause.
   OMPClause *ActOnOpenMPNumTeamsClause(Expr *NumTeams, SourceLocation StartLoc,
                                        SourceLocation LParenLoc,
Index: include/clang/Serialization/ASTBitCodes.h
===================================================================
--- include/clang/Serialization/ASTBitCodes.h
+++ include/clang/Serialization/ASTBitCodes.h
@@ -1443,6 +1443,8 @@
       STMT_OMP_ATOMIC_DIRECTIVE,
       STMT_OMP_TARGET_DIRECTIVE,
       STMT_OMP_TARGET_DATA_DIRECTIVE,
+      STMT_OMP_TARGET_ENTER_DATA_DIRECTIVE,
+      STMT_OMP_TARGET_EXIT_DATA_DIRECTIVE,
       STMT_OMP_TEAMS_DIRECTIVE,
       STMT_OMP_TASKGROUP_DIRECTIVE,
       STMT_OMP_CANCELLATION_POINT_DIRECTIVE,
Index: lib/AST/OpenMPClause.cpp
===================================================================
--- lib/AST/OpenMPClause.cpp
+++ lib/AST/OpenMPClause.cpp
@@ -400,10 +400,12 @@
                                    SourceLocation EndLoc, ArrayRef<Expr *> VL,
                                    OpenMPMapClauseKind TypeModifier,
                                    OpenMPMapClauseKind Type,
+                                   bool TypeIsImplicit,
                                    SourceLocation TypeLoc) {
   void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(VL.size()));
   OMPMapClause *Clause = new (Mem) OMPMapClause(
-      TypeModifier, Type, TypeLoc, StartLoc, LParenLoc, EndLoc, VL.size());
+      TypeModifier, Type, TypeIsImplicit, TypeLoc, StartLoc, LParenLoc, EndLoc,
+      VL.size());
   Clause->setVarRefs(VL);
   Clause->setMapTypeModifier(TypeModifier);
   Clause->setMapType(Type);
Index: lib/AST/StmtOpenMP.cpp
===================================================================
--- lib/AST/StmtOpenMP.cpp
+++ lib/AST/StmtOpenMP.cpp
@@ -718,6 +718,50 @@
   return new (Mem) OMPTargetDataDirective(N);
 }
 
+OMPTargetEnterDataDirective *OMPTargetEnterDataDirective::Create(
+    const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+    ArrayRef<OMPClause *> Clauses) {
+  void *Mem =
+      C.Allocate(llvm::RoundUpToAlignment(sizeof(OMPTargetEnterDataDirective),
+                                          llvm::alignOf<OMPClause *>()) +
+                 sizeof(OMPClause *) * Clauses.size());
+  OMPTargetEnterDataDirective *Dir =
+      new (Mem) OMPTargetEnterDataDirective(StartLoc, EndLoc, Clauses.size());
+  Dir->setClauses(Clauses);
+  return Dir;
+}
+
+OMPTargetEnterDataDirective *OMPTargetEnterDataDirective::CreateEmpty(
+    const ASTContext &C, unsigned N, EmptyShell) {
+  void *Mem =
+      C.Allocate(llvm::RoundUpToAlignment(sizeof(OMPTargetEnterDataDirective),
+                                          llvm::alignOf<OMPClause *>()) +
+                 sizeof(OMPClause *) * N);
+  return new (Mem) OMPTargetEnterDataDirective(N);
+}
+
+OMPTargetExitDataDirective *OMPTargetExitDataDirective::Create(
+    const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+    ArrayRef<OMPClause *> Clauses) {
+  void *Mem =
+      C.Allocate(llvm::RoundUpToAlignment(sizeof(OMPTargetExitDataDirective),
+                                          llvm::alignOf<OMPClause *>()) +
+                 sizeof(OMPClause *) * Clauses.size());
+  OMPTargetExitDataDirective *Dir =
+      new (Mem) OMPTargetExitDataDirective(StartLoc, EndLoc, Clauses.size());
+  Dir->setClauses(Clauses);
+  return Dir;
+}
+
+OMPTargetExitDataDirective *OMPTargetExitDataDirective::CreateEmpty(
+    const ASTContext &C, unsigned N, EmptyShell) {
+  void *Mem =
+      C.Allocate(llvm::RoundUpToAlignment(sizeof(OMPTargetExitDataDirective),
+                                          llvm::alignOf<OMPClause *>()) +
+                 sizeof(OMPClause *) * N);
+  return new (Mem) OMPTargetExitDataDirective(N);
+}
+
 OMPTeamsDirective *OMPTeamsDirective::Create(const ASTContext &C,
                                              SourceLocation StartLoc,
                                              SourceLocation EndLoc,
Index: lib/AST/StmtPrinter.cpp
===================================================================
--- lib/AST/StmtPrinter.cpp
+++ lib/AST/StmtPrinter.cpp
@@ -1051,6 +1051,18 @@
   PrintOMPExecutableDirective(Node);
 }
 
+void StmtPrinter::VisitOMPTargetEnterDataDirective(
+    OMPTargetEnterDataDirective *Node) {
+  Indent() << "#pragma omp target enter data ";
+  PrintOMPExecutableDirective(Node);
+}
+
+void StmtPrinter::VisitOMPTargetExitDataDirective(
+    OMPTargetExitDataDirective *Node) {
+  Indent() << "#pragma omp target exit data ";
+  PrintOMPExecutableDirective(Node);
+}
+
 void StmtPrinter::VisitOMPTeamsDirective(OMPTeamsDirective *Node) {
   Indent() << "#pragma omp teams ";
   PrintOMPExecutableDirective(Node);
Index: lib/AST/StmtProfile.cpp
===================================================================
--- lib/AST/StmtProfile.cpp
+++ lib/AST/StmtProfile.cpp
@@ -584,6 +584,16 @@
   VisitOMPExecutableDirective(S);
 }
 
+void StmtProfiler::VisitOMPTargetEnterDataDirective(
+    const OMPTargetEnterDataDirective *S) {
+  VisitOMPExecutableDirective(S);
+}
+
+void StmtProfiler::VisitOMPTargetExitDataDirective(
+    const OMPTargetExitDataDirective *S) {
+  VisitOMPExecutableDirective(S);
+}
+
 void StmtProfiler::VisitOMPTeamsDirective(const OMPTeamsDirective *S) {
   VisitOMPExecutableDirective(S);
 }
Index: lib/Basic/OpenMPKinds.cpp
===================================================================
--- lib/Basic/OpenMPKinds.cpp
+++ lib/Basic/OpenMPKinds.cpp
@@ -398,6 +398,26 @@
       break;
     }
     break;
+  case OMPD_target_enter_data:
+    switch (CKind) {
+#define OPENMP_TARGET_ENTER_DATA_CLAUSE(Name)                                  \
+  case OMPC_##Name:                                                            \
+    return true;
+#include "clang/Basic/OpenMPKinds.def"
+    default:
+      break;
+    }
+    break;
+  case OMPD_target_exit_data:
+    switch (CKind) {
+#define OPENMP_TARGET_EXIT_DATA_CLAUSE(Name)                                   \
+  case OMPC_##Name:                                                            \
+    return true;
+#include "clang/Basic/OpenMPKinds.def"
+    default:
+      break;
+    }
+    break;
   case OMPD_teams:
     switch (CKind) {
 #define OPENMP_TEAMS_CLAUSE(Name)                                              \
Index: lib/CodeGen/CGStmt.cpp
===================================================================
--- lib/CodeGen/CGStmt.cpp
+++ lib/CodeGen/CGStmt.cpp
@@ -256,6 +256,12 @@
   case Stmt::OMPTargetDataDirectiveClass:
     EmitOMPTargetDataDirective(cast<OMPTargetDataDirective>(*S));
     break;
+  case Stmt::OMPTargetEnterDataDirectiveClass:
+    EmitOMPTargetEnterDataDirective(cast<OMPTargetEnterDataDirective>(*S));
+    break;
+  case Stmt::OMPTargetExitDataDirectiveClass:
+    EmitOMPTargetExitDataDirective(cast<OMPTargetExitDataDirective>(*S));
+    break;
   case Stmt::OMPTaskLoopDirectiveClass:
     EmitOMPTaskLoopDirective(cast<OMPTaskLoopDirective>(*S));
     break;
Index: lib/CodeGen/CGStmtOpenMP.cpp
===================================================================
--- lib/CodeGen/CGStmtOpenMP.cpp
+++ lib/CodeGen/CGStmtOpenMP.cpp
@@ -2660,6 +2660,16 @@
       [&CS](CodeGenFunction &CGF) { CGF.EmitStmt(CS->getCapturedStmt()); });
 }
 
+void CodeGenFunction::EmitOMPTargetEnterDataDirective(
+    const OMPTargetEnterDataDirective &S) {
+  llvm_unreachable("CodeGen for 'omp target enter data' is not supported.");
+}
+
+void CodeGenFunction::EmitOMPTargetExitDataDirective(
+    const OMPTargetExitDataDirective &S) {
+  llvm_unreachable("CodeGen for 'omp target exit data' is not supported.");
+}
+
 void CodeGenFunction::EmitOMPTaskLoopDirective(const OMPTaskLoopDirective &S) {
   // emit the code inside the construct for now
   auto CS = cast<CapturedStmt>(S.getAssociatedStmt());
Index: lib/CodeGen/CodeGenFunction.h
===================================================================
--- lib/CodeGen/CodeGenFunction.h
+++ lib/CodeGen/CodeGenFunction.h
@@ -2335,6 +2335,8 @@
   void EmitOMPAtomicDirective(const OMPAtomicDirective &S);
   void EmitOMPTargetDirective(const OMPTargetDirective &S);
   void EmitOMPTargetDataDirective(const OMPTargetDataDirective &S);
+  void EmitOMPTargetEnterDataDirective(const OMPTargetEnterDataDirective &S);
+  void EmitOMPTargetExitDataDirective(const OMPTargetExitDataDirective &S);
   void EmitOMPTeamsDirective(const OMPTeamsDirective &S);
   void
   EmitOMPCancellationPointDirective(const OMPCancellationPointDirective &S);
Index: lib/Parse/ParseOpenMP.cpp
===================================================================
--- lib/Parse/ParseOpenMP.cpp
+++ lib/Parse/ParseOpenMP.cpp
@@ -34,6 +34,12 @@
       {OMPD_unknown /*cancellation*/, OMPD_unknown /*point*/,
        OMPD_cancellation_point},
       {OMPD_target, OMPD_unknown /*data*/, OMPD_target_data},
+      {OMPD_target, OMPD_unknown /*enter/exit*/,
+       OMPD_unknown /*target enter/exit*/},
+      {OMPD_unknown /*target enter*/, OMPD_unknown /*data*/,
+       OMPD_target_enter_data},
+      {OMPD_unknown /*target exit*/, OMPD_unknown /*data*/,
+       OMPD_target_exit_data},
       {OMPD_for, OMPD_simd, OMPD_for_simd},
       {OMPD_parallel, OMPD_for, OMPD_parallel_for},
       {OMPD_parallel_for, OMPD_simd, OMPD_parallel_for_simd},
@@ -49,8 +55,12 @@
   for (unsigned i = 0; i < llvm::array_lengthof(F); ++i) {
     if (!Tok.isAnnotation() && DKind == OMPD_unknown) {
       TokenMatched =
-          (i == 0) &&
-          !P.getPreprocessor().getSpelling(Tok).compare("cancellation");
+          ((i == 0) &&
+           !P.getPreprocessor().getSpelling(Tok).compare("cancellation")) ||
+          ((i == 3) &&
+           !P.getPreprocessor().getSpelling(Tok).compare("enter")) ||
+          ((i == 4) &&
+           !P.getPreprocessor().getSpelling(Tok).compare("exit"));
     } else {
       TokenMatched = DKind == F[i][0] && DKind != OMPD_unknown;
     }
@@ -67,7 +77,11 @@
         TokenMatched =
             ((i == 0) &&
              !P.getPreprocessor().getSpelling(Tok).compare("point")) ||
-            ((i == 1) && !P.getPreprocessor().getSpelling(Tok).compare("data"));
+            ((i == 1 || i == 3 || i == 4) &&
+             !P.getPreprocessor().getSpelling(Tok).compare("data")) ||
+            ((i == 2) &&
+             (!P.getPreprocessor().getSpelling(Tok).compare("enter") ||
+              !P.getPreprocessor().getSpelling(Tok).compare("exit")));
       } else {
         TokenMatched = SDKind == F[i][1] && SDKind != OMPD_unknown;
       }
@@ -138,6 +152,8 @@
   case OMPD_cancellation_point:
   case OMPD_cancel:
   case OMPD_target_data:
+  case OMPD_target_enter_data:
+  case OMPD_target_exit_data:
   case OMPD_taskloop:
   case OMPD_taskloop_simd:
   case OMPD_distribute:
@@ -162,7 +178,7 @@
 ///         'barrier' | 'taskwait' | 'flush' | 'ordered' | 'atomic' |
 ///         'for simd' | 'parallel for simd' | 'target' | 'target data' |
 ///         'taskgroup' | 'teams' | 'taskloop' | 'taskloop simd' {clause} |
-///         'distribute'
+///         'distribute' | 'target enter data' | 'target exit data'
 ///         annot_pragma_openmp_end
 ///
 StmtResult
@@ -213,6 +229,8 @@
   case OMPD_taskwait:
   case OMPD_cancellation_point:
   case OMPD_cancel:
+  case OMPD_target_enter_data:
+  case OMPD_target_exit_data:
     if (!StandAloneAllowed) {
       Diag(Tok, diag::err_omp_immediate_directive)
           << getOpenMPDirectiveName(DKind) << 0;
@@ -866,6 +884,7 @@
   OpenMPLinearClauseKind LinearModifier = OMPC_LINEAR_val;
   OpenMPMapClauseKind MapType = OMPC_MAP_unknown;
   OpenMPMapClauseKind MapTypeModifier = OMPC_MAP_unknown;
+  bool MapTypeIsImplicit = false;
   bool MapTypeModifierSpecified = false;
   bool UnexpectedId = false;
   SourceLocation DepLinMapLoc;
@@ -916,7 +935,8 @@
             Kind, llvm::None, /*TailExpr=*/nullptr, Loc, LOpen,
             /*ColonLoc=*/SourceLocation(), Tok.getLocation(),
             ReductionIdScopeSpec, DeclarationNameInfo(), DepKind,
-            LinearModifier, MapTypeModifier, MapType, DepLinMapLoc);
+            LinearModifier, MapTypeModifier, MapType, MapTypeIsImplicit,
+            DepLinMapLoc);
       }
     }
     if (Tok.is(tok::colon)) {
@@ -980,9 +1000,11 @@
           ConsumeToken();
         } else {
           MapType = OMPC_MAP_tofrom;
+          MapTypeIsImplicit = true;
         }
       } else {
         MapType = OMPC_MAP_tofrom;
+        MapTypeIsImplicit = true;
       }
     } else {
       UnexpectedId = true;
@@ -1063,6 +1085,7 @@
       ReductionIdScopeSpec,
       ReductionId.isValid() ? Actions.GetNameFromUnqualifiedId(ReductionId)
                             : DeclarationNameInfo(),
-      DepKind, LinearModifier, MapTypeModifier, MapType, DepLinMapLoc);
+      DepKind, LinearModifier, MapTypeModifier, MapType, MapTypeIsImplicit,
+      DepLinMapLoc);
 }
 
Index: lib/Sema/SemaOpenMP.cpp
===================================================================
--- lib/Sema/SemaOpenMP.cpp
+++ lib/Sema/SemaOpenMP.cpp
@@ -1610,6 +1610,8 @@
   case OMPD_cancellation_point:
   case OMPD_cancel:
   case OMPD_flush:
+  case OMPD_target_enter_data:
+  case OMPD_target_exit_data:
     llvm_unreachable("OpenMP Directive is not allowed");
   case OMPD_unknown:
     llvm_unreachable("Unknown OpenMP directive");
@@ -2697,6 +2699,16 @@
                                          EndLoc);
     AllowedNameModifiers.push_back(OMPD_target_data);
     break;
+  case OMPD_target_enter_data:
+    Res = ActOnOpenMPTargetEnterDataDirective(ClausesWithImplicit, StartLoc,
+                                              EndLoc);
+    AllowedNameModifiers.push_back(OMPD_target_enter_data);
+    break;
+  case OMPD_target_exit_data:
+    Res = ActOnOpenMPTargetExitDataDirective(ClausesWithImplicit, StartLoc,
+                                             EndLoc);
+    AllowedNameModifiers.push_back(OMPD_target_exit_data);
+    break;
   case OMPD_taskloop:
     Res = ActOnOpenMPTaskLoopDirective(ClausesWithImplicit, AStmt, StartLoc,
                                        EndLoc, VarsWithInheritedDSA);
@@ -5448,6 +5460,34 @@
                                         AStmt);
 }
 
+StmtResult Sema::ActOnOpenMPTargetEnterDataDirective(
+    ArrayRef<OMPClause *> Clauses, SourceLocation StartLoc,
+    SourceLocation EndLoc) {
+  // OpenMP [2.10.2, Restrictions, p. 99]
+  // At least one map clause must appear on the directive.
+  if (!HasMapClause(Clauses)) {
+    Diag(StartLoc, diag::err_omp_no_map_for_directive) <<
+        getOpenMPDirectiveName(OMPD_target_enter_data);
+    return StmtError();
+  }
+
+  return OMPTargetEnterDataDirective::Create(Context, StartLoc, EndLoc, Clauses);
+}
+
+StmtResult Sema::ActOnOpenMPTargetExitDataDirective(
+    ArrayRef<OMPClause *> Clauses, SourceLocation StartLoc,
+    SourceLocation EndLoc) {
+  // OpenMP [2.10.3, Restrictions, p. 102]
+  // At least one map clause must appear on the directive.
+  if (!HasMapClause(Clauses)) {
+    Diag(StartLoc, diag::err_omp_no_map_for_directive) <<
+        getOpenMPDirectiveName(OMPD_target_exit_data);
+    return StmtError();
+  }
+
+  return OMPTargetExitDataDirective::Create(Context, StartLoc, EndLoc, Clauses);
+}
+
 StmtResult Sema::ActOnOpenMPTeamsDirective(ArrayRef<OMPClause *> Clauses,
                                            Stmt *AStmt, SourceLocation StartLoc,
                                            SourceLocation EndLoc) {
@@ -5863,6 +5903,19 @@
   return ICE;
 }
 
+bool Sema::HasMapClause(ArrayRef<OMPClause *> Clauses) {
+  for (ArrayRef<OMPClause *>::iterator I = Clauses.begin(), E = Clauses.end();
+       I != E; ++I) {
+    if (*I) {
+      OMPClause *Clause = *I;
+      if (Clause->getClauseKind() == OMPC_map)
+        return true;
+    }
+  }
+
+  return false;
+}
+
 OMPClause *Sema::ActOnOpenMPSafelenClause(Expr *Len, SourceLocation StartLoc,
                                           SourceLocation LParenLoc,
                                           SourceLocation EndLoc) {
@@ -6371,7 +6424,8 @@
     SourceLocation EndLoc, CXXScopeSpec &ReductionIdScopeSpec,
     const DeclarationNameInfo &ReductionId, OpenMPDependClauseKind DepKind,
     OpenMPLinearClauseKind LinKind, OpenMPMapClauseKind MapTypeModifier, 
-    OpenMPMapClauseKind MapType, SourceLocation DepLinMapLoc) {
+    OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
+    SourceLocation DepLinMapLoc) {
   OMPClause *Res = nullptr;
   switch (Kind) {
   case OMPC_private:
@@ -6412,8 +6466,9 @@
                                   StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_map:
-    Res = ActOnOpenMPMapClause(MapTypeModifier, MapType, DepLinMapLoc, ColonLoc,
-                               VarList, StartLoc, LParenLoc, EndLoc);
+    Res = ActOnOpenMPMapClause(MapTypeModifier, MapType, IsMapTypeImplicit,
+                               DepLinMapLoc, ColonLoc, VarList, StartLoc,
+                               LParenLoc, EndLoc);
     break;
   case OMPC_if:
   case OMPC_final:
@@ -8332,8 +8387,9 @@
 
 OMPClause *Sema::ActOnOpenMPMapClause(
     OpenMPMapClauseKind MapTypeModifier, OpenMPMapClauseKind MapType,
-    SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
-    SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc) {
+    bool IsMapTypeImplicit, SourceLocation MapLoc, SourceLocation ColonLoc,
+    ArrayRef<Expr *> VarList, SourceLocation StartLoc, SourceLocation LParenLoc,
+    SourceLocation EndLoc) {
   SmallVector<Expr *, 4> Vars;
 
   for (auto &RE : VarList) {
@@ -8440,6 +8496,37 @@
                            DSAStack, Type))
       continue;
 
+    // target enter data
+    // OpenMP [2.10.2, Restrictions, p. 99]
+    // A map-type must be specified in all map clauses and must be either
+    // to or alloc.
+    OpenMPDirectiveKind DKind = DSAStack->getCurrentDirective();
+    if (DKind == OMPD_target_enter_data &&
+        !(MapType == OMPC_MAP_to || MapType == OMPC_MAP_alloc)) {
+      Diag(StartLoc, diag::err_omp_invalid_map_type_for_directive) <<
+          static_cast<unsigned>(IsMapTypeImplicit) <<
+          getOpenMPSimpleClauseTypeName(OMPC_map, MapType) <<
+          getOpenMPDirectiveName(DKind);
+      // Proceed to add the variable in a map clause anyway, to prevent
+      // further spurious messages
+    }
+
+    // target exit_data
+    // OpenMP [2.10.3, Restrictions, p. 102]
+    // A map-type must be specified in all map clauses and must be either
+    // from, release, or delete.
+    DKind = DSAStack->getCurrentDirective();
+    if (DKind == OMPD_target_exit_data &&
+        !(MapType == OMPC_MAP_from || MapType == OMPC_MAP_release ||
+          MapType == OMPC_MAP_delete)) {
+      Diag(StartLoc, diag::err_omp_invalid_map_type_for_directive) <<
+          static_cast<unsigned>(IsMapTypeImplicit) <<
+          getOpenMPSimpleClauseTypeName(OMPC_map, MapType) <<
+          getOpenMPDirectiveName(DKind);
+      // Proceed to add the variable in a map clause anyway, to prevent
+      // further spurious messages
+    }
+
     Vars.push_back(RE);
     MI.RefExpr = RE;
     DSAStack->addMapInfoForVar(VD, MI);
@@ -8448,7 +8535,8 @@
     return nullptr;
 
   return OMPMapClause::Create(Context, StartLoc, LParenLoc, EndLoc, Vars,
-                              MapTypeModifier, MapType, MapLoc);
+                              MapTypeModifier, MapType, IsMapTypeImplicit,
+                              MapLoc);
 }
 
 OMPClause *Sema::ActOnOpenMPNumTeamsClause(Expr *NumTeams, 
Index: lib/Sema/TreeTransform.h
===================================================================
--- lib/Sema/TreeTransform.h
+++ lib/Sema/TreeTransform.h
@@ -1657,10 +1657,11 @@
   /// Subclasses may override this routine to provide different behavior.
   OMPClause *RebuildOMPMapClause(
       OpenMPMapClauseKind MapTypeModifier, OpenMPMapClauseKind MapType,
-      SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
-      SourceLocation StartLoc, SourceLocation LParenLoc,
-      SourceLocation EndLoc) {
-    return getSema().ActOnOpenMPMapClause(MapTypeModifier, MapType, MapLoc,
+      bool IsMapTypeImplicit, SourceLocation MapLoc, SourceLocation ColonLoc,
+      ArrayRef<Expr *> VarList, SourceLocation StartLoc,
+      SourceLocation LParenLoc, SourceLocation EndLoc) {
+    return getSema().ActOnOpenMPMapClause(MapTypeModifier, MapType,
+                                          IsMapTypeImplicit, MapLoc,
                                           ColonLoc, VarList,StartLoc,
                                           LParenLoc, EndLoc);
   }
@@ -7353,6 +7354,28 @@
 }
 
 template <typename Derived>
+StmtResult TreeTransform<Derived>::TransformOMPTargetEnterDataDirective(
+    OMPTargetEnterDataDirective *D) {
+  DeclarationNameInfo DirName;
+  getDerived().getSema().StartOpenMPDSABlock(OMPD_target_enter_data, DirName,
+                                             nullptr, D->getLocStart());
+  StmtResult Res = getDerived().TransformOMPExecutableDirective(D);
+  getDerived().getSema().EndOpenMPDSABlock(Res.get());
+  return Res;
+}
+
+template <typename Derived>
+StmtResult TreeTransform<Derived>::TransformOMPTargetExitDataDirective(
+    OMPTargetExitDataDirective *D) {
+  DeclarationNameInfo DirName;
+  getDerived().getSema().StartOpenMPDSABlock(OMPD_target_exit_data, DirName,
+                                             nullptr, D->getLocStart());
+  StmtResult Res = getDerived().TransformOMPExecutableDirective(D);
+  getDerived().getSema().EndOpenMPDSABlock(Res.get());
+  return Res;
+}
+
+template <typename Derived>
 StmtResult
 TreeTransform<Derived>::TransformOMPTeamsDirective(OMPTeamsDirective *D) {
   DeclarationNameInfo DirName;
@@ -7800,9 +7823,9 @@
     Vars.push_back(EVar.get());
   }
   return getDerived().RebuildOMPMapClause(
-      C->getMapTypeModifier(), C->getMapType(), C->getMapLoc(),
-      C->getColonLoc(), Vars, C->getLocStart(), C->getLParenLoc(),
-      C->getLocEnd());
+      C->getMapTypeModifier(), C->getMapType(), C->isImplicitMapType(),
+      C->getMapLoc(), C->getColonLoc(), Vars, C->getLocStart(),
+      C->getLParenLoc(), C->getLocEnd());
 }
 
 template <typename Derived>
Index: lib/Serialization/ASTReaderStmt.cpp
===================================================================
--- lib/Serialization/ASTReaderStmt.cpp
+++ lib/Serialization/ASTReaderStmt.cpp
@@ -2445,6 +2445,20 @@
   VisitOMPExecutableDirective(D);
 }
 
+void ASTStmtReader::VisitOMPTargetEnterDataDirective(
+    OMPTargetEnterDataDirective *D) {
+  VisitStmt(D);
+  ++Idx;
+  VisitOMPExecutableDirective(D);
+}
+
+void ASTStmtReader::VisitOMPTargetExitDataDirective(
+    OMPTargetExitDataDirective *D) {
+  VisitStmt(D);
+  ++Idx;
+  VisitOMPExecutableDirective(D);
+}
+
 void ASTStmtReader::VisitOMPTeamsDirective(OMPTeamsDirective *D) {
   VisitStmt(D);
   // The NumClauses field was read in ReadStmtFromStream.
@@ -3088,6 +3102,16 @@
           Context, Record[ASTStmtReader::NumStmtFields], Empty);
       break;
 
+    case STMT_OMP_TARGET_ENTER_DATA_DIRECTIVE:
+      S = OMPTargetEnterDataDirective::CreateEmpty(
+          Context, Record[ASTStmtReader::NumStmtFields], Empty);
+      break;
+
+    case STMT_OMP_TARGET_EXIT_DATA_DIRECTIVE:
+      S = OMPTargetExitDataDirective::CreateEmpty(
+          Context, Record[ASTStmtReader::NumStmtFields], Empty);
+      break;
+
     case STMT_OMP_TEAMS_DIRECTIVE:
       S = OMPTeamsDirective::CreateEmpty(
           Context, Record[ASTStmtReader::NumStmtFields], Empty);
Index: lib/Serialization/ASTWriterStmt.cpp
===================================================================
--- lib/Serialization/ASTWriterStmt.cpp
+++ lib/Serialization/ASTWriterStmt.cpp
@@ -2212,6 +2212,22 @@
   Code = serialization::STMT_OMP_TARGET_DATA_DIRECTIVE;
 }
 
+void ASTStmtWriter::VisitOMPTargetEnterDataDirective(
+    OMPTargetEnterDataDirective *D) {
+  VisitStmt(D);
+  Record.push_back(D->getNumClauses());
+  VisitOMPExecutableDirective(D);
+  Code = serialization::STMT_OMP_TARGET_ENTER_DATA_DIRECTIVE;
+}
+
+void ASTStmtWriter::VisitOMPTargetExitDataDirective(
+    OMPTargetExitDataDirective *D) {
+  VisitStmt(D);
+  Record.push_back(D->getNumClauses());
+  VisitOMPExecutableDirective(D);
+  Code = serialization::STMT_OMP_TARGET_EXIT_DATA_DIRECTIVE;
+}
+
 void ASTStmtWriter::VisitOMPTaskyieldDirective(OMPTaskyieldDirective *D) {
   VisitStmt(D);
   VisitOMPExecutableDirective(D);
Index: lib/StaticAnalyzer/Core/ExprEngine.cpp
===================================================================
--- lib/StaticAnalyzer/Core/ExprEngine.cpp
+++ lib/StaticAnalyzer/Core/ExprEngine.cpp
@@ -830,6 +830,8 @@
     case Stmt::OMPAtomicDirectiveClass:
     case Stmt::OMPTargetDirectiveClass:
     case Stmt::OMPTargetDataDirectiveClass:
+    case Stmt::OMPTargetEnterDataDirectiveClass:
+    case Stmt::OMPTargetExitDataDirectiveClass:
     case Stmt::OMPTeamsDirectiveClass:
     case Stmt::OMPCancellationPointDirectiveClass:
     case Stmt::OMPCancelDirectiveClass:
Index: test/OpenMP/target_enter_data_ast_print.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_enter_data_ast_print.cpp
@@ -0,0 +1,100 @@
+// RUN: %clang_cc1 -verify -fopenmp -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+template <typename T, int C>
+T tmain(T argc, T *argv) {
+  T i, j, b, c, d, e, x[20];
+
+  i = argc;
+#pragma omp target enter data map(to: i)
+
+#pragma omp target enter data map(to: i) if (target enter data: j > 0)
+
+#pragma omp target enter data map(to: i) if (b)
+
+#pragma omp target enter data map(to: c)
+
+#pragma omp target enter data map(to: c) if(b>e)
+
+#pragma omp target enter data map(alloc: x[0:10], c)
+
+#pragma omp target enter data map(to: c) map(alloc: d)
+
+#pragma omp target enter data map(always,alloc: e)
+
+  return 0;
+}
+
+// CHECK: template <typename T = int, int C = 5> int tmain(int argc, int *argv) {
+// CHECK-NEXT: int i, j, b, c, d, e, x[20];
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target enter data map(to: i)
+// CHECK-NEXT: #pragma omp target enter data map(to: i) if(target enter data: j > 0)
+// CHECK-NEXT: #pragma omp target enter data map(to: i) if(b)
+// CHECK-NEXT: #pragma omp target enter data map(to: c)
+// CHECK-NEXT: #pragma omp target enter data map(to: c) if(b > e)
+// CHECK-NEXT: #pragma omp target enter data map(alloc: x[0:10],c)
+// CHECK-NEXT: #pragma omp target enter data map(to: c) map(alloc: d)
+// CHECK-NEXT: #pragma omp target enter data map(always,alloc: e)
+// CHECK: template <typename T = char, int C = 1> char tmain(char argc, char *argv) {
+// CHECK-NEXT: char i, j, b, c, d, e, x[20];
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target enter data map(to: i)
+// CHECK-NEXT: #pragma omp target enter data map(to: i) if(target enter data: j > 0)
+// CHECK-NEXT: #pragma omp target enter data map(to: i) if(b)
+// CHECK-NEXT: #pragma omp target enter data map(to: c)
+// CHECK-NEXT: #pragma omp target enter data map(to: c) if(b > e)
+// CHECK-NEXT: #pragma omp target enter data map(alloc: x[0:10],c)
+// CHECK-NEXT: #pragma omp target enter data map(to: c) map(alloc: d)
+// CHECK-NEXT: #pragma omp target enter data map(always,alloc: e)
+// CHECK: template <typename T, int C> T tmain(T argc, T *argv) {
+// CHECK-NEXT: T i, j, b, c, d, e, x[20];
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target enter data map(to: i)
+// CHECK-NEXT: #pragma omp target enter data map(to: i) if(target enter data: j > 0)
+// CHECK-NEXT: #pragma omp target enter data map(to: i) if(b)
+// CHECK-NEXT: #pragma omp target enter data map(to: c)
+// CHECK-NEXT: #pragma omp target enter data map(to: c) if(b > e)
+// CHECK-NEXT: #pragma omp target enter data map(alloc: x[0:10],c)
+// CHECK-NEXT: #pragma omp target enter data map(to: c) map(alloc: d)
+// CHECK-NEXT: #pragma omp target enter data map(always,alloc: e)
+
+int main (int argc, char **argv) {
+  int b = argc, c, d, e, f, g, x[20];
+  static int a;
+// CHECK: static int a;
+
+#pragma omp target enter data map(to: a)
+// CHECK:      #pragma omp target enter data map(to: a)
+  a=2;
+// CHECK-NEXT: a = 2;
+#pragma omp target enter data map(to: a) if (target enter data: b)
+// CHECK: #pragma omp target enter data map(to: a) if(target enter data: b)
+
+#pragma omp target enter data map(to: a) if (b > g)
+// CHECK: #pragma omp target enter data map(to: a) if(b > g)
+
+#pragma omp target enter data map(to: c)
+// CHECK-NEXT: #pragma omp target enter data map(to: c)
+
+#pragma omp target enter data map(alloc: c) if(b>g)
+// CHECK-NEXT: #pragma omp target enter data map(alloc: c) if(b > g)
+
+#pragma omp target enter data map(to: x[0:10], c)
+// CHECK-NEXT: #pragma omp target enter data map(to: x[0:10],c)
+
+#pragma omp target enter data map(to: c) map(alloc: d)
+// CHECK-NEXT: #pragma omp target enter data map(to: c) map(alloc: d)
+
+#pragma omp target enter data map(always,alloc: e)
+// CHECK-NEXT: #pragma omp target enter data map(always,alloc: e)
+
+  return tmain<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);
+}
+
+#endif
Index: test/OpenMP/target_enter_data_device_messages.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_enter_data_device_messages.cpp
@@ -0,0 +1,29 @@
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+  return argc;
+}
+
+struct S1; // expected-note {{declared here}}
+
+int main(int argc, char **argv) {
+  int i;
+  #pragma omp target enter data map(to: i) device // expected-error {{expected '(' after 'device'}}
+  #pragma omp target enter data map(to: i) device ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  #pragma omp target enter data map(to: i) device () // expected-error {{expected expression}}
+  #pragma omp target enter data map(to: i) device (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  #pragma omp target enter data map(to: i) device (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target enter data' are ignored}}
+#pragma omp target enter data map(to: i) device (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+  #pragma omp target enter data map(to: i) device (argc + argc)
+  #pragma omp target enter data map(to: i) device (argc), device (argc+1) // expected-error {{directive '#pragma omp target enter data' cannot contain more than one 'device' clause}}
+  #pragma omp target enter data map(to: i) device (S1) // expected-error {{'S1' does not refer to a value}}
+  #pragma omp target enter data map(to: i) device (-2) // expected-error {{argument to 'device' clause must be a non-negative integer value}}
+  #pragma omp target enter data map(to: i) device (-10u)
+  #pragma omp target enter data map(to: i) device (3.14) // expected-error {{expression must have integral or unscoped enumeration type, not 'double'}}
+  foo();
+
+  return 0;
+}
Index: test/OpenMP/target_enter_data_if_messages.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_enter_data_if_messages.cpp
@@ -0,0 +1,35 @@
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+  return argc;
+}
+
+struct S1; // expected-note {{declared here}}
+
+int main(int argc, char **argv) {
+  int i;
+  #pragma omp target enter data map(to: i) if // expected-error {{expected '(' after 'if'}}
+  #pragma omp target enter data map(to: i) if ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  #pragma omp target enter data map(to: i) if () // expected-error {{expected expression}}
+  #pragma omp target enter data map(to: i) if (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  #pragma omp target enter data map(to: i) if (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target enter data' are ignored}}
+  #pragma omp target enter data map(to: i) if (argc > 0 ? argv[1] : argv[2])
+  #pragma omp target enter data map(to: i) if (argc + argc)
+  #pragma omp target enter data map(to: i) if (foobool(argc)), if (true) // expected-error {{directive '#pragma omp target enter data' cannot contain more than one 'if' clause}}
+  #pragma omp target enter data map(to: i) if (S1) // expected-error {{'S1' does not refer to a value}}
+  #pragma omp target enter data map(to: i) if (argv[1]=2) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  #pragma omp target enter data map(to: i) if(target data : true) // expected-error {{directive name modifier 'target data' is not allowed for '#pragma omp target enter data'}}
+  #pragma omp target enter data map(to: i) if(target enter data : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  #pragma omp target enter data map(to: i) if(target enter data : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  #pragma omp target enter data map(to: i) if(target enter data : argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  #pragma omp target enter data map(to: i) if(target enter data : argc)
+  #pragma omp target enter data map(to: i) if(target enter data : argc) if (for:argc) // expected-error {{directive name modifier 'for' is not allowed for '#pragma omp target enter data'}}
+  #pragma omp target enter data map(to: i) if(target enter data : argc) if (target enter data:argc) // expected-error {{directive '#pragma omp target enter data' cannot contain more than one 'if' clause with 'target enter data' name modifier}}
+  #pragma omp target enter data map(to: i) if(target enter data : argc) if (argc) // expected-error {{no more 'if' clause is allowed}} expected-note {{previous clause with directive name modifier specified here}}
+  foo();
+
+  return 0;
+}
Index: test/OpenMP/target_enter_data_map_messages.c
===================================================================
--- /dev/null
+++ test/OpenMP/target_enter_data_map_messages.c
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s
+
+int main(int argc, char **argv) {
+
+  int r;
+  #pragma omp target enter data map(r) // expected-error {{map type must be specified for '#pragma omp target enter data'}}
+
+  #pragma omp target enter data // expected-error {{expected at least one map clause for '#pragma omp target enter data'}}
+
+  #pragma omp target enter data map(tofrom: r) // expected-error {{map type 'tofrom' is not allowed for '#pragma omp target enter data'}}
+
+  #pragma omp target enter data map(always, to: r)
+  #pragma omp target enter data map(always, alloc: r)
+  #pragma omp target enter data map(always, from: r) // expected-error {{map type 'from' is not allowed for '#pragma omp target enter data'}}
+  #pragma omp target enter data map(release: r) // expected-error {{map type 'release' is not allowed for '#pragma omp target enter data'}}
+  #pragma omp target enter data map(delete: r) // expected-error {{map type 'delete' is not allowed for '#pragma omp target enter data'}}
+
+  return 0;
+}
Index: test/OpenMP/target_exit_data_ast_print.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_exit_data_ast_print.cpp
@@ -0,0 +1,100 @@
+// RUN: %clang_cc1 -verify -fopenmp -ast-print %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+template <typename T, int C>
+T tmain(T argc, T *argv) {
+  T i, j, b, c, d, e, x[20];
+
+  i = argc;
+#pragma omp target exit data map(from: i)
+
+#pragma omp target exit data map(from: i) if (target exit data: j > 0)
+
+#pragma omp target exit data map(from: i) if (b)
+
+#pragma omp target exit data map(from: c)
+
+#pragma omp target exit data map(from: c) if(b>e)
+
+#pragma omp target exit data map(release: x[0:10], c)
+
+#pragma omp target exit data map(from: c) map(release: d)
+
+#pragma omp target exit data map(always,release: e)
+
+  return 0;
+}
+
+// CHECK: template <typename T = int, int C = 5> int tmain(int argc, int *argv) {
+// CHECK-NEXT: int i, j, b, c, d, e, x[20];
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target exit data map(from: i)
+// CHECK-NEXT: #pragma omp target exit data map(from: i) if(target exit data: j > 0)
+// CHECK-NEXT: #pragma omp target exit data map(from: i) if(b)
+// CHECK-NEXT: #pragma omp target exit data map(from: c)
+// CHECK-NEXT: #pragma omp target exit data map(from: c) if(b > e)
+// CHECK-NEXT: #pragma omp target exit data map(release: x[0:10],c)
+// CHECK-NEXT: #pragma omp target exit data map(from: c) map(release: d)
+// CHECK-NEXT: #pragma omp target exit data map(always,release: e)
+// CHECK: template <typename T = char, int C = 1> char tmain(char argc, char *argv) {
+// CHECK-NEXT: char i, j, b, c, d, e, x[20];
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target exit data map(from: i)
+// CHECK-NEXT: #pragma omp target exit data map(from: i) if(target exit data: j > 0)
+// CHECK-NEXT: #pragma omp target exit data map(from: i) if(b)
+// CHECK-NEXT: #pragma omp target exit data map(from: c)
+// CHECK-NEXT: #pragma omp target exit data map(from: c) if(b > e)
+// CHECK-NEXT: #pragma omp target exit data map(release: x[0:10],c)
+// CHECK-NEXT: #pragma omp target exit data map(from: c) map(release: d)
+// CHECK-NEXT: #pragma omp target exit data map(always,release: e)
+// CHECK: template <typename T, int C> T tmain(T argc, T *argv) {
+// CHECK-NEXT: T i, j, b, c, d, e, x[20];
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target exit data map(from: i)
+// CHECK-NEXT: #pragma omp target exit data map(from: i) if(target exit data: j > 0)
+// CHECK-NEXT: #pragma omp target exit data map(from: i) if(b)
+// CHECK-NEXT: #pragma omp target exit data map(from: c)
+// CHECK-NEXT: #pragma omp target exit data map(from: c) if(b > e)
+// CHECK-NEXT: #pragma omp target exit data map(release: x[0:10],c)
+// CHECK-NEXT: #pragma omp target exit data map(from: c) map(release: d)
+// CHECK-NEXT: #pragma omp target exit data map(always,release: e)
+
+int main (int argc, char **argv) {
+  int b = argc, c, d, e, f, g, x[20];
+  static int a;
+// CHECK: static int a;
+
+#pragma omp target exit data map(from: a)
+// CHECK:      #pragma omp target exit data map(from: a)
+  a=2;
+// CHECK-NEXT: a = 2;
+#pragma omp target exit data map(from: a) if (target exit data: b)
+// CHECK: #pragma omp target exit data map(from: a) if(target exit data: b)
+
+#pragma omp target exit data map(from: a) if (b > g)
+// CHECK: #pragma omp target exit data map(from: a) if(b > g)
+
+#pragma omp target exit data map(from: c)
+// CHECK-NEXT: #pragma omp target exit data map(from: c)
+
+#pragma omp target exit data map(release: c) if(b>g)
+// CHECK-NEXT: #pragma omp target exit data map(release: c) if(b > g)
+
+#pragma omp target exit data map(from: x[0:10], c)
+// CHECK-NEXT: #pragma omp target exit data map(from: x[0:10],c)
+
+#pragma omp target exit data map(from: c) map(release: d)
+// CHECK-NEXT: #pragma omp target exit data map(from: c) map(release: d)
+
+#pragma omp target exit data map(always,release: e)
+// CHECK-NEXT: #pragma omp target exit data map(always,release: e)
+
+  return tmain<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);
+}
+
+#endif
Index: test/OpenMP/target_exit_data_device_messages.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_exit_data_device_messages.cpp
@@ -0,0 +1,29 @@
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+  return argc;
+}
+
+struct S1; // expected-note {{declared here}}
+
+int main(int argc, char **argv) {
+  int i;
+  #pragma omp target exit data map(from: i) device // expected-error {{expected '(' after 'device'}}
+  #pragma omp target exit data map(from: i) device ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  #pragma omp target exit data map(from: i) device () // expected-error {{expected expression}}
+  #pragma omp target exit data map(from: i) device (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  #pragma omp target exit data map(from: i) device (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target exit data' are ignored}}
+#pragma omp target exit data map(from: i) device (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+  #pragma omp target exit data map(from: i) device (argc + argc)
+  #pragma omp target exit data map(from: i) device (argc), device (argc+1) // expected-error {{directive '#pragma omp target exit data' cannot contain more than one 'device' clause}}
+  #pragma omp target exit data map(from: i) device (S1) // expected-error {{'S1' does not refer to a value}}
+  #pragma omp target exit data map(from: i) device (-2) // expected-error {{argument to 'device' clause must be a non-negative integer value}}
+  #pragma omp target exit data map(from: i) device (-10u)
+  #pragma omp target exit data map(from: i) device (3.14) // expected-error {{expression must have integral or unscoped enumeration type, not 'double'}}
+  foo();
+
+  return 0;
+}
Index: test/OpenMP/target_exit_data_if_messages.cpp
===================================================================
--- /dev/null
+++ test/OpenMP/target_exit_data_if_messages.cpp
@@ -0,0 +1,35 @@
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+  return argc;
+}
+
+struct S1; // expected-note {{declared here}}
+
+int main(int argc, char **argv) {
+  int i;
+  #pragma omp target exit data map(from: i) if // expected-error {{expected '(' after 'if'}}
+  #pragma omp target exit data map(from: i) if ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  #pragma omp target exit data map(from: i) if () // expected-error {{expected expression}}
+  #pragma omp target exit data map(from: i) if (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  #pragma omp target exit data map(from: i) if (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target exit data' are ignored}}
+  #pragma omp target exit data map(from: i) if (argc > 0 ? argv[1] : argv[2])
+  #pragma omp target exit data map(from: i) if (argc + argc)
+  #pragma omp target exit data map(from: i) if (foobool(argc)), if (true) // expected-error {{directive '#pragma omp target exit data' cannot contain more than one 'if' clause}}
+  #pragma omp target exit data map(from: i) if (S1) // expected-error {{'S1' does not refer to a value}}
+  #pragma omp target exit data map(from: i) if (argv[1]=2) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  #pragma omp target exit data map(from: i) if(target data : true) // expected-error {{directive name modifier 'target data' is not allowed for '#pragma omp target exit data'}}
+  #pragma omp target exit data map(from: i) if(target exit data : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  #pragma omp target exit data map(from: i) if(target exit data : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  #pragma omp target exit data map(from: i) if(target exit data : argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  #pragma omp target exit data map(from: i) if(target exit data : argc)
+  #pragma omp target exit data map(from: i) if(target exit data : argc) if (for:argc) // expected-error {{directive name modifier 'for' is not allowed for '#pragma omp target exit data'}}
+  #pragma omp target exit data map(from: i) if(target exit data : argc) if (target exit data:argc) // expected-error {{directive '#pragma omp target exit data' cannot contain more than one 'if' clause with 'target exit data' name modifier}}
+  #pragma omp target exit data map(from: i) if(target exit data : argc) if (argc) // expected-error {{no more 'if' clause is allowed}} expected-note {{previous clause with directive name modifier specified here}}
+  foo();
+
+  return 0;
+}
Index: test/OpenMP/target_exit_data_map_messages.c
===================================================================
--- /dev/null
+++ test/OpenMP/target_exit_data_map_messages.c
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s
+
+int main(int argc, char **argv) {
+
+  int r;
+  #pragma omp target exit data map(r) // expected-error {{map type must be specified for '#pragma omp target exit data'}}
+
+  #pragma omp target exit data // expected-error {{expected at least one map clause for '#pragma omp target exit data'}}
+
+  #pragma omp target exit data map(tofrom: r) // expected-error {{map type 'tofrom' is not allowed for '#pragma omp target exit data'}}
+
+  #pragma omp target exit data map(always, from: r)
+  #pragma omp target exit data map(delete: r)
+  #pragma omp target exit data map(release: r)
+  #pragma omp target exit data map(always, alloc: r) // expected-error {{map type 'alloc' is not allowed for '#pragma omp target exit data'}}
+  #pragma omp target exit data map(to: r) // expected-error {{map type 'to' is not allowed for '#pragma omp target exit data'}}
+
+  return 0;
+}
Index: tools/libclang/CIndex.cpp
===================================================================
--- tools/libclang/CIndex.cpp
+++ tools/libclang/CIndex.cpp
@@ -1949,6 +1949,8 @@
   void VisitOMPAtomicDirective(const OMPAtomicDirective *D);
   void VisitOMPTargetDirective(const OMPTargetDirective *D);
   void VisitOMPTargetDataDirective(const OMPTargetDataDirective *D);
+  void VisitOMPTargetEnterDataDirective(const OMPTargetEnterDataDirective *D);
+  void VisitOMPTargetExitDataDirective(const OMPTargetExitDataDirective *D);
   void VisitOMPTeamsDirective(const OMPTeamsDirective *D);
   void VisitOMPTaskLoopDirective(const OMPTaskLoopDirective *D);
   void VisitOMPTaskLoopSimdDirective(const OMPTaskLoopSimdDirective *D);
@@ -2624,6 +2626,16 @@
   VisitOMPExecutableDirective(D);
 }
 
+void EnqueueVisitor::VisitOMPTargetEnterDataDirective(
+    const OMPTargetEnterDataDirective *D) {
+  VisitOMPExecutableDirective(D);
+}
+
+void EnqueueVisitor::VisitOMPTargetExitDataDirective(
+    const OMPTargetExitDataDirective *D) {
+  VisitOMPExecutableDirective(D);
+}
+
 void EnqueueVisitor::VisitOMPTeamsDirective(const OMPTeamsDirective *D) {
   VisitOMPExecutableDirective(D);
 }
@@ -4506,6 +4518,10 @@
     return cxstring::createRef("OMPTargetDirective");
   case CXCursor_OMPTargetDataDirective:
     return cxstring::createRef("OMPTargetDataDirective");
+  case CXCursor_OMPTargetEnterDataDirective:
+    return cxstring::createRef("OMPTargetEnterDataDirective");
+  case CXCursor_OMPTargetExitDataDirective:
+    return cxstring::createRef("OMPTargetExitDataDirective");
   case CXCursor_OMPTeamsDirective:
     return cxstring::createRef("OMPTeamsDirective");
   case CXCursor_OMPCancellationPointDirective:
Index: tools/libclang/CXCursor.cpp
===================================================================
--- tools/libclang/CXCursor.cpp
+++ tools/libclang/CXCursor.cpp
@@ -600,6 +600,12 @@
   case Stmt::OMPTargetDataDirectiveClass:
     K = CXCursor_OMPTargetDataDirective;
     break;
+  case Stmt::OMPTargetEnterDataDirectiveClass:
+    K = CXCursor_OMPTargetEnterDataDirective;
+    break;
+  case Stmt::OMPTargetExitDataDirectiveClass:
+    K = CXCursor_OMPTargetExitDataDirective;
+    break;
   case Stmt::OMPTeamsDirectiveClass:
     K = CXCursor_OMPTeamsDirective;
     break;