[OpenMP] Parsing and sema support for target update directive

[sfantao]

Summary

This patch is to add parsing and sema support for target update directive. Support for the to and from clauses will be added by a different patch. This patch also adds support for other clauses that are already implemented upstream and apply to target update, e.g. device and if.

This patch is based on the original post by Kelvin Li.

[ABataev]

Kelvin, thanks for the patch!
One more comment: what about nesting of regions? You need to add a corresponding code and tests

include/clang/AST/OpenMPClause.h

3196 ↗

(On Diff #44176)

New clauses must be added in separate patches after commit of the new directive.

lib/CodeGen/CGStmtOpenMP.cpp

3540

Maybe just do not emit anything, to not crash a compiler.

lib/Parse/ParseOpenMP.cpp

244–245

Probably, we need to add local enumeric for these constants (0, 1, 2 etc.)

[kkwli0]

include/clang/AST/OpenMPClause.h

3196 ↗

(On Diff #44176)

Will do

lib/Parse/ParseOpenMP.cpp

244–245

Will use the latest infrastructure.

[kkwli0]

I rebase the patch to the latest trunk and make use of the infrastructure form parsing and sema.

[sfantao]

Hi Kelvin,

Thanks for working on this! I have some minor comments inlined.

lib/Sema/SemaOpenMP.cpp

10091

I'd add an assertion here to make sure CKind is To/Form.

10147

Same thing.

10778

Given that these checks are the same, I think it would be better to outline the map clause implementation to some static function and pass the clause kind along. Only a few map-specific things would have to be guarded with the kind.

10841

You should update the comment. The conflicts are not with the map clause but with to/from clauses of the same constructs. You can still say, that the concerns here are the same as in map clauses so you can reuse the same checks.

10935

Same thing here.

[kkwli0]

Addressed the comments from the last review: added assert calls and outline the common code in ActOnOpenMPToClause, ActOnOpenMPFromClause and ActOnOpenMPMapClause to a static function.

[sfantao]

Thanks Kelvin,

I'm fine with the patch. Let's wait for Alexey to see if he has any concern.

Thanks again!
Samuel

[ABataev]

include/clang/AST/OpenMPClause.h

3466 ↗

(On Diff #50024)

New clauses must be added in separate patches

[kkwli0]

Thanks for the review.

As request, I will split this patch into two:

1. to and from clause (without test cases)
2. target update directive (with test cases)

I will use this review for the target update directive.

[kkwli0]

Patch for the to and from clauses is in http://reviews.llvm.org/D18488.

[sfantao]

Add parsing and sema support only for target update and already existing clauses (if and device).

[ABataev]

Add a test for nesting of regions here

lib/Parse/ParseOpenMP.cpp

42

Please, add a comma after 'OMPD_update' to reduce number of changes in the next modification of this enum.

[sfantao]

• Add nestings test.

[ABataev]

LG

include/clang-c/Index.h

enum CXCursorKind {
   /** \brief OpenMP taskloop directive.
    */
   CXCursor_OMPTaskLoopDirective          = 258,
 
   /** \brief OpenMP taskloop simd directive.
    */
   CXCursor_OMPTaskLoopSimdDirective      = 259,
 
-   /** \brief OpenMP distribute directive.
    */
   CXCursor_OMPDistributeDirective        = 260,
 
   /** \brief OpenMP target enter data directive.
    */
   CXCursor_OMPTargetEnterDataDirective   = 261,
 
   /** \brief OpenMP target exit data directive.
    */
   CXCursor_OMPTargetExitDataDirective    = 262,
 
   /** \brief OpenMP target parallel directive.
    */
   CXCursor_OMPTargetParallelDirective    = 263,
 
   /** \brief OpenMP target parallel for directive.
    */
   CXCursor_OMPTargetParallelForDirective = 264,
 
-  CXCursor_LastStmt                   = CXCursor_OMPTargetParallelForDirective,
+  /** \brief OpenMP target update directive.
+   */
+  CXCursor_OMPTargetUpdateDirective      = 265,
+
+  CXCursor_LastStmt                      = CXCursor_OMPTargetUpdateDirective,
 
   /**
    * \brief Cursor that represents the translation unit itself.
    *
    * The translation unit cursor exists primarily to act as the root
    * cursor for traversing the contents of a translation unit.
    */
   CXCursor_TranslationUnit               = 300,

include/clang/AST/RecursiveASTVisitor.h

DEF_TRAVERSE_STMT(OMPTargetParallelDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPTargetParallelForDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPTeamsDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
+DEF_TRAVERSE_STMT(OMPTargetUpdateDirective,
+                  { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
 DEF_TRAVERSE_STMT(OMPTaskLoopDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPTaskLoopSimdDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPDistributeDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })

include/clang/AST/StmtOpenMP.h

public:
   ///
   /// \param C AST context.
   /// \param StartLoc Starting location of the directive kind.
   /// \param EndLoc Ending Location of the directive.
   /// \param CollapsedNum Number of collapsed loops.
   /// \param Clauses List of clauses.
   /// \param AssociatedStmt Statement, associated with the directive.
   /// \param Exprs Helper expressions for CodeGen.
-    ///
   static OMPDistributeDirective *
   Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
          unsigned CollapsedNum, ArrayRef<OMPClause *> Clauses,
          Stmt *AssociatedStmt, const HelperExprs &Exprs);
 
   /// \brief Creates an empty directive with the place
   /// for \a NumClauses clauses.
   ///
   /// \param C AST context.
   /// \param CollapsedNum Number of collapsed nested loops.
   /// \param NumClauses Number of clauses.
   ///
   static OMPDistributeDirective *CreateEmpty(const ASTContext &C,
                                              unsigned NumClauses,
                                              unsigned CollapsedNum, EmptyShell);
 
   static bool classof(const Stmt *T) {
     return T->getStmtClass() == OMPDistributeDirectiveClass;
   }
 };
 
+/// \brief This represents '#pragma omp target update' directive.
+///
+/// \code
+/// #pragma omp target update to(a) from(b) device(1)
+/// \endcode
+/// In this example directive '#pragma omp target update' has clause 'to' with
+/// argument 'a', clause 'from' with argument 'b' and clause 'device' with
+/// argument '1'.
+///
+class OMPTargetUpdateDirective : 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.
+  ///
+  OMPTargetUpdateDirective(SourceLocation StartLoc, SourceLocation EndLoc,
+                           unsigned NumClauses)
+      : OMPExecutableDirective(this, OMPTargetUpdateDirectiveClass,
+                               OMPD_target_update, StartLoc, EndLoc, NumClauses,
+                               0) {}
+
+  /// \brief Build an empty directive.
+  ///
+  /// \param NumClauses Number of clauses.
+  ///
+  explicit OMPTargetUpdateDirective(unsigned NumClauses)
+      : OMPExecutableDirective(this, OMPTargetUpdateDirectiveClass,
+                               OMPD_target_update, SourceLocation(),
+                               SourceLocation(), NumClauses, 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.
+  ///
+  static OMPTargetUpdateDirective *Create(const ASTContext &C,
+                                          SourceLocation StartLoc,
+                                          SourceLocation EndLoc,
+                                          ArrayRef<OMPClause *> Clauses);
+
+  /// \brief Creates an empty directive with the place for \a NumClauses
+  /// clauses.
+  ///
+  /// \param C AST context.
+  /// \param NumClauses The number of clauses.
+  ///
+  static OMPTargetUpdateDirective *CreateEmpty(const ASTContext &C,
+                                               unsigned NumClauses, EmptyShell);
+
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == OMPTargetUpdateDirectiveClass;
+  }
+};
+
 } // end namespace clang
 
 #endif

include/clang/Basic/DiagnosticSemaKinds.td

 def err_omp_variable_in_map_and_dsa : Error<
   "%0 variable cannot be in a map clause in '#pragma omp %1' directive">;
 def err_omp_param_or_this_in_clause : Error<
   "expected reference to one of the parameters of function %0%select{| or 'this'}1">;
 def err_omp_expected_uniform_param : Error<
   "expected a reference to a parameter specified in a 'uniform' clause">;
 def err_omp_expected_int_param : Error<
   "expected a reference to an integer-typed parameter">;
+def err_omp_at_least_one_motion_clause_required : Error<
+  "expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'">;
 } // end of OpenMP category
 
 let CategoryName = "Related Result Type Issue" in {
 // Objective-C related result type compatibility
 def warn_related_result_type_compatibility_class : Warning<
   "method is expected to return an instance of its class type "
   "%diff{$, but is declared to return $|"
   ", but is declared to return different type}0,1">;

include/clang/Basic/OpenMPKinds.def

 #define OPENMP_TARGET_EXIT_DATA_CLAUSE(Name)
 #endif
 #ifndef OPENMP_TARGET_PARALLEL_CLAUSE
 #  define OPENMP_TARGET_PARALLEL_CLAUSE(Name)
 #endif
 #ifndef OPENMP_TARGET_PARALLEL_FOR_CLAUSE
 #  define OPENMP_TARGET_PARALLEL_FOR_CLAUSE(Name)
 #endif
+#ifndef OPENMP_TARGET_UPDATE_CLAUSE
+#  define OPENMP_TARGET_UPDATE_CLAUSE(Name)
+#endif
 #ifndef OPENMP_TEAMS_CLAUSE
 #  define OPENMP_TEAMS_CLAUSE(Name)
 #endif
 #ifndef OPENMP_CANCEL_CLAUSE
 #  define OPENMP_CANCEL_CLAUSE(Name)
 #endif
 #ifndef OPENMP_ORDERED_CLAUSE
 #  define OPENMP_ORDERED_CLAUSE(Name)

 OPENMP_DIRECTIVE(target)
 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(target_parallel, "target parallel")
 OPENMP_DIRECTIVE_EXT(target_parallel_for, "target parallel for")
+OPENMP_DIRECTIVE_EXT(target_update, "target update")
 OPENMP_DIRECTIVE_EXT(parallel_for, "parallel for")
 OPENMP_DIRECTIVE_EXT(parallel_for_simd, "parallel for simd")
 OPENMP_DIRECTIVE_EXT(parallel_sections, "parallel sections")
 OPENMP_DIRECTIVE_EXT(for_simd, "for simd")
 OPENMP_DIRECTIVE_EXT(cancellation_point, "cancellation point")
 OPENMP_DIRECTIVE_EXT(declare_reduction, "declare reduction")
 OPENMP_DIRECTIVE_EXT(declare_simd, "declare simd")
 OPENMP_DIRECTIVE(taskloop)

 OPENMP_TARGET_PARALLEL_FOR_CLAUSE(proc_bind)
 OPENMP_TARGET_PARALLEL_FOR_CLAUSE(shared)
 OPENMP_TARGET_PARALLEL_FOR_CLAUSE(reduction)
 OPENMP_TARGET_PARALLEL_FOR_CLAUSE(collapse)
 OPENMP_TARGET_PARALLEL_FOR_CLAUSE(schedule)
 OPENMP_TARGET_PARALLEL_FOR_CLAUSE(ordered)
 OPENMP_TARGET_PARALLEL_FOR_CLAUSE(linear)
 
+// Clauses allowed for OpenMP directive 'target update'.
+// TODO More clauses for 'target update' directive.
+OPENMP_TARGET_UPDATE_CLAUSE(if)
+OPENMP_TARGET_UPDATE_CLAUSE(device)
+
 // Clauses allowed for OpenMP directive 'teams'.
 // TODO More clauses for 'teams' directive.
 OPENMP_TEAMS_CLAUSE(default)
 OPENMP_TEAMS_CLAUSE(private)
 OPENMP_TEAMS_CLAUSE(firstprivate)
 OPENMP_TEAMS_CLAUSE(shared)
 OPENMP_TEAMS_CLAUSE(reduction)
 OPENMP_TEAMS_CLAUSE(num_teams)

 #undef OPENMP_SIMD_CLAUSE
 #undef OPENMP_FOR_CLAUSE
 #undef OPENMP_FOR_SIMD_CLAUSE
 #undef OPENMP_MAP_KIND
 #undef OPENMP_DISTRIBUTE_CLAUSE
 #undef OPENMP_DIST_SCHEDULE_KIND
 #undef OPENMP_DEFAULTMAP_KIND
 #undef OPENMP_DEFAULTMAP_MODIFIER
+#undef OPENMP_TARGET_UPDATE_CLAUSE

include/clang/Basic/StmtNodes.td

 def OMPOrderedDirective : DStmt<OMPExecutableDirective>;
 def OMPAtomicDirective : DStmt<OMPExecutableDirective>;
 def OMPTargetDirective : DStmt<OMPExecutableDirective>;
 def OMPTargetDataDirective : DStmt<OMPExecutableDirective>;
 def OMPTargetEnterDataDirective : DStmt<OMPExecutableDirective>;
 def OMPTargetExitDataDirective : DStmt<OMPExecutableDirective>;
 def OMPTargetParallelDirective : DStmt<OMPExecutableDirective>;
 def OMPTargetParallelForDirective : DStmt<OMPExecutableDirective>;
+def OMPTargetUpdateDirective : DStmt<OMPExecutableDirective>;
 def OMPTeamsDirective : DStmt<OMPExecutableDirective>;
 def OMPCancellationPointDirective : DStmt<OMPExecutableDirective>;
 def OMPCancelDirective : DStmt<OMPExecutableDirective>;
 def OMPTaskLoopDirective : DStmt<OMPLoopDirective>;
 def OMPTaskLoopSimdDirective : DStmt<OMPLoopDirective>;
 def OMPDistributeDirective : DStmt<OMPLoopDirective>;

include/clang/Sema/Sema.h

StmtResult ActOnOpenMPTaskLoopSimdDirective(
       SourceLocation EndLoc,
       llvm::DenseMap<ValueDecl *, Expr *> &VarsWithImplicitDSA);
   /// \brief Called on well-formed '\#pragma omp distribute' after parsing
   /// of the associated statement.
   StmtResult ActOnOpenMPDistributeDirective(
       ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
       SourceLocation EndLoc,
       llvm::DenseMap<ValueDecl *, Expr *> &VarsWithImplicitDSA);
+  /// \brief Called on well-formed '\#pragma omp target update'.
+  StmtResult ActOnOpenMPTargetUpdateDirective(ArrayRef<OMPClause *> Clauses,
+                                              SourceLocation StartLoc,
+                                              SourceLocation EndLoc);
 
   /// Checks correctness of linear modifiers.
   bool CheckOpenMPLinearModifier(OpenMPLinearClauseKind LinKind,
                                  SourceLocation LinLoc);
   /// Checks that the specified declaration matches requirements for the linear
   /// decls.
   bool CheckOpenMPLinearDecl(ValueDecl *D, SourceLocation ELoc,
                              OpenMPLinearClauseKind LinKind, QualType Type);

include/clang/Serialization/ASTBitCodes.h

enum StmtCode {
       STMT_OMP_TARGET_PARALLEL_FOR_DIRECTIVE,
       STMT_OMP_TEAMS_DIRECTIVE,
       STMT_OMP_TASKGROUP_DIRECTIVE,
       STMT_OMP_CANCELLATION_POINT_DIRECTIVE,
       STMT_OMP_CANCEL_DIRECTIVE,
       STMT_OMP_TASKLOOP_DIRECTIVE,
       STMT_OMP_TASKLOOP_SIMD_DIRECTIVE,
       STMT_OMP_DISTRIBUTE_DIRECTIVE,
+      STMT_OMP_TARGET_UPDATE_DIRECTIVE,
       EXPR_OMP_ARRAY_SECTION,
 
       // ARC
       EXPR_OBJC_BRIDGED_CAST,     // ObjCBridgedCastExpr
       
       STMT_MS_DEPENDENT_EXISTS,   // MSDependentExistsStmt
       EXPR_LAMBDA                 // LambdaExpr
     };

lib/AST/StmtOpenMP.cpp

OMPDistributeDirective::CreateEmpty(const ASTContext &C, unsigned NumClauses,
                                     unsigned CollapsedNum, EmptyShell) {
   unsigned Size = llvm::alignTo(sizeof(OMPDistributeDirective),
                                 llvm::alignOf<OMPClause *>());
   void *Mem = C.Allocate(Size + sizeof(OMPClause *) * NumClauses +
                          sizeof(Stmt *) *
                              numLoopChildren(CollapsedNum, OMPD_distribute));
   return new (Mem) OMPDistributeDirective(CollapsedNum, NumClauses);
 }
+
+OMPTargetUpdateDirective *
+OMPTargetUpdateDirective::Create(const ASTContext &C, SourceLocation StartLoc,
+                                 SourceLocation EndLoc,
+                                 ArrayRef<OMPClause *> Clauses) {
+  unsigned Size = llvm::alignTo(sizeof(OMPTargetUpdateDirective),
+                                llvm::alignOf<OMPClause *>());
+  void *Mem = C.Allocate(Size + sizeof(OMPClause *) * Clauses.size());
+  OMPTargetUpdateDirective *Dir =
+      new (Mem) OMPTargetUpdateDirective(StartLoc, EndLoc, Clauses.size());
+  Dir->setClauses(Clauses);
+  return Dir;
+}
+
+OMPTargetUpdateDirective *
+OMPTargetUpdateDirective::CreateEmpty(const ASTContext &C, unsigned NumClauses,
+                                      EmptyShell) {
+  unsigned Size = llvm::alignTo(sizeof(OMPTargetUpdateDirective),
+                                llvm::alignOf<OMPClause *>());
+  void *Mem = C.Allocate(Size + sizeof(OMPClause *) * NumClauses);
+  return new (Mem) OMPTargetUpdateDirective(NumClauses);
+}

lib/AST/StmtPrinter.cpp

void StmtPrinter::VisitOMPTaskLoopSimdDirective(
   PrintOMPExecutableDirective(Node);
 }
 
 void StmtPrinter::VisitOMPDistributeDirective(OMPDistributeDirective *Node) {
   Indent() << "#pragma omp distribute ";
   PrintOMPExecutableDirective(Node);
 }
 
+void StmtPrinter::VisitOMPTargetUpdateDirective(
+    OMPTargetUpdateDirective *Node) {
+  Indent() << "#pragma omp target update ";
+  PrintOMPExecutableDirective(Node);
+}
+
 //===----------------------------------------------------------------------===//
 //  Expr printing methods.
 //===----------------------------------------------------------------------===//
 
 void StmtPrinter::VisitDeclRefExpr(DeclRefExpr *Node) {
   if (auto *OCED = dyn_cast<OMPCapturedExprDecl>(Node->getDecl())) {
     OCED->getInit()->IgnoreImpCasts()->printPretty(OS, nullptr, Policy);
     return;

lib/AST/StmtProfile.cpp

void OMPClauseProfiler::VisitOMPDistScheduleClause(
     const OMPDistScheduleClause *C) {
   VistOMPClauseWithPreInit(C);
   if (auto *S = C->getChunkSize())
     Profiler->VisitStmt(S);
 }
 
 void OMPClauseProfiler::VisitOMPDefaultmapClause(const OMPDefaultmapClause *) {}
 
+void StmtProfiler::VisitOMPTargetUpdateDirective(
+    const OMPTargetUpdateDirective *S) {
+  VisitOMPExecutableDirective(S);
+}
+
 void StmtProfiler::VisitExpr(const Expr *S) {
   VisitStmt(S);
 }
 
 void StmtProfiler::VisitDeclRefExpr(const DeclRefExpr *S) {
   VisitExpr(S);
   if (!Canonical)
     VisitNestedNameSpecifier(S->getQualifier());

lib/Basic/OpenMPKinds.cpp

 #define OPENMP_TARGET_PARALLEL_FOR_CLAUSE(Name)                                \
   case OMPC_##Name:                                                            \
     return true;
 #include "clang/Basic/OpenMPKinds.def"
     default:
       break;
     }
     break;
+  case OMPD_target_update:
+    switch (CKind) {
+#define OPENMP_TARGET_UPDATE_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)                                              \
   case OMPC_##Name:                                                            \
     return true;
 #include "clang/Basic/OpenMPKinds.def"
     default:
       break;
bool clang::isOpenMPTargetExecutionDirective(OpenMPDirectiveKind DKind) {
   // TODO add next directives.
   return DKind == OMPD_target || DKind == OMPD_target_parallel ||
          DKind == OMPD_target_parallel_for;
 }
 
 bool clang::isOpenMPTargetDataManagementDirective(OpenMPDirectiveKind DKind) {
   // TODO add target update directive check.
   return DKind == OMPD_target_data || DKind == OMPD_target_enter_data ||
-         DKind == OMPD_target_exit_data;
+         DKind == OMPD_target_exit_data || DKind == OMPD_target_update;
 }
 
 bool clang::isOpenMPTeamsDirective(OpenMPDirectiveKind DKind) {
   return DKind == OMPD_teams; // TODO add next directives.
 }
 
 bool clang::isOpenMPSimdDirective(OpenMPDirectiveKind DKind) {
   return DKind == OMPD_simd || DKind == OMPD_for_simd ||

lib/CodeGen/CGStmt.cpp

case Stmt::OMPTaskLoopDirectiveClass:
     EmitOMPTaskLoopDirective(cast<OMPTaskLoopDirective>(*S));
     break;
   case Stmt::OMPTaskLoopSimdDirectiveClass:
     EmitOMPTaskLoopSimdDirective(cast<OMPTaskLoopSimdDirective>(*S));
     break;
   case Stmt::OMPDistributeDirectiveClass:
     EmitOMPDistributeDirective(cast<OMPDistributeDirective>(*S));
     break;
+  case Stmt::OMPTargetUpdateDirectiveClass:
+    EmitOMPTargetUpdateDirective(cast<OMPTargetUpdateDirective>(*S));
+    break;
   }
 }
 
 bool CodeGenFunction::EmitSimpleStmt(const Stmt *S) {
   switch (S->getStmtClass()) {
   default: return false;
   case Stmt::NullStmtClass: break;
   case Stmt::CompoundStmtClass: EmitCompoundStmt(cast<CompoundStmt>(*S)); break;

lib/CodeGen/CGStmtOpenMP.cpp

 void CodeGenFunction::EmitOMPTaskLoopDirective(const OMPTaskLoopDirective &S) {
   EmitOMPTaskLoopBasedDirective(S);
 }
 
 void CodeGenFunction::EmitOMPTaskLoopSimdDirective(
     const OMPTaskLoopSimdDirective &S) {
   EmitOMPTaskLoopBasedDirective(S);
 }
+
+// Generate the instructions for '#pragma omp target update' directive.
+void CodeGenFunction::EmitOMPTargetUpdateDirective(
+    const OMPTargetUpdateDirective &S) {
+  // TODO: codegen for target update
+}

lib/CodeGen/CodeGenFunction.h

#endif
   void EmitOMPTaskgroupDirective(const OMPTaskgroupDirective &S);
   void EmitOMPFlushDirective(const OMPFlushDirective &S);
   void EmitOMPOrderedDirective(const OMPOrderedDirective &S);
   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 EmitOMPTargetUpdateDirective(const OMPTargetUpdateDirective &S);
   void EmitOMPTargetParallelDirective(const OMPTargetParallelDirective &S);
   void
   EmitOMPTargetParallelForDirective(const OMPTargetParallelForDirective &S);
   void EmitOMPTeamsDirective(const OMPTeamsDirective &S);
   void
   EmitOMPCancellationPointDirective(const OMPCancellationPointDirective &S);
   void EmitOMPCancelDirective(const OMPCancelDirective &S);
   void EmitOMPTaskLoopBasedDirective(const OMPLoopDirective &S);

lib/Parse/ParseOpenMP.cpp

enum OpenMPDirectiveKindEx {
   OMPD_declare,
   OMPD_end,
   OMPD_end_declare,
   OMPD_enter,
   OMPD_exit,
   OMPD_point,
   OMPD_reduction,
   OMPD_target_enter,
-  OMPD_target_exit
+  OMPD_target_exit,
+  OMPD_update,
 };
 
 class ThreadprivateListParserHelper final {
   SmallVector<Expr *, 4> Identifiers;
   Parser *P;
 
 public:
   ThreadprivateListParserHelper(Parser *P) : P(P) {}
return llvm::StringSwitch<unsigned>(S)
       .Case("cancellation", OMPD_cancellation)
       .Case("data", OMPD_data)
       .Case("declare", OMPD_declare)
       .Case("end", OMPD_end)
       .Case("enter", OMPD_enter)
       .Case("exit", OMPD_exit)
       .Case("point", OMPD_point)
       .Case("reduction", OMPD_reduction)
+      .Case("update", OMPD_update)
       .Default(OMPD_unknown);
 }
 
 static OpenMPDirectiveKind ParseOpenMPDirectiveKind(Parser &P) {
   // Array of foldings: F[i][0] F[i][1] ===> F[i][2].
   // E.g.: OMPD_for OMPD_simd ===> OMPD_for_simd
   // TODO: add other combined directives in topological order.
   static const unsigned F[][3] = {
     { OMPD_cancellation, OMPD_point, OMPD_cancellation_point },
     { OMPD_declare, OMPD_reduction, OMPD_declare_reduction },
     { OMPD_declare, OMPD_simd, OMPD_declare_simd },
     { OMPD_declare, OMPD_target, OMPD_declare_target },
     { OMPD_end, OMPD_declare, OMPD_end_declare },
     { OMPD_end_declare, OMPD_target, OMPD_end_declare_target },
     { OMPD_target, OMPD_data, OMPD_target_data },
     { OMPD_target, OMPD_enter, OMPD_target_enter },
     { OMPD_target, OMPD_exit, OMPD_target_exit },
+    { OMPD_target, OMPD_update, OMPD_target_update },
     { OMPD_target_enter, OMPD_data, OMPD_target_enter_data },
     { OMPD_target_exit, OMPD_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 },
     { OMPD_parallel, OMPD_sections, OMPD_parallel_sections },
     { OMPD_taskloop, OMPD_simd, OMPD_taskloop_simd },
     { OMPD_target, OMPD_parallel, OMPD_target_parallel },
if (TR.isUsable()) {
       if (!ReductionType.isNull()) {
         ReductionTypes.push_back(
             std::make_pair(ReductionType, Range.getBegin()));
       }
     } else {
       SkipUntil(tok::comma, tok::colon, tok::annot_pragma_openmp_end,
                 StopBeforeMatch);
     }
 
     if (Tok.is(tok::colon) || Tok.is(tok::annot_pragma_openmp_end))
       break;
 
     // Consume ','.
     if (ExpectAndConsume(tok::comma)) {
       IsCorrect = false;
       if (Tok.is(tok::annot_pragma_openmp_end)) {
         Diag(Tok.getLocation(), diag::err_expected_type);
         return DeclGroupPtrTy();
Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
   case OMPD_target_enter_data:
   case OMPD_target_exit_data:
   case OMPD_target_parallel:
   case OMPD_target_parallel_for:
   case OMPD_taskloop:
   case OMPD_taskloop_simd:
   case OMPD_distribute:
   case OMPD_end_declare_target:
+  case OMPD_target_update:
     Diag(Tok, diag::err_omp_unexpected_directive)
         << getOpenMPDirectiveName(DKind);
     break;
   }
   while (Tok.isNot(tok::annot_pragma_openmp_end))
     ConsumeAnyToken();
   ConsumeAnyToken();
   return nullptr;

 ///       executable-directive:
 ///         annot_pragma_openmp 'parallel' | 'simd' | 'for' | 'sections' |
 ///         'section' | 'single' | 'master' | 'critical' [ '(' <name> ')' ] |
 ///         'parallel for' | 'parallel sections' | 'task' | 'taskyield' |
 ///         'barrier' | 'taskwait' | 'flush' | 'ordered' | 'atomic' |
 ///         'for simd' | 'parallel for simd' | 'target' | 'target data' |
 ///         'taskgroup' | 'teams' | 'taskloop' | 'taskloop simd' |
 ///         'distribute' | 'target enter data' | 'target exit data' |
-///         'target parallel' | 'target parallel for' {clause}
+///         'target parallel' | 'target parallel for' |
+///         'target update' {clause}
 ///         annot_pragma_openmp_end
 ///
 StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(
     AllowedContsructsKind Allowed) {
   assert(Tok.is(tok::annot_pragma_openmp) && "Not an OpenMP directive!");
   ParenBraceBracketBalancer BalancerRAIIObj(*this);
   SmallVector<OMPClause *, 5> Clauses;
   SmallVector<llvm::PointerIntPair<OMPClause *, 1, bool>, OMPC_unknown + 1>
case OMPD_flush:
     }
   case OMPD_taskyield:
   case OMPD_barrier:
   case OMPD_taskwait:
   case OMPD_cancellation_point:
   case OMPD_cancel:
   case OMPD_target_enter_data:
   case OMPD_target_exit_data:
+  case OMPD_target_update:
     if (Allowed == ACK_StatementsOpenMPNonStandalone) {
       Diag(Tok, diag::err_omp_immediate_directive)
           << getOpenMPDirectiveName(DKind) << 0;
     }
     HasAssociatedStatement = false;
     // Fall through for further analysis.
   case OMPD_parallel:
   case OMPD_simd:

lib/Sema/SemaOpenMP.cpp

void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
   case OMPD_cancel:
   case OMPD_flush:
   case OMPD_target_enter_data:
   case OMPD_target_exit_data:
   case OMPD_declare_reduction:
   case OMPD_declare_simd:
   case OMPD_declare_target:
   case OMPD_end_declare_target:
+  case OMPD_target_update:
     llvm_unreachable("OpenMP Directive is not allowed");
   case OMPD_unknown:
     llvm_unreachable("Unknown OpenMP directive");
   }
 }
 
 static OMPCapturedExprDecl *buildCaptureDecl(Sema &S, IdentifierInfo *Id,
                                              Expr *CaptureExpr, bool WithInit,
case OMPD_taskloop_simd:
     Res = ActOnOpenMPTaskLoopSimdDirective(ClausesWithImplicit, AStmt, StartLoc,
                                            EndLoc, VarsWithInheritedDSA);
     AllowedNameModifiers.push_back(OMPD_taskloop);
     break;
   case OMPD_distribute:
     Res = ActOnOpenMPDistributeDirective(ClausesWithImplicit, AStmt, StartLoc,
                                          EndLoc, VarsWithInheritedDSA);
     break;
+  case OMPD_target_update:
+    assert(!AStmt && "Statement is not allowed for target update");
+    Res =
+        ActOnOpenMPTargetUpdateDirective(ClausesWithImplicit, StartLoc, EndLoc);
+    AllowedNameModifiers.push_back(OMPD_target_update);
+    break;
   case OMPD_declare_target:
   case OMPD_end_declare_target:
   case OMPD_threadprivate:
   case OMPD_declare_reduction:
   case OMPD_declare_simd:
     llvm_unreachable("OpenMP Directive is not allowed");
   case OMPD_unknown:
     llvm_unreachable("Unknown OpenMP 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::ActOnOpenMPTargetUpdateDirective(ArrayRef<OMPClause *> Clauses,
+                                                  SourceLocation StartLoc,
+                                                  SourceLocation EndLoc) {
+  // TODO: Set this flag accordingly when we add support for the 'to' and 'from'
+  // clauses.
+  bool seenMotionClause = false;
+
+  if (!seenMotionClause) {
+    Diag(StartLoc, diag::err_omp_at_least_one_motion_clause_required);
+    return StmtError();
+  }
+  return OMPTargetUpdateDirective::Create(Context, StartLoc, EndLoc, Clauses);
+}
+
 StmtResult Sema::ActOnOpenMPTeamsDirective(ArrayRef<OMPClause *> Clauses,
                                            Stmt *AStmt, SourceLocation StartLoc,
                                            SourceLocation EndLoc) {
   if (!AStmt)
     return StmtError();
 
   CapturedStmt *CS = cast<CapturedStmt>(AStmt);
   // 1.2.2 OpenMP Language Terminology
bool FoundError = DSAS->checkMappableExprComponentListsForDecl(
         //  original storage.
         //
         // If the expressions are exactly the same or one is a subset of the
         // other, it means they are sharing storage.
         if (CI == CE && SI == SE) {
           if (CurrentRegionOnly) {
             SemaRef.Diag(ELoc, diag::err_omp_map_shared_storage) << ERange;
             SemaRef.Diag(RE->getExprLoc(), diag::note_used_here)
                 << RE->getSourceRange();
             return true;
           } else {
             // If we find the same expression in the enclosing data environment,
             // that is legal.
             IsEnclosedByDataEnvironmentExpr = true;
             return false;
           }
         }
bool FoundError = DSAS->checkMappableExprComponentListsForDecl(
         //  List items of map clauses in the same construct must not share
         //  original storage.
         //
         // An expression is a subset of the other.
         if (CurrentRegionOnly && (CI == CE || SI == SE)) {
           SemaRef.Diag(ELoc, diag::err_omp_map_shared_storage) << ERange;
           SemaRef.Diag(RE->getExprLoc(), diag::note_used_here)
               << RE->getSourceRange();
           return true;
         }
 
         // The current expression uses the same base as other expression in the
         // data environment but does not contain it completely.
         if (!CurrentRegionOnly && SI != SE)
           EnclosingExpr = RE;
 
         // The current expression is a subset of the expression in the data
for (auto &RE : VarList) {
     Vars.push_back(RE);
 
     // Store the components in the stack so that they can be used to check
     // against other clauses later on.
     DSAStack->addMappableExpressionComponents(CurDeclaration, CurComponents);
 
     // Save the components and declaration to create the clause. For purposes of
     // the clause creation, any component list that has has base 'this' uses
-    // null has
+    // null as base declaration.
     ClauseComponents.resize(ClauseComponents.size() + 1);
     ClauseComponents.back().append(CurComponents.begin(), CurComponents.end());
     ClauseBaseDeclarations.push_back(isa<MemberExpr>(BE) ? nullptr
                                                          : CurDeclaration);
   }
 
   // We need to produce a map clause even if we don't have variables so that
   // other diagnostics related with non-existing map clauses are accurate.
OMPClause *Sema::ActOnOpenMPDefaultmapClause(
   return new (Context)
       OMPDefaultmapClause(StartLoc, LParenLoc, MLoc, KindLoc, EndLoc, Kind, M);
 }
 
 bool Sema::ActOnStartOpenMPDeclareTargetDirective(SourceLocation Loc) {
   DeclContext *CurLexicalContext = getCurLexicalContext();
   if (!CurLexicalContext->isFileContext() &&
       !CurLexicalContext->isExternCContext() &&
       !CurLexicalContext->isExternCXXContext()) {
     Diag(Loc, diag::err_omp_region_not_file_context);
     return false;
   }
   if (IsInOpenMPDeclareTargetContext) {
     Diag(Loc, diag::err_omp_enclosed_declare_target);
     return false;
   }
 
if (!ND->hasAttr<OMPDeclareTargetDeclAttr>()) {
       if (ASTMutationListener *ML = Context.getASTMutationListener())
         ML->DeclarationMarkedOpenMPDeclareTarget(ND, A);
       checkDeclIsAllowedInOpenMPTarget(nullptr, ND);
     } else if (ND->getAttr<OMPDeclareTargetDeclAttr>()->getMapType() != MT) {
       Diag(Id.getLoc(), diag::err_omp_declare_target_to_and_link)
           << Id.getName();
     }
   } else
     Diag(Id.getLoc(), diag::err_omp_invalid_target_decl) << Id.getName();
 }
 
 static void checkDeclInTargetContext(SourceLocation SL, SourceRange SR,
                                      Sema &SemaRef, Decl *D) {
   if (!D)
     return;
   Decl *LD = nullptr;
   if (isa<TagDecl>(D)) {
void Sema::checkDeclIsAllowedInOpenMPTarget(Expr *E, Decl *D) {
   SourceRange SR = E ? E->getSourceRange() : D->getSourceRange();
   SourceLocation SL = E ? E->getLocStart() : D->getLocation();
   // 2.10.6: threadprivate variable cannot appear in a declare target directive.
   if (VarDecl *VD = dyn_cast<VarDecl>(D)) {
     if (DSAStack->isThreadPrivate(VD)) {
       Diag(SL, diag::err_omp_threadprivate_in_target);
       ReportOriginalDSA(*this, DSAStack, VD, DSAStack->getTopDSA(VD, false));
       return;
     }
   }
   if (ValueDecl *VD = dyn_cast<ValueDecl>(D)) {
     // Problem if any with var declared with incomplete type will be reported
     // as normal, so no need to check it here.
     if ((E || !VD->getType()->isIncompleteType()) &&
         !checkValueDeclInTarget(SL, SR, *this, DSAStack, VD)) {
       // Mark decl as declared target to prevent further diagnostic.
       if (isa<VarDecl>(VD) || isa<FunctionDecl>(VD)) {

lib/Sema/TreeTransform.h

StmtResult TreeTransform<Derived>::TransformOMPTargetParallelForDirective(
   getDerived().getSema().StartOpenMPDSABlock(OMPD_target_parallel_for, DirName,
                                              nullptr, D->getLocStart());
   StmtResult Res = getDerived().TransformOMPExecutableDirective(D);
   getDerived().getSema().EndOpenMPDSABlock(Res.get());
   return Res;
 }
 
 template <typename Derived>
+StmtResult TreeTransform<Derived>::TransformOMPTargetUpdateDirective(
+    OMPTargetUpdateDirective *D) {
+  DeclarationNameInfo DirName;
+  getDerived().getSema().StartOpenMPDSABlock(OMPD_target_update, 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;
   getDerived().getSema().StartOpenMPDSABlock(OMPD_teams, DirName, nullptr,
                                              D->getLocStart());
   StmtResult Res = getDerived().TransformOMPExecutableDirective(D);
   getDerived().getSema().EndOpenMPDSABlock(Res.get());
   return Res;

lib/Serialization/ASTReaderStmt.cpp

 void ASTStmtReader::VisitOMPTaskLoopSimdDirective(OMPTaskLoopSimdDirective *D) {
   VisitOMPLoopDirective(D);
 }
 
 void ASTStmtReader::VisitOMPDistributeDirective(OMPDistributeDirective *D) {
   VisitOMPLoopDirective(D);
 }
 
+void ASTStmtReader::VisitOMPTargetUpdateDirective(OMPTargetUpdateDirective *D) {
+  VisitStmt(D);
+  ++Idx;
+  VisitOMPExecutableDirective(D);
+}
+
 //===----------------------------------------------------------------------===//
 // ASTReader Implementation
 //===----------------------------------------------------------------------===//
 
 Stmt *ASTReader::ReadStmt(ModuleFile &F) {
   switch (ReadingKind) {
   case Read_None:
     llvm_unreachable("should not call this when not reading anything");
while (true) {
     case STMT_OMP_TARGET_PARALLEL_FOR_DIRECTIVE: {
       unsigned NumClauses = Record[ASTStmtReader::NumStmtFields];
       unsigned CollapsedNum = Record[ASTStmtReader::NumStmtFields + 1];
       S = OMPTargetParallelForDirective::CreateEmpty(Context, NumClauses,
                                                      CollapsedNum, Empty);
       break;
     }
 
+    case STMT_OMP_TARGET_UPDATE_DIRECTIVE:
+      S = OMPTargetUpdateDirective::CreateEmpty(
+          Context, Record[ASTStmtReader::NumStmtFields], Empty);
+      break;
+
     case STMT_OMP_TEAMS_DIRECTIVE:
       S = OMPTeamsDirective::CreateEmpty(
           Context, Record[ASTStmtReader::NumStmtFields], Empty);
       break;
 
     case STMT_OMP_CANCELLATION_POINT_DIRECTIVE:
       S = OMPCancellationPointDirective::CreateEmpty(Context, Empty);
       break;

lib/Serialization/ASTWriterStmt.cpp

void ASTStmtWriter::VisitOMPTaskLoopSimdDirective(OMPTaskLoopSimdDirective *D) {
   Code = serialization::STMT_OMP_TASKLOOP_SIMD_DIRECTIVE;
 }
 
 void ASTStmtWriter::VisitOMPDistributeDirective(OMPDistributeDirective *D) {
   VisitOMPLoopDirective(D);
   Code = serialization::STMT_OMP_DISTRIBUTE_DIRECTIVE;
 }
 
+void ASTStmtWriter::VisitOMPTargetUpdateDirective(OMPTargetUpdateDirective *D) {
+  VisitStmt(D);
+  Record.push_back(D->getNumClauses());
+  VisitOMPExecutableDirective(D);
+  Code = serialization::STMT_OMP_TARGET_UPDATE_DIRECTIVE;
+}
+
 //===----------------------------------------------------------------------===//
 // ASTWriter Implementation
 //===----------------------------------------------------------------------===//
 
 unsigned ASTWriter::RecordSwitchCaseID(SwitchCase *S) {
   assert(SwitchCaseIDs.find(S) == SwitchCaseIDs.end() &&
          "SwitchCase recorded twice");
   unsigned NextID = SwitchCaseIDs.size();

lib/StaticAnalyzer/Core/ExprEngine.cpp

switch (S->getStmtClass()) {
     case Stmt::OMPOrderedDirectiveClass:
     case Stmt::OMPAtomicDirectiveClass:
     case Stmt::OMPTargetDirectiveClass:
     case Stmt::OMPTargetDataDirectiveClass:
     case Stmt::OMPTargetEnterDataDirectiveClass:
     case Stmt::OMPTargetExitDataDirectiveClass:
     case Stmt::OMPTargetParallelDirectiveClass:
     case Stmt::OMPTargetParallelForDirectiveClass:
+    case Stmt::OMPTargetUpdateDirectiveClass:
     case Stmt::OMPTeamsDirectiveClass:
     case Stmt::OMPCancellationPointDirectiveClass:
     case Stmt::OMPCancelDirectiveClass:
     case Stmt::OMPTaskLoopDirectiveClass:
     case Stmt::OMPTaskLoopSimdDirectiveClass:
     case Stmt::OMPDistributeDirectiveClass:
       llvm_unreachable("Stmt should not be in analyzer evaluation loop");
 

test/OpenMP/nesting_of_regions.cpp

for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp parallel
   {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'parallel' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp parallel
+  {
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  }
 
 // SIMD DIRECTIVE
 #pragma omp simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target update // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+  }
 
 // FOR DIRECTIVE
 #pragma omp for
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'for' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp for
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'for' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp for
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  }
 
 // FOR SIMD DIRECTIVE
 #pragma omp for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp for simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target update // expected-error {{OpenMP constructs may not be nested inside a simd region}} 
+    bar();
+  }
 
 // SECTIONS DIRECTIVE
 #pragma omp sections
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'sections' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp sections
   {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'sections' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp sections
+  {
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  }
 
 // SECTION DIRECTIVE
 #pragma omp section // expected-error {{orphaned 'omp section' directives are prohibited, it must be closely nested to a sections region}}
   {
     bar();
   }
 #pragma omp sections
   {
#pragma omp taskloop
   }
 #pragma omp sections
   {
 #pragma omp section
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'section' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp sections
+  {
+#pragma omp section
+    {
+      bar();
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    }
+  }
 
 // SINGLE DIRECTIVE
 #pragma omp single
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'single' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp single
   {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'single' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp single
+  {
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    bar();
+  }
 
 // MASTER DIRECTIVE
 #pragma omp master
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'master' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp master
   {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'master' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp master
+  {
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    bar();
+  }
 
 // CRITICAL DIRECTIVE
 #pragma omp critical
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'critical' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp critical
   {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'critical' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp critical
+  {
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    bar();
+  }
 
 // PARALLEL FOR DIRECTIVE
 #pragma omp parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'parallel for' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'parallel for' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp parallel for
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target update //expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  }
 
 // PARALLEL FOR SIMD DIRECTIVE
 #pragma omp parallel for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp parallel for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp parallel for simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target update // expected-error {{OpenMP constructs may not be nested inside a simd region}} 
+    bar();
+  }
 
 // PARALLEL SECTIONS DIRECTIVE
 #pragma omp parallel sections
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'parallel sections' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp parallel sections
   {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'parallel sections' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp parallel sections
+  {
+#pragma omp target update //expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  }
 
 // TASK DIRECTIVE
 #pragma omp task
 #pragma omp for // expected-error {{region cannot be closely nested inside 'task' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp task
 #pragma omp simd
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp task
   {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'task' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp task
+  {
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    bar();
+  }
 
 // ORDERED DIRECTIVE
 #pragma omp ordered
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'ordered' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp ordered
   {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'ordered' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp ordered
+  {
+    bar();
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    bar();
+  }
 
 // ATOMIC DIRECTIVE
 #pragma omp atomic
   // expected-error@+2 {{the statement for 'atomic' must be an expression statement of form '++x;', '--x;', 'x++;', 'x--;', 'x binop= expr;', 'x = x binop expr' or 'x = expr binop x', where x is an l-value expression with scalar type}}
   // expected-note@+1 {{expected an expression statement}}
   {
 #pragma omp for // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
     for (int i = 0; i < 10; ++i)

 #pragma omp atomic
   // expected-error@+2 {{the statement for 'atomic' must be an expression statement of form '++x;', '--x;', 'x++;', 'x--;', 'x binop= expr;', 'x = x binop expr' or 'x = expr binop x', where x is an l-value expression with scalar type}}
   // expected-note@+1 {{expected an expression statement}}
   {
 #pragma omp distribute // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp atomic
+  // expected-error@+2 {{the statement for 'atomic' must be an expression statement of form '++x;', '--x;', 'x++;', 'x--;', 'x binop= expr;', 'x = x binop expr' or 'x = expr binop x', where x is an l-value expression with scalar type}}
+  // expected-note@+1 {{expected an expression statement}}
+  {
+#pragma omp target update // expected-error {{OpenMP constructs may not be nested inside an atomic region}} 
+    bar();
+  }
 
 // TARGET DIRECTIVE
 #pragma omp target
 #pragma omp parallel
   bar();
 #pragma omp target
 #pragma omp for
   for (int i = 0; i < 10; ++i)

 #pragma omp target
   {
 #pragma omp target enter data map(to: a) // expected-error {{region cannot be nested inside 'target' region}}
   }
 #pragma omp target
   {
 #pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target' region}}
   }
+#pragma omp target
+  {
+#pragma omp target update // expected-error {{region cannot be nested inside 'target' region}} 
+  }
 
 // TARGET PARALLEL DIRECTIVE
 #pragma omp target parallel
 #pragma omp parallel
   bar();
 #pragma omp target parallel
 #pragma omp for
   for (int i = 0; i < 10; ++i)

 #pragma omp target parallel
   {
 #pragma omp target enter data map(to: a) // expected-error {{region cannot be nested inside 'target parallel' region}}
   }
 #pragma omp target parallel
   {
 #pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target parallel' region}}
   }
+#pragma omp target parallel
+  {
+#pragma omp target update // expected-error {{region cannot be nested inside 'target parallel' region}} 
+  }
 
 // TARGET PARALLEL FOR DIRECTIVE
 #pragma omp target parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'target parallel for' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp target parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'target parallel for' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp target parallel for
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target update // expected-error {{region cannot be nested inside 'target parallel for' region}} 
+  }
 
 // TEAMS DIRECTIVE
 #pragma omp target
 #pragma omp teams
 #pragma omp parallel
   bar();
 #pragma omp target
 #pragma omp teams

 #pragma omp target
 #pragma omp teams
 #pragma omp distribute
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp distribute
   for (int j = 0; j < 10; ++j)
     ;        
+#pragma omp target
+#pragma omp teams
+  {
+#pragma omp target update // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp target update' directive into a parallel region?}} 
+  }
 
 // TASKLOOP DIRECTIVE
 #pragma omp taskloop
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'taskloop' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
#pragma omp teams // expected-error {{region cannot be closely nested inside 'taskloop' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
     ++a;
   }
 #pragma omp taskloop
   for (int i = 0; i < 10; ++i) {
 #pragma omp taskloop
   for (int i = 0; i < 10; ++i)
     ++a;
   }
+#pragma omp taskloop
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    bar();
+  }
+
+
 // DISTRIBUTE DIRECTIVE
 #pragma omp target
 #pragma omp teams
 #pragma omp distribute
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'distribute' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
#pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target' region}}
   }
 #pragma omp target
 #pragma omp teams
 #pragma omp distribute
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams // expected-error {{region cannot be closely nested inside 'distribute' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
     ++a;
   }
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target update // expected-error {{region cannot be nested inside 'target' region}} 
+    ++a;
+  }
 }
 
 void foo() {
   int a = 0;
 // PARALLEL DIRECTIVE
 #pragma omp parallel
 #pragma omp for
   for (int i = 0; i < 10; ++i)
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp parallel
   {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'parallel' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp parallel
+  {
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    a++;
+  }
 
 // SIMD DIRECTIVE
 #pragma omp simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target update // expected-error {{OpenMP constructs may not be nested inside a simd region}} 
+    a++;
+  }
 
 // FOR DIRECTIVE
 #pragma omp for
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'for' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp for
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'for' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp for
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    ++a;
+  }
 
 // FOR SIMD DIRECTIVE
 #pragma omp for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp for simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target update // expected-error {{OpenMP constructs may not be nested inside a simd region}} 
+    ++a;
+  }
 
 // SECTIONS DIRECTIVE
 #pragma omp sections
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'sections' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp sections
   {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'sections' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp sections
+  {
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  }
 
 // SECTION DIRECTIVE
 #pragma omp section // expected-error {{orphaned 'omp section' directives are prohibited, it must be closely nested to a sections region}}
   {
     bar();
   }
 #pragma omp sections
   {
#pragma omp taskloop
   }
 #pragma omp sections
   {
 #pragma omp section
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'section' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp sections
+  {
+#pragma omp section
+    {
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+      a++;
+    }
+  }
 
 // SINGLE DIRECTIVE
 #pragma omp single
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'single' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp single
   {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'single' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp single
+  {
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    a++;
+  }
 
 // MASTER DIRECTIVE
 #pragma omp master
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'master' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp master
   {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'master' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp master
+  {
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    ++a;
+  }
 
 // CRITICAL DIRECTIVE
 #pragma omp critical
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'critical' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp critical
   {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'critical' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp critical
+  {
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    a++;
+  }
 
 // PARALLEL FOR DIRECTIVE
 #pragma omp parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'parallel for' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'parallel for' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp parallel for
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    a++;
+  }
 
 // PARALLEL FOR SIMD DIRECTIVE
 #pragma omp parallel for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp parallel for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp parallel for simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target update // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+    a++;
+  }
 
 // PARALLEL SECTIONS DIRECTIVE
 #pragma omp parallel sections
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'parallel sections' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp parallel sections
   {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'parallel sections' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp parallel sections
+  {
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  }
 
 // TASK DIRECTIVE
 #pragma omp task
 #pragma omp for // expected-error {{region cannot be closely nested inside 'task' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp task
 #pragma omp simd
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp task
   {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'task' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp task
+  {
+#pragma omp target update // // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    a++;
+  }
 
 // ATOMIC DIRECTIVE
 #pragma omp atomic
   // expected-error@+2 {{the statement for 'atomic' must be an expression statement of form '++x;', '--x;', 'x++;', 'x--;', 'x binop= expr;', 'x = x binop expr' or 'x = expr binop x', where x is an l-value expression with scalar type}}
   // expected-note@+1 {{expected an expression statement}}
   {
 #pragma omp for // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
     for (int i = 0; i < 10; ++i)

 #pragma omp atomic
   // expected-error@+2 {{the statement for 'atomic' must be an expression statement of form '++x;', '--x;', 'x++;', 'x--;', 'x binop= expr;', 'x = x binop expr' or 'x = expr binop x', where x is an l-value expression with scalar type}}
   // expected-note@+1 {{expected an expression statement}}
   {
 #pragma omp distribute // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp atomic
+  // expected-error@+2 {{the statement for 'atomic' must be an expression statement of form '++x;', '--x;', 'x++;', 'x--;', 'x binop= expr;', 'x = x binop expr' or 'x = expr binop x', where x is an l-value expression with scalar type}}
+  // expected-note@+1 {{expected an expression statement}}
+  {
+#pragma omp target update // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
+  }
 
 // TARGET DIRECTIVE
 #pragma omp target
 #pragma omp parallel
   bar();
 #pragma omp target
 #pragma omp for
   for (int i = 0; i < 10; ++i)
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp target
   { 
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'target' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp atomic
+  // expected-error@+2 {{the statement for 'atomic' must be an expression statement of form '++x;', '--x;', 'x++;', 'x--;', 'x binop= expr;', 'x = x binop expr' or 'x = expr binop x', where x is an l-value expression with scalar type}}
+  // expected-note@+1 {{expected an expression statement}}
+  {
+#pragma omp target update // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
+    a++;
+  }
 
 // TARGET PARALLEL DIRECTIVE
 #pragma omp target parallel
 #pragma omp parallel
   bar();
 #pragma omp target parallel
 #pragma omp for
   for (int i = 0; i < 10; ++i)

 #pragma omp target parallel
   {
 #pragma omp target enter data map(to: a) // expected-error {{region cannot be nested inside 'target parallel' region}}
   }
 #pragma omp target parallel
   {
 #pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target parallel' region}}
   }
+#pragma omp target parallel
+  {
+#pragma omp target update // expected-error {{region cannot be nested inside 'target parallel' region}}
+  }
+
 
 // TARGET PARALLEL FOR DIRECTIVE
 #pragma omp target parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'target parallel for' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
     ++a;
   }
 #pragma omp target parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'target parallel for' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp target parallel for
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target update // expected-error {{region cannot be nested inside 'target parallel for' region}}
+    a++;
+  }
 
 // TEAMS DIRECTIVE
 #pragma omp target
 #pragma omp teams
 #pragma omp parallel
   bar();
 #pragma omp target
 #pragma omp teams

 #pragma omp target
 #pragma omp teams
 #pragma omp distribute
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp distribute
   for (int j = 0; j < 10; ++j)
     ;
+#pragma omp target
+#pragma omp teams
+  {
+#pragma omp target update // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp target update' directive into a parallel region?}}
+    ++a;
+  }
 
 // TASKLOOP DIRECTIVE
 #pragma omp taskloop
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'taskloop' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
#pragma omp teams // expected-error {{region cannot be closely nested inside 'taskloop' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
     ++a;
   }
 #pragma omp taskloop
   for (int i = 0; i < 10; ++i) {
 #pragma omp taskloop
   for (int i = 0; i < 10; ++i)
     ++a;
   }
+#pragma omp taskloop
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    ++a;
+  }
 
 // DISTRIBUTE DIRECTIVE
 #pragma omp target
 #pragma omp teams
 #pragma omp distribute
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'distribute' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)
#pragma omp target enter data map(to: a) // expected-error {{region cannot be nested inside 'target' region}}
   }
 #pragma omp target
 #pragma omp teams
 #pragma omp distribute
   for (int i = 0; i < 10; ++i) {
 #pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target update // expected-error {{region cannot be nested inside 'target' region}}
+    ++a;
+  }
 }

test/OpenMP/target_update_device_messages.cpp

+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+  return argc;
+}
+
+struct S1; // expected-note 2 {{declared here}}
+
+template <class T, class S>
+int tmain(T argc, S **argv) {
+  int i;
+#pragma omp target update device // expected-error {{expected '(' after 'device'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device () // expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device (argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device (argc > 0 ? argv[1] : argv[2]) // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device (argc + argc)  // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device (argc), device (argc+1) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'device' clause}}  expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device (S1) // expected-error {{'S1' does not refer to a value}}  expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device (3.14) // expected-error {{expression must have integral or unscoped enumeration type, not 'double'}}  expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device (-2) // expected-error {{argument to 'device' clause must be a non-negative integer value}}  expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+}
+
+int main(int argc, char **argv) {
+  int j;
+#pragma omp target update device // expected-error {{expected '(' after 'device'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device () // expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}  expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device (argc + argc) //  expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device (argc), device (argc+1) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'device' clause}}  expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device (S1) // expected-error {{'S1' does not refer to a value}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device (-2) // expected-error {{argument to 'device' clause must be a non-negative integer value}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update device (3.14) // expected-error {{expression must have integral or unscoped enumeration type, not 'double'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+  return tmain(argc, argv);
+}

test/OpenMP/target_update_if_messages.cpp

+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+  return argc;
+}
+
+struct S1; // expected-note {{declared here}}
+
+template <class T, class S> // expected-note {{declared here}}
+int tmain(T argc, S **argv) {
+  int n;
+#pragma omp target update if // expected-error {{expected '(' after 'if'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if () // expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if (argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if (argc > 0 ? argv[1] : argv[2]) //  expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if (foobool(argc)), if (true) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'if' clause}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if (S) // expected-error {{'S' does not refer to a value}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if (argv[1]=2) // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if (argc argc) // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if(argc) //  expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if(target update // expected-warning {{missing ':' after directive name modifier - ignoring}} expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if(target update : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if(target update : argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if(target update : argc) // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if(target update : argc) if (for:argc) // expected-error {{directive name modifier 'for' is not allowed for '#pragma omp target update'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if(target update : argc) if (target update:argc) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'if' clause with 'target update' name modifier}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if(target update : argc) if (argc) // expected-error {{no more 'if' clause is allowed}} expected-note {{previous clause with directive name modifier specified here}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  return 0;
+}
+
+int main(int argc, char **argv) {
+  int m;
+#pragma omp target update if // expected-error {{expected '(' after 'if'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if () // expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if (argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if (argc > 0 ? argv[1] : argv[2]) // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if (foobool(argc)), if (true) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'if' clause}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if (S1) // expected-error {{'S1' does not refer to a value}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if (argv[1]=2) // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if (argc argc) // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if (1 0) // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if(if(tmain(argc, argv) // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if(target update // expected-warning {{missing ':' after directive name modifier - ignoring}} expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if(target update : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if(target update : argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if(target update : argc) // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if(target update : argc) if (for:argc) // expected-error {{directive name modifier 'for' is not allowed for '#pragma omp target update'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if(target update : argc) if (target update:argc)  // expected-error {{directive '#pragma omp target update' cannot contain more than one 'if' clause with 'target update' name modifier}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update if(target update : argc) if (argc) // expected-error {{no more 'if' clause is allowed}} expected-note {{previous clause with directive name modifier specified here}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  return tmain(argc, argv);
+}

test/OpenMP/target_update_messages.cpp

+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+  return argc;
+}
+
+struct S1; // Aexpected-note {{declared here}}
+
+template <class T, class S> // Aexpected-note {{declared here}}
+int tmain(T argc, S **argv) {
+  int n;
+  return 0;
+}
+
+int main(int argc, char **argv) {
+  int m;
+  #pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  #pragma omp target update { // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  #pragma omp target update ( // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  #pragma omp target update [ // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  #pragma omp target update ] // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  #pragma omp target update ) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+  #pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  {
+    foo();
+  }
+  return tmain(argc, argv);
+}

tools/libclang/CIndex.cpp

CXString clang_getCursorKindSpelling(enum CXCursorKind Kind) {
   case CXCursor_OMPTargetEnterDataDirective:
     return cxstring::createRef("OMPTargetEnterDataDirective");
   case CXCursor_OMPTargetExitDataDirective:
     return cxstring::createRef("OMPTargetExitDataDirective");
   case CXCursor_OMPTargetParallelDirective:
     return cxstring::createRef("OMPTargetParallelDirective");
   case CXCursor_OMPTargetParallelForDirective:
     return cxstring::createRef("OMPTargetParallelForDirective");
+  case CXCursor_OMPTargetUpdateDirective:
+    return cxstring::createRef("OMPTargetUpdateDirective");
   case CXCursor_OMPTeamsDirective:
     return cxstring::createRef("OMPTeamsDirective");
   case CXCursor_OMPCancellationPointDirective:
     return cxstring::createRef("OMPCancellationPointDirective");
   case CXCursor_OMPCancelDirective:
     return cxstring::createRef("OMPCancelDirective");
   case CXCursor_OMPTaskLoopDirective:
     return cxstring::createRef("OMPTaskLoopDirective");

tools/libclang/CXCursor.cpp

case Stmt::OMPTargetExitDataDirectiveClass:
     K = CXCursor_OMPTargetExitDataDirective;
     break;
   case Stmt::OMPTargetParallelDirectiveClass:
     K = CXCursor_OMPTargetParallelDirective;
     break;
   case Stmt::OMPTargetParallelForDirectiveClass:
     K = CXCursor_OMPTargetParallelForDirective;
     break;
+  case Stmt::OMPTargetUpdateDirectiveClass:
+    K = CXCursor_OMPTargetUpdateDirective;
+    break;
   case Stmt::OMPTeamsDirectiveClass:
     K = CXCursor_OMPTeamsDirective;
     break;
   case Stmt::OMPCancellationPointDirectiveClass:
     K = CXCursor_OMPCancellationPointDirective;
     break;
   case Stmt::OMPCancelDirectiveClass:
     K = CXCursor_OMPCancelDirective;