[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

3256

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

lib/CodeGen/CGStmtOpenMP.cpp

3020

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

lib/Parse/ParseOpenMP.cpp

212–213

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

[kkwli0]

include/clang/AST/OpenMPClause.h

3256

Will do

lib/Parse/ParseOpenMP.cpp

212–213

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

9392

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

9452

Same thing.

10045

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.

10108

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.

10202

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

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

40

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/OpenMPClause.h

public:
 
   static bool classof(const OMPClause *T) {
     return T->getClauseKind() == OMPC_hint;
   }
 
   child_range children() { return child_range(&Hint, &Hint + 1); }
 };
 
 /// \brief This represents 'dist_schedule' clause in the '#pragma omp ...'
 /// directive.
 ///
 /// \code
 /// #pragma omp distribute dist_schedule(static, 3)
 /// \endcode
 /// In this example directive '#pragma omp distribute' has 'dist_schedule'
 /// clause with arguments 'static' and '3'.
 ///
public:
   static bool classof(const OMPClause *T) {
     return T->getClauseKind() == OMPC_defaultmap;
   }
 
   child_range children() {
     return child_range(child_iterator(), child_iterator());
   }
 };
+
+/// \brief This represents clause 'from' in the '#pragma omp ...'
+/// directives.
+///
+/// \code
+/// #pragma omp target update from(a,b)
+/// \endcode
+/// In this example directive '#pragma omp target update' has clause 'from'
+/// with the variables 'a' and 'b'.
+///
+class OMPFromClause final : public OMPVarListClause<OMPFromClause>, 
+                            private llvm::TrailingObjects<OMPFromClause, Expr *> {
+  friend TrailingObjects;
+  friend OMPVarListClause;
+  friend class OMPClauseReader;
+
+  /// \brief Build clause with number of variables \a N.
+  ///
+  /// \param StartLoc Starting location of the clause.
+  /// \param EndLoc Ending location of the clause.
+  /// \param N Number of the variables in the clause.
+  ///
+  explicit OMPFromClause(SourceLocation StartLoc, SourceLocation LParenLoc,
+                         SourceLocation EndLoc, unsigned N)
+    : OMPVarListClause<OMPFromClause>(OMPC_from, StartLoc, LParenLoc, EndLoc, N)
+            {}
+  /// \brief Build an empty clause.
+  ///
+  /// \param N Number of variables.
+  ///
+  explicit OMPFromClause(unsigned N)
+    : OMPVarListClause<OMPFromClause>(OMPC_from, SourceLocation(),
+                                      SourceLocation(), SourceLocation(), N) {}
+
+public:
+  /// \brief Creates clause with a list of variables \a VL.
+  ///
+  /// \param C AST context.
+  /// \brief StartLoc Starting location of the clause.
+  /// \brief EndLoc Ending location of the clause.
+  /// \param VL List of references to the variables.
+  ///
+  static OMPFromClause *Create(const ASTContext &C, SourceLocation StartLoc,
+                               SourceLocation LParenLoc, SourceLocation EndLoc,
+                               ArrayRef<Expr *> VL);
+  /// \brief Creates an empty clause with the place for \a N variables.
+  ///
+  /// \param C AST context.
+  /// \param N The number of variables.
+  ///
+  static OMPFromClause *CreateEmpty(const ASTContext &C, unsigned N);
+
+  static bool classof(const OMPClause *T) {
+    return T->getClauseKind() == OMPC_from;
+  }
+
+  child_range children() {
+    return child_range(reinterpret_cast<Stmt **>(varlist_begin()),
+                       reinterpret_cast<Stmt **>(varlist_end()));
+  }
+};
+
+/// \brief This represents clause 'to' in the '#pragma omp ...'
+/// directives.
+///
+/// \code
+/// #pragma omp target update to(a,b)
+/// \endcode
+/// In this example directive '#pragma omp target update' has clause 'to'
+/// with the variables 'a' and 'b'.
+///
+class OMPToClause final : public OMPVarListClause<OMPToClause>,
+                          private llvm::TrailingObjects<OMPToClause, Expr *> {
+  friend TrailingObjects;
+  friend OMPVarListClause;
+  friend class OMPClauseReader;
+
+  /// \brief Build clause with number of variables \a N.
+  ///
+  /// \param StartLoc Starting location of the clause.
+  /// \param EndLoc Ending location of the clause.
+  /// \param N Number of the variables in the clause.
+  ///
+  explicit OMPToClause(SourceLocation StartLoc, SourceLocation LParenLoc,
+                       SourceLocation EndLoc, unsigned N)
+    : OMPVarListClause<OMPToClause>(OMPC_to, StartLoc, LParenLoc, EndLoc, N) {}
+  /// \brief Build an empty clause.
+  ///
+  /// \param N Number of variables.
+  ///
+  explicit OMPToClause(unsigned N)
+    : OMPVarListClause<OMPToClause>(OMPC_to, SourceLocation(), SourceLocation(),
+                                    SourceLocation(), N) {}
+
+public:
+  /// \brief Creates clause with a list of variables \a VL.
+  ///
+  /// \param C AST context.
+  /// \brief StartLoc Starting location of the clause.
+  /// \brief EndLoc Ending location of the clause.
+  /// \param VL List of references to the variables.
+  ///
+  static OMPToClause *Create(const ASTContext &C, SourceLocation StartLoc,
+                             SourceLocation LParenLoc, SourceLocation EndLoc,
+                             ArrayRef<Expr *> VL);
+  /// \brief Creates an empty clause with the place for \a N variables.
+  ///
+  /// \param C AST context.
+  /// \param N The number of variables.
+  ///
+  static OMPToClause *CreateEmpty(const ASTContext &C, unsigned N);
+
+  static bool classof(const OMPClause *T) {
+    return T->getClauseKind() == OMPC_to;
+  }
+
+  child_range children() {
+    return child_range(reinterpret_cast<Stmt **>(varlist_begin()),
+                       reinterpret_cast<Stmt **>(varlist_end()));
+  }
+};
+
 } // end namespace clang
 
 #endif // LLVM_CLANG_AST_OPENMPCLAUSE_H

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)); })

 }
 
 template <typename Derived>
 bool
 RecursiveASTVisitor<Derived>::VisitOMPDefaultmapClause(OMPDefaultmapClause *C) {
   return true;
 }
 
+template <typename Derived>
+bool RecursiveASTVisitor<Derived>::VisitOMPFromClause(OMPFromClause *C) {
+  TRY_TO(VisitOMPClauseList(C));
+  return true;
+}
+
+template <typename Derived>
+bool RecursiveASTVisitor<Derived>::VisitOMPToClause(OMPToClause *C) {
+  TRY_TO(VisitOMPClauseList(C));
+  return true;
+}
+
 // FIXME: look at the following tricky-seeming exprs to see if we
 // need to recurse on anything.  These are ones that have methods
 // returning decls or qualtypes or nestednamespecifier -- though I'm
 // not sure if they own them -- or just seemed very complicated, or
 // had lots of sub-types to explore.
 //
 // VisitOverloadExpr and its children: recurse on template args? etc?
 

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_unexpected_clause_value : Error<
   "expected %0 in OpenMP clause '%1'">;
 def err_omp_expected_var_name_member_expr : Error<
   "expected variable name%select{| or data member of current class}0">;
 def err_omp_expected_var_name_member_expr_or_array_item : Error<
   "expected variable name%select{|, data member of current class}0, array element or array section">;
 def err_omp_expected_named_var_member_or_array_expression: Error<
   "expected expression containing only member accesses and/or array sections based on named variables">;
-def err_omp_bit_fields_forbidden_in_map_clause : Error<
-  "bit fields cannot be used to specify storage in a map clause">;
+def err_omp_bit_fields_forbidden_in_clause : Error<
+  "bit fields cannot be used to specify storage in a '%0' clause">;
 def err_omp_array_section_in_rightmost_expression : Error<
   "array section can only be associated with the rightmost variable in a map clause expression">;
 def err_omp_union_type_not_allowed : Error<
   "mapped storage cannot be derived from a union">;
 def err_omp_expected_access_to_data_field : Error<
   "expected access to data field">;
 def err_omp_multiple_array_items_in_map_clause : Error<
   "multiple array elements associated with the same variable are not allowed in map clauses of the same construct">;

 def err_omp_reduction_ref_type_arg : Error<
   "argument of OpenMP clause 'reduction' must reference the same object in all threads">;
 def err_omp_clause_not_arithmetic_type_arg : Error<
   "arguments of OpenMP clause 'reduction' for 'min' or 'max' must be of %select{scalar|arithmetic}0 type">;
 def err_omp_clause_floating_type_arg : Error<
   "arguments of OpenMP clause 'reduction' with bitwise operators cannot be of floating type">;
 def err_omp_once_referenced : Error<
   "variable can appear only once in OpenMP '%0' clause">;
+def err_omp_once_referenced_in_target_update : Error<
+  "variable can appear only once in OpenMP 'target update' construct">;
 def note_omp_referenced : Note<
   "previously referenced here">;
 def err_omp_reduction_in_task : Error<
   "reduction variables may not be accessed in an explicit task">;
 def err_omp_reduction_id_not_compatible : Error<
   "list item of type %0 is not valid for specified reduction operation: unable to provide default initialization value">;
 def err_omp_prohibited_region : Error<
   "region cannot be%select{| closely}0 nested inside '%1' region"

 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<
   "mappable type cannot contain static members">;
-def err_omp_threadprivate_in_map : Error<
-  "threadprivate variables are not allowed in map clause">;
+def err_omp_threadprivate_in_clause : Error<
+  "threadprivate variables are not allowed in '%0' clause">;
 def err_omp_wrong_ordered_loop_count : Error<
   "the parameter of the 'ordered' clause must be greater than or equal to the parameter of the 'collapse' clause">;
 def note_collapse_loop_count : Note<
   "parameter of the 'collapse' clause">;
 def err_omp_grainsize_num_tasks_mutually_exclusive : Error<
   "'%0' and '%1' clause are mutually exclusive and may not appear on the same directive">;
 def note_omp_previous_grainsize_num_tasks : Note<
   "'%0' clause is specified here">;

 def err_omp_unexpected_schedule_modifier : Error<
   "modifier '%0' cannot be used along with modifier '%1'">;
 def err_omp_schedule_nonmonotonic_static : Error<
   "'nonmonotonic' modifier can only be specified with 'dynamic' or 'guided' schedule kind">;
 def err_omp_schedule_nonmonotonic_ordered : Error<
   "'schedule' clause with 'nonmonotonic' modifier cannot be specified if an 'ordered' clause is specified">;
 def err_omp_ordered_simd : Error<
   "'ordered' clause with a parameter can not be specified in '#pragma omp %0' directive">;
+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(taskloop)
 OPENMP_DIRECTIVE_EXT(taskloop_simd, "taskloop simd")

 OPENMP_CLAUSE(thread_limit, OMPThreadLimitClause)
 OPENMP_CLAUSE(priority, OMPPriorityClause)
 OPENMP_CLAUSE(grainsize, OMPGrainsizeClause)
 OPENMP_CLAUSE(nogroup, OMPNogroupClause)
 OPENMP_CLAUSE(num_tasks, OMPNumTasksClause)
 OPENMP_CLAUSE(hint, OMPHintClause)
 OPENMP_CLAUSE(dist_schedule, OMPDistScheduleClause)
 OPENMP_CLAUSE(defaultmap, OMPDefaultmapClause)
+OPENMP_CLAUSE(to, OMPToClause)
+OPENMP_CLAUSE(from, OMPFromClause)
 
 // Clauses allowed for OpenMP directive 'parallel'.
 OPENMP_PARALLEL_CLAUSE(if)
 OPENMP_PARALLEL_CLAUSE(num_threads)
 OPENMP_PARALLEL_CLAUSE(default)
 OPENMP_PARALLEL_CLAUSE(proc_bind)
 OPENMP_PARALLEL_CLAUSE(private)
 OPENMP_PARALLEL_CLAUSE(firstprivate)

 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)
+OPENMP_TARGET_UPDATE_CLAUSE(to)
+OPENMP_TARGET_UPDATE_CLAUSE(from)
+
 // 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);
 
   OMPClause *ActOnOpenMPSingleExprClause(OpenMPClauseKind Kind,
                                          Expr *Expr,
                                          SourceLocation StartLoc,
                                          SourceLocation LParenLoc,
                                          SourceLocation EndLoc);
   /// \brief Called on well-formed 'if' clause.
   OMPClause *ActOnOpenMPIfClause(OpenMPDirectiveKind NameModifier,
OMPClause *ActOnOpenMPDistScheduleClause(
       OpenMPDistScheduleClauseKind Kind, Expr *ChunkSize,
       SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation KindLoc,
       SourceLocation CommaLoc, SourceLocation EndLoc);
   /// \brief Called on well-formed 'defaultmap' clause.
   OMPClause *ActOnOpenMPDefaultmapClause(
       OpenMPDefaultmapClauseModifier M, OpenMPDefaultmapClauseKind Kind,
       SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation MLoc,
       SourceLocation KindLoc, SourceLocation EndLoc);
+  /// \brief Called on well-formed 'to' clause.
+  OMPClause *ActOnOpenMPToClause(ArrayRef<Expr *> VarList,
+                                 SourceLocation StartLoc,
+                                 SourceLocation LParenLoc,
+                                 SourceLocation EndLoc);
+  /// \brief Called on well-formed 'from' clause.
+  OMPClause *ActOnOpenMPFromClause(ArrayRef<Expr *> VarList,
+                                   SourceLocation StartLoc,
+                                   SourceLocation LParenLoc,
+                                   SourceLocation EndLoc);
 
   /// \brief The kind of conversion being performed.
   enum CheckedConversionKind {
     /// \brief An implicit conversion.
     CCK_ImplicitConversion,
     /// \brief A C-style cast.
     CCK_CStyleCast,
     /// \brief A functional-style cast.

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/OpenMPClause.cpp

const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
   case OMPC_num_teams:
   case OMPC_thread_limit:
   case OMPC_priority:
   case OMPC_grainsize:
   case OMPC_nogroup:
   case OMPC_num_tasks:
   case OMPC_hint:
   case OMPC_defaultmap:
+  case OMPC_to:
+  case OMPC_from:
   case OMPC_unknown:
     break;
   }
 
   return nullptr;
 }
 
 OMPClauseWithPostUpdate *OMPClauseWithPostUpdate::get(OMPClause *C) {
const OMPClauseWithPostUpdate *OMPClauseWithPostUpdate::get(const OMPClause *C) {
   case OMPC_num_teams:
   case OMPC_thread_limit:
   case OMPC_priority:
   case OMPC_grainsize:
   case OMPC_nogroup:
   case OMPC_num_tasks:
   case OMPC_hint:
   case OMPC_defaultmap:
+  case OMPC_to:
+  case OMPC_from:
   case OMPC_unknown:
     break;
   }
 
   return nullptr;
 }
 
 void OMPPrivateClause::setPrivateCopies(ArrayRef<Expr *> VL) {
OMPMapClause *OMPMapClause::Create(const ASTContext &C, SourceLocation StartLoc,
   Clause->setMapLoc(TypeLoc);
   return Clause;
 }
 
 OMPMapClause *OMPMapClause::CreateEmpty(const ASTContext &C, unsigned N) {
   void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(N));
   return new (Mem) OMPMapClause(N);
 }
+
+OMPToClause *OMPToClause::Create(const ASTContext &C, SourceLocation StartLoc,
+                                 SourceLocation LParenLoc, SourceLocation EndLoc,
+                                 ArrayRef<Expr *> VL) {
+  void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(VL.size()));
+  OMPToClause *Clause = new (Mem) OMPToClause(StartLoc, LParenLoc, EndLoc,
+                                              VL.size());
+  Clause->setVarRefs(VL);
+  return Clause;
+}
+
+OMPToClause *OMPToClause::CreateEmpty(const ASTContext &C, unsigned N) {
+  void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(N));
+  return new (Mem) OMPToClause(N);
+}
+
+OMPFromClause *
+OMPFromClause::Create(const ASTContext &C, SourceLocation StartLoc,
+                      SourceLocation LParenLoc, SourceLocation EndLoc,
+                      ArrayRef<Expr *> VL) {
+  void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(VL.size()));
+  OMPFromClause *Clause = new (Mem) OMPFromClause(StartLoc, LParenLoc, EndLoc,
+                                                  VL.size());
+  Clause->setVarRefs(VL);
+  return Clause;
+}
+
+OMPFromClause *OMPFromClause::CreateEmpty(const ASTContext &C, unsigned N) {
+  void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(N));
+  return new (Mem) OMPFromClause(N);
+}

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

if (Node->getMapType() != OMPC_MAP_unknown) {
       OS << getOpenMPSimpleClauseTypeName(OMPC_map, Node->getMapType());
       OS << ':';
     }
     VisitOMPClauseList(Node, ' ');
     OS << ")";
   }
 }
 
+void OMPClausePrinter::VisitOMPToClause(OMPToClause *Node) {
+  if (!Node->varlist_empty()) {
+    OS << "to";
+    VisitOMPClauseList(Node, '(');
+    OS << ")";
+  }
+}
+
+void OMPClausePrinter::VisitOMPFromClause(OMPFromClause *Node) {
+  if (!Node->varlist_empty()) {
+    OS << "from";
+    VisitOMPClauseList(Node, '(');
+    OS << ")";
+  }
+}
+
 void OMPClausePrinter::VisitOMPDistScheduleClause(OMPDistScheduleClause *Node) {
   OS << "dist_schedule(" << getOpenMPSimpleClauseTypeName(
                            OMPC_dist_schedule, Node->getDistScheduleKind());
   if (auto *E = Node->getChunkSize()) {
     OS << ", ";
     E->printPretty(OS, nullptr, Policy);
   }
   OS << ")";
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::VisitOMPGrainsizeClause(const OMPGrainsizeClause *C) {
   Profiler->VisitStmt(C->getGrainsize());
 }
 void OMPClauseProfiler::VisitOMPNumTasksClause(const OMPNumTasksClause *C) {
   Profiler->VisitStmt(C->getNumTasks());
 }
 void OMPClauseProfiler::VisitOMPHintClause(const OMPHintClause *C) {
   Profiler->VisitStmt(C->getHint());
 }
+void OMPClauseProfiler::VisitOMPToClause(const OMPToClause *C) {
+  VisitOMPClauseList(C);
+}
+void OMPClauseProfiler::VisitOMPFromClause(const OMPFromClause *C) {
+  VisitOMPClauseList(C);
+}
 }
 
 void
 StmtProfiler::VisitOMPExecutableDirective(const OMPExecutableDirective *S) {
   VisitStmt(S);
   OMPClauseProfiler P(this);
   ArrayRef<OMPClause *> Clauses = S->clauses();
   for (ArrayRef<OMPClause *>::iterator I = Clauses.begin(), E = Clauses.end();
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

#include "clang/Basic/OpenMPKinds.def"
   case OMPC_simd:
   case OMPC_num_teams:
   case OMPC_thread_limit:
   case OMPC_priority:
   case OMPC_grainsize:
   case OMPC_nogroup:
   case OMPC_num_tasks:
   case OMPC_hint:
+  case OMPC_to:
+  case OMPC_from:
     break;
   }
   llvm_unreachable("Invalid OpenMP simple clause kind");
 }
 
 const char *clang::getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind,
                                                  unsigned Type) {
   switch (Kind) {
#include "clang/Basic/OpenMPKinds.def"
   case OMPC_simd:
   case OMPC_num_teams:
   case OMPC_thread_limit:
   case OMPC_priority:
   case OMPC_grainsize:
   case OMPC_nogroup:
   case OMPC_num_tasks:
   case OMPC_hint:
+  case OMPC_to:
+  case OMPC_from:
     break;
   }
   llvm_unreachable("Invalid OpenMP simple clause kind");
 }
 
 bool clang::isAllowedClauseForDirective(OpenMPDirectiveKind DKind,
                                         OpenMPClauseKind CKind) {
   assert(DKind <= OMPD_unknown);

 #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

static void EmitOMPAtomicExpr(CodeGenFunction &CGF, OpenMPClauseKind Kind,
   case OMPC_thread_limit:
   case OMPC_priority:
   case OMPC_grainsize:
   case OMPC_nogroup:
   case OMPC_num_tasks:
   case OMPC_hint:
   case OMPC_dist_schedule:
   case OMPC_defaultmap:
+  case OMPC_to:
+  case OMPC_from:
     llvm_unreachable("Clause is not allowed in 'omp atomic'.");
   }
 }
 
 void CodeGenFunction::EmitOMPAtomicDirective(const OMPAtomicDirective &S) {
   bool IsSeqCst = S.getSingleClause<OMPSeqCstClause>();
   OpenMPClauseKind Kind = OMPC_unknown;
   for (auto *C : S.clauses()) {
void CodeGenFunction::EmitOMPTaskLoopSimdDirective(
     const OMPTaskLoopSimdDirective &S) {
   // emit the code inside the construct for now
   auto CS = cast<CapturedStmt>(S.getAssociatedStmt());
   CGM.getOpenMPRuntime().emitInlinedDirective(
       *this, OMPD_taskloop_simd,
       [&CS](CodeGenFunction &CGF) { CGF.EmitStmt(CS->getCapturedStmt()); });
 }
 
+// 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 EmitOMPTaskLoopDirective(const OMPTaskLoopDirective &S);

lib/Parse/ParseOpenMP.cpp

enum OpenMPDirectiveKindEx {
   OMPD_cancellation = OMPD_unknown + 1,
   OMPD_data,
   OMPD_declare,
   OMPD_enter,
   OMPD_exit,
   OMPD_point,
   OMPD_reduction,
   OMPD_target_enter,
-  OMPD_target_exit
+  OMPD_target_exit,
+  OMPD_update
 };
 } // namespace
 
 // Map token string to extended OMP token kind that are
 // OpenMPDirectiveKind + OpenMPDirectiveKindEx.
 static unsigned getOpenMPDirectiveKindEx(StringRef S) {
   auto DKind = getOpenMPDirectiveKind(S);
   if (DKind != OMPD_unknown)
     return DKind;
 
   return llvm::StringSwitch<unsigned>(S)
       .Case("cancellation", OMPD_cancellation)
       .Case("data", OMPD_data)
       .Case("declare", OMPD_declare)
       .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_target, OMPD_data, OMPD_target_data },
     { OMPD_target, OMPD_enter, OMPD_target_enter },
     { OMPD_target, OMPD_exit, OMPD_target_exit },
     { OMPD_target_enter, OMPD_data, OMPD_target_enter_data },
     { OMPD_target_exit, OMPD_data, OMPD_target_exit_data },
+    { OMPD_target, OMPD_update, OMPD_target_update },
     { 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 },
     { OMPD_target_parallel, OMPD_for, OMPD_target_parallel_for }
   };
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::ParseOpenMPDeclarativeDirective(AccessSpecifier AS) {
   case OMPD_target_data:
   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_target_update:
     Diag(Tok, diag::err_omp_unexpected_directive)
         << getOpenMPDirectiveName(DKind);
     break;
   }
   while (Tok.isNot(tok::annot_pragma_openmp_end))
     ConsumeAnyToken();
   ConsumeAnyToken();
   return nullptr;

 ///         annot_pragma_openmp_end
 ///
 ///       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' |
+///         '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<Expr *, 5> Identifiers;
   SmallVector<OMPClause *, 5> Clauses;
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:

 ///       default-clause | private-clause | firstprivate-clause | shared-clause
 ///       | linear-clause | aligned-clause | collapse-clause |
 ///       lastprivate-clause | reduction-clause | proc_bind-clause |
 ///       schedule-clause | copyin-clause | copyprivate-clause | untied-clause |
 ///       mergeable-clause | flush-clause | read-clause | write-clause |
 ///       update-clause | capture-clause | seq_cst-clause | device-clause |
 ///       simdlen-clause | threads-clause | simd-clause | num_teams-clause |
 ///       thread_limit-clause | priority-clause | grainsize-clause |
-///       nogroup-clause | num_tasks-clause | hint-clause
+///       nogroup-clause | num_tasks-clause | hint-clause | to-clause | 
+///       from-clause
 ///
 OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
                                      OpenMPClauseKind CKind, bool FirstClause) {
   OMPClause *Clause = nullptr;
   bool ErrorFound = false;
   // Check if clause is allowed for the given directive.
   if (CKind != OMPC_unknown && !isAllowedClauseForDirective(DKind, CKind)) {
     Diag(Tok, diag::err_omp_unexpected_clause) << getOpenMPClauseName(CKind)
OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
   case OMPC_reduction:
   case OMPC_linear:
   case OMPC_aligned:
   case OMPC_copyin:
   case OMPC_copyprivate:
   case OMPC_flush:
   case OMPC_depend:
   case OMPC_map:
+  case OMPC_to:
+  case OMPC_from:
     Clause = ParseOpenMPVarListClause(DKind, CKind);
     break;
   case OMPC_unknown:
     Diag(Tok, diag::warn_omp_extra_tokens_at_eol)
         << getOpenMPDirectiveName(DKind);
     SkipUntil(tok::annot_pragma_openmp_end, StopBeforeMatch);
     break;
   case OMPC_threadprivate:
if (ReductionIdScopeSpec.isEmpty()) {
     }
   }
   return P.ParseUnqualifiedId(ReductionIdScopeSpec, /*EnteringContext*/ false,
                               /*AllowDestructorName*/ false,
                               /*AllowConstructorName*/ false, nullptr,
                               TemplateKWLoc, ReductionId);
 }
 
-/// \brief Parsing of OpenMP clause 'private', 'firstprivate', 'lastprivate',
-/// 'shared', 'copyin', 'copyprivate', 'flush' or 'reduction'.
+/// \brief Parsing of OpenMP clauses
 ///
 ///    private-clause:
 ///       'private' '(' list ')'
 ///    firstprivate-clause:
 ///       'firstprivate' '(' list ')'
 ///    lastprivate-clause:
 ///       'lastprivate' '(' list ')'
 ///    shared-clause:
 ///       'shared' '(' list ')'
 ///    linear-clause:
 ///       'linear' '(' linear-list [ ':' linear-step ] ')'
 ///    aligned-clause:
 ///       'aligned' '(' list [ ':' alignment ] ')'
 ///    reduction-clause:
 ///       'reduction' '(' reduction-identifier ':' list ')'
 ///    copyprivate-clause:
 ///       'copyprivate' '(' list ')'
 ///    flush-clause:
 ///       'flush' '(' list ')'
 ///    depend-clause:
 ///       'depend' '(' in | out | inout : list | source ')'
 ///    map-clause:
 ///       'map' '(' [ [ always , ]
 ///          to | from | tofrom | alloc | release | delete ':' ] list ')';
+///    to-clause:
+///       'to' '(' list ')'
+///    from-clause:
+///       'from' '(' list ')'
 ///
 /// For 'linear' clause linear-list may have the following forms:
 ///  list
 ///  modifier(list)
 /// where modifier is 'val' (C) or 'ref', 'val' or 'uval'(C++).
 OMPClause *Parser::ParseOpenMPVarListClause(OpenMPDirectiveKind DKind,
                                             OpenMPClauseKind Kind) {
   SourceLocation Loc = Tok.getLocation();

lib/Sema/SemaOpenMP.cpp

void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
   case OMPD_barrier:
   case OMPD_taskwait:
   case OMPD_cancellation_point:
   case OMPD_cancel:
   case OMPD_flush:
   case OMPD_target_enter_data:
   case OMPD_target_exit_data:
   case OMPD_declare_reduction:
+  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_threadprivate:
   case OMPD_declare_reduction:
     llvm_unreachable("OpenMP Directive is not allowed");
   case OMPD_unknown:
     llvm_unreachable("Unknown OpenMP directive");
   }
 
   for (auto P : VarsWithInheritedDSA) {
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) {
+  bool seenMotionClause = false;
+  for (auto *C : Clauses) {
+    if (C->getClauseKind() == OMPC_to || C->getClauseKind() == OMPC_from)
+      seenMotionClause = true;
+  }
+  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
OMPClause *Sema::ActOnOpenMPSingleExprClause(OpenMPClauseKind Kind, Expr *Expr,
   case OMPC_seq_cst:
   case OMPC_depend:
   case OMPC_threads:
   case OMPC_simd:
   case OMPC_map:
   case OMPC_nogroup:
   case OMPC_dist_schedule:
   case OMPC_defaultmap:
+  case OMPC_to:
+  case OMPC_from:
   case OMPC_unknown:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;
 }
 
 OMPClause *Sema::ActOnOpenMPIfClause(OpenMPDirectiveKind NameModifier,
                                      Expr *Condition, SourceLocation StartLoc,
OMPClause *Sema::ActOnOpenMPSimpleClause(
   case OMPC_thread_limit:
   case OMPC_priority:
   case OMPC_grainsize:
   case OMPC_nogroup:
   case OMPC_num_tasks:
   case OMPC_hint:
   case OMPC_dist_schedule:
   case OMPC_defaultmap:
+  case OMPC_to:
+  case OMPC_from:
   case OMPC_unknown:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;
 }
 
 static std::string
 getListOfPossibleValues(OpenMPClauseKind K, unsigned First, unsigned Last,
OMPClause *Sema::ActOnOpenMPSingleExprWithArgClause(
   case OMPC_map:
   case OMPC_num_teams:
   case OMPC_thread_limit:
   case OMPC_priority:
   case OMPC_grainsize:
   case OMPC_nogroup:
   case OMPC_num_tasks:
   case OMPC_hint:
+  case OMPC_to:
+  case OMPC_from:
   case OMPC_unknown:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;
 }
 
 static bool checkScheduleModifiers(Sema &S, OpenMPScheduleClauseModifier M1,
                                    OpenMPScheduleClauseModifier M2,
OMPClause *Sema::ActOnOpenMPClause(OpenMPClauseKind Kind,
   case OMPC_num_teams:
   case OMPC_thread_limit:
   case OMPC_priority:
   case OMPC_grainsize:
   case OMPC_num_tasks:
   case OMPC_hint:
   case OMPC_dist_schedule:
   case OMPC_defaultmap:
+  case OMPC_to:
+  case OMPC_from:
   case OMPC_unknown:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;
 }
 
 OMPClause *Sema::ActOnOpenMPNowaitClause(SourceLocation StartLoc,
                                          SourceLocation EndLoc) {
case OMPC_depend:
     Res = ActOnOpenMPDependClause(DepKind, DepLinMapLoc, ColonLoc, VarList, 
                                   StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_map:
     Res = ActOnOpenMPMapClause(MapTypeModifier, MapType, IsMapTypeImplicit,
                                DepLinMapLoc, ColonLoc, VarList, StartLoc,
                                LParenLoc, EndLoc);
     break;
+  case OMPC_to:
+    Res = ActOnOpenMPToClause(VarList, StartLoc, LParenLoc, EndLoc);
+    break;
+  case OMPC_from:
+    Res = ActOnOpenMPFromClause(VarList, StartLoc, LParenLoc, EndLoc);
+    break;
   case OMPC_if:
   case OMPC_final:
   case OMPC_num_threads:
   case OMPC_safelen:
   case OMPC_simdlen:
   case OMPC_collapse:
   case OMPC_default:
   case OMPC_proc_bind:
if (!RD->isInvalidDecl() &&
       return false;
   }
   return true;
 }
 
 // Return the expression of the base of the map clause or null if it cannot
 // be determined and do all the necessary checks to see if the expression is
 // valid as a standalone map clause expression.
-static Expr *CheckMapClauseExpressionBase(Sema &SemaRef, Expr *E) {
+static Expr *CheckMapClauseExpressionBase(Sema &SemaRef, Expr *E,
+                                          OpenMPClauseKind CKind) {
   SourceLocation ELoc = E->getExprLoc();
   SourceRange ERange = E->getSourceRange();
 
   // The base of elements of list in a map clause have to be either:
   //  - a reference to variable or field.
   //  - a member expression.
   //  - an array expression.
   //
if (auto *CurE = dyn_cast<MemberExpr>(E)) {
       }
 
       auto *FD = cast<FieldDecl>(CurE->getMemberDecl());
 
       // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C/C++, p.3]
       //  A bit-field cannot appear in a map clause.
       //
       if (FD->isBitField()) {
-        SemaRef.Diag(ELoc, diag::err_omp_bit_fields_forbidden_in_map_clause)
-            << CurE->getSourceRange();
+        SemaRef.Diag(ELoc, diag::err_omp_bit_fields_forbidden_in_clause) <<
+            CurE->getSourceRange() << getOpenMPClauseName(CKind);
         break;
       }
 
       // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C++, p.1]
       //  If the type of a list item is a reference to a type T then the type
       //  will be considered to be T for all purposes of this clause.
       QualType CurType = BaseE->getType();
       if (CurType->isReferenceType())
static Expr *CheckMapClauseExpressionBase(Sema &SemaRef, Expr *E,
   }
 
   return RelevantExpr;
 }
 
 // Return true if expression E associated with value VD has conflicts with other
 // map information.
 static bool CheckMapConflicts(Sema &SemaRef, DSAStackTy *DSAS, ValueDecl *VD,
-                              Expr *E, bool CurrentRegionOnly) {
+                              Expr *E, bool CurrentRegionOnly, 
+                              OpenMPClauseKind CKind) {
   assert(VD && E);
 
   // Types used to organize the components of a valid map clause.
   typedef std::pair<Expr *, ValueDecl *> MapExpressionComponent;
   typedef SmallVector<MapExpressionComponent, 4> MapExpressionComponents;
 
   // Helper to extract the components in the map clause expression E and store
   // them into MEC. This assumes that E is a valid map clause expression, i.e.
bool FoundError =
         // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.4]
         //  List items of map clauses in the same construct must not share
         //  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) {
+            if (CKind == OMPC_map)
               SemaRef.Diag(ELoc, diag::err_omp_map_shared_storage) << ERange;
+            else
+              SemaRef.Diag(ELoc, 
+                           diag::err_omp_once_referenced_in_target_update)
+                << 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 =
         }
 
         // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.4]
         //  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)) {
+          if (CKind == OMPC_map)
             SemaRef.Diag(ELoc, diag::err_omp_map_shared_storage) << ERange;
+          else
+            SemaRef.Diag(ELoc, 
+                         diag::err_omp_once_referenced_in_target_update)
+              << 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 completelly.
         if (!CurrentRegionOnly && SI != SE)
           EnclosingExpr = RE;
 
for (auto &RE : VarList) {
 
     if (!RE->IgnoreParenImpCasts()->isLValue()) {
       Diag(ELoc, diag::err_omp_expected_named_var_member_or_array_expression)
           << RE->getSourceRange();
       continue;
     }
 
     // Obtain the array or member expression bases if required.
-    auto *BE = CheckMapClauseExpressionBase(*this, SimpleExpr);
+    auto *BE = CheckMapClauseExpressionBase(*this, SimpleExpr, OMPC_map);
     if (!BE)
       continue;
 
     // If the base is a reference to a variable, we rely on that variable for
     // the following checks. If it is a 'this' expression we rely on the field.
     ValueDecl *D = nullptr;
     if (auto *DRE = dyn_cast<DeclRefExpr>(BE)) {
       D = DRE->getDecl();
for (auto &RE : VarList) {
 
     assert((VD || FD) && "Only variables or fields are expected here!");
     (void)FD;
 
     // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.10]
     //  threadprivate variables cannot appear in a map clause.
     if (VD && DSAStack->isThreadPrivate(VD)) {
       auto DVar = DSAStack->getTopDSA(VD, false);
-      Diag(ELoc, diag::err_omp_threadprivate_in_map);
+      Diag(ELoc, diag::err_omp_threadprivate_in_clause) << "map";
       ReportOriginalDSA(*this, DSAStack, VD, DVar);
       continue;
     }
 
     // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.9]
     //  A list item cannot appear in both a map clause and a data-sharing
     //  attribute clause on the same construct.
     //
     // TODO: Implement this check - it cannot currently be tested because of
     // missing implementation of the other data sharing clauses in target
     // directives.
 
     // Check conflicts with other map clause expressions. We check the conflicts
     // with the current construct separately from the enclosing data
     // environment, because the restrictions are different.
     if (CheckMapConflicts(*this, DSAStack, D, SimpleExpr,
-                          /*CurrentRegionOnly=*/true))
+                          /*CurrentRegionOnly=*/true, OMPC_map))
       break;
     if (CheckMapConflicts(*this, DSAStack, D, SimpleExpr,
-                          /*CurrentRegionOnly=*/false))
+                          /*CurrentRegionOnly=*/false, OMPC_map))
       break;
 
     // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C++, p.1]
     //  If the type of a list item is a reference to a type T then the type will
     //  be considered to be T for all purposes of this clause.
     QualType Type = D->getType();
     if (Type->isReferenceType())
       Type = Type->getPointeeType();
if (M != OMPC_DEFAULTMAP_MODIFIER_tofrom ||
     Diag(Loc, diag::err_omp_unexpected_clause_value)
         << Value << getOpenMPClauseName(OMPC_defaultmap);
     return nullptr;
   }
 
   return new (Context)
       OMPDefaultmapClause(StartLoc, LParenLoc, MLoc, KindLoc, EndLoc, Kind, M);
 }
+
+OMPClause *Sema::ActOnOpenMPToClause(ArrayRef<Expr *> VarList,
+                                     SourceLocation StartLoc,
+                                     SourceLocation LParenLoc,
+                                     SourceLocation EndLoc) {
+  SmallVector<Expr *, 4> Vars;
+  for (auto &RE : VarList) {
+    assert(RE && "Null expr in omp to clause");
+    if (isa<DependentScopeDeclRefExpr>(RE)) {
+      // It will be analyzed later.
+      Vars.push_back(RE);
+      continue;
+    }
+    SourceLocation ELoc = RE->getExprLoc();
+
+    auto *VE = RE->IgnoreParenLValueCasts();
+
+    if (VE->isValueDependent() || VE->isTypeDependent() ||
+        VE->isInstantiationDependent() ||
+        VE->containsUnexpandedParameterPack()) {
+      // We can only analyze this information once the missing information is
+      // resolved.
+      Vars.push_back(RE);
+      continue;
+    }
+
+    auto *SimpleExpr = RE->IgnoreParenCasts();
+
+    if (!RE->IgnoreParenImpCasts()->isLValue()) {
+      Diag(ELoc, diag::err_omp_expected_named_var_member_or_array_expression)
+          << RE->getSourceRange();
+      continue;
+    }
+
+    // Obtain the array or member expression bases if required.
+    auto *BE = CheckMapClauseExpressionBase(*this, SimpleExpr, OMPC_to);
+    if (!BE)
+      continue;
+
+    // If the base is a reference to a variable, we rely on that variable for
+    // the following checks. if it is a 'this' expression we rely on the field.
+    ValueDecl *D = nullptr;
+    if (auto *DRE = dyn_cast<DeclRefExpr>(BE)) {
+      D = DRE -> getDecl();
+    } else {
+      auto *ME = cast<MemberExpr>(BE);
+      assert(isa<CXXThisExpr>(ME->getBase()) && "Unexpected expression!");
+      D = ME-> getMemberDecl();
+    }
+    assert(D && "Null decl on from clause.");
+
+    auto *VD = dyn_cast<VarDecl>(D);
+    auto *FD = dyn_cast<FieldDecl>(D);
+
+    assert((VD || FD) && "Only variables or fields are expected here!");
+
+    // OpenMP 4.5 [2.10.5, target update Construct]
+    // threadprivate variables cannot appear in a from clause.
+    if (VD && DSAStack->isThreadPrivate(VD)) {
+      auto DVar = DSAStack->getTopDSA(VD, false);
+      Diag(ELoc, diag::err_omp_threadprivate_in_clause) << "to";
+      ReportOriginalDSA(*this, DSAStack, VD, DVar);
+      continue;
+    }
+
+    // Check conflicts with other map clause expressions. We check the conflicts
+    // with the current construct separately from the enclosing data
+    // environment, because the restrictions are different.
+    if (CheckMapConflicts(*this, DSAStack, D, SimpleExpr,
+                          /*CurrentRegionOnly=*/true, OMPC_to))
+      break;
+
+    // OpenMP 4.5 [2.10.5, target update Construct]
+    // If the type of a list item is a reference to a type T then the type will
+    // be considered to be T for all purpose of this clause.
+    auto Type = D->getType();
+    if (Type->isReferenceType())
+      Type = Type->getPointeeType();
+
+    // OpenMP 4.5 [2.10.5, target update Construct, Restrictions, p.4]
+    // A list item in a to or from clause must have a mappable type.
+    if (!CheckTypeMappable(VE->getExprLoc(), VE->getSourceRange(), *this,
+                           DSAStack, Type))
+      continue;
+
+    Vars.push_back(RE);
+    DSAStack->addExprToVarMapInfo(D, RE);
+  }
+
+  if (Vars.empty())
+    return nullptr;
+
+  return OMPToClause::Create(Context, StartLoc, LParenLoc, EndLoc, Vars);
+}
+
+OMPClause *Sema::ActOnOpenMPFromClause(ArrayRef<Expr *> VarList,
+                                       SourceLocation StartLoc,
+                                       SourceLocation LParenLoc,
+                                       SourceLocation EndLoc) {
+  SmallVector<Expr *, 4> Vars;
+  for (auto &RE : VarList) {
+    assert(RE && "Null expr in omp from clause");
+    if (isa<DependentScopeDeclRefExpr>(RE)) {
+      // It will be analyzed later.
+      Vars.push_back(RE);
+      continue;
+    }
+    SourceLocation ELoc = RE->getExprLoc();
+
+    auto *VE = RE->IgnoreParenLValueCasts();
+
+    if (VE->isValueDependent() || VE->isTypeDependent() ||
+        VE->isInstantiationDependent() ||
+        VE->containsUnexpandedParameterPack()) {
+      // We can only analyze this information once the missing information is
+      // resolved.
+      Vars.push_back(RE);
+      continue;
+    }
+
+    auto *SimpleExpr = RE->IgnoreParenCasts();
+
+    if (!RE->IgnoreParenImpCasts()->isLValue()) {
+      Diag(ELoc, diag::err_omp_expected_named_var_member_or_array_expression)
+          << RE->getSourceRange();
+      continue;
+    }
+
+    // Obtain the array or member expression bases if required.
+    auto *BE = CheckMapClauseExpressionBase(*this, SimpleExpr, OMPC_from);
+    if (!BE)
+      continue;
+
+    // If the base is a reference to a variable, we rely on that variable for
+    // the following checks. if it is a 'this' expression we rely on the field.
+    ValueDecl *D = nullptr;
+    if (auto *DRE = dyn_cast<DeclRefExpr>(BE)) {
+      D = DRE -> getDecl();
+    } else {
+      auto *ME = cast<MemberExpr>(BE);
+      assert(isa<CXXThisExpr>(ME->getBase()) && "Unexpected expression!");
+      D = ME-> getMemberDecl();
+    }
+    assert(D && "Null decl on from clause.");
+
+    auto *VD = dyn_cast<VarDecl>(D);
+    auto *FD = dyn_cast<FieldDecl>(D);
+
+    assert((VD || FD) && "Only variables or fields are expected here!");
+
+    // OpenMP 4.5 [2.10.5, target update Construct]
+    // threadprivate variables cannot appear in a from clause.
+    if (VD && DSAStack->isThreadPrivate(VD)) {
+      auto DVar = DSAStack->getTopDSA(VD, false);
+      Diag(ELoc, diag::err_omp_threadprivate_in_clause) << "from";
+      ReportOriginalDSA(*this, DSAStack, VD, DVar);
+      continue;
+    }
+
+    // Check conflicts with other map clause expressions. We check the conflicts
+    // with the current construct separately from the enclosing data
+    // environment, because the restrictions are different.
+    if (CheckMapConflicts(*this, DSAStack, D, SimpleExpr,
+                          /*CurrentRegionOnly=*/true, OMPC_from))
+      break;
+
+    // OpenMP 4.5 [2.10.5, target update Construct]
+    // If the type of a list item is a reference to a type T then the type will
+    // be considered to be T for all purpose of this clause.
+    auto Type = D->getType();
+    if (Type->isReferenceType())
+      Type = Type->getPointeeType();
+
+    // OpenMP 4.5 [2.10.5, target update Construct, Restrictions, p.4]
+    // A list item in a to or from clause must have a mappable type.
+    if (!CheckTypeMappable(VE->getExprLoc(), VE->getSourceRange(), *this,
+                           DSAStack, Type))
+      continue;
+
+    Vars.push_back(RE);
+    DSAStack->addExprToVarMapInfo(D, RE);
+  }
+
+  if (Vars.empty())
+    return nullptr;
+
+  return OMPFromClause::Create(Context, StartLoc, LParenLoc, EndLoc, Vars);
+}

lib/Sema/TreeTransform.h

   RebuildOMPDistScheduleClause(OpenMPDistScheduleClauseKind Kind,
                                Expr *ChunkSize, SourceLocation StartLoc,
                                SourceLocation LParenLoc, SourceLocation KindLoc,
                                SourceLocation CommaLoc, SourceLocation EndLoc) {
     return getSema().ActOnOpenMPDistScheduleClause(
         Kind, ChunkSize, StartLoc, LParenLoc, KindLoc, CommaLoc, EndLoc);
   }
 
+  /// \brief Build a new OpenMP 'to' clause.
+  ///
+  /// By default, performs semantic analysis to build the new statement.
+  /// Subclasses may override this routine to provide different behavior.
+  OMPClause *RebuildOMPToClause(ArrayRef<Expr *> VarList,
+                                SourceLocation StartLoc,
+                                SourceLocation LParenLoc,
+                                SourceLocation EndLoc) {
+    return getSema().ActOnOpenMPToClause(VarList, StartLoc, LParenLoc, EndLoc);
+  }
+
+  /// \brief Build a new OpenMP 'from' clause.
+  ///
+  /// By default, performs semantic analysis to build the new statement.
+  /// Subclasses may override this routine to provide different behavior.
+  OMPClause *RebuildOMPFromClause(ArrayRef<Expr *> VarList,
+                                  SourceLocation StartLoc,
+                                  SourceLocation LParenLoc,
+                                  SourceLocation EndLoc) {
+    return getSema().ActOnOpenMPFromClause(VarList, StartLoc, LParenLoc,
+                                           EndLoc);
+  }
+
   /// \brief Rebuild the operand to an Objective-C \@synchronized statement.
   ///
   /// By default, performs semantic analysis to build the new statement.
   /// Subclasses may override this routine to provide different behavior.
   ExprResult RebuildObjCAtSynchronizedOperand(SourceLocation atLoc,
                                               Expr *object) {
     return getSema().ActOnObjCAtSynchronizedOperand(atLoc, object);
   }
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;

 }
 
 template <typename Derived>
 OMPClause *
 TreeTransform<Derived>::TransformOMPDefaultmapClause(OMPDefaultmapClause *C) {
   return C;
 }
 
+template <typename Derived>
+OMPClause *TreeTransform<Derived>::TransformOMPToClause(OMPToClause *C) {
+  llvm::SmallVector<Expr *, 16> Vars;
+  Vars.reserve(C->varlist_size());
+  for (auto *VE : C->varlists()) {
+    ExprResult EVar = getDerived().TransformExpr(cast<Expr>(VE));
+    if (EVar.isInvalid())
+      return 0;
+    Vars.push_back(EVar.get());
+  }
+  return getDerived().RebuildOMPToClause(Vars, C->getLocStart(),
+                                         C->getLParenLoc(), C->getLocEnd());
+}
+
+template <typename Derived>
+OMPClause *TreeTransform<Derived>::TransformOMPFromClause(OMPFromClause *C) {
+  llvm::SmallVector<Expr *, 16> Vars;
+  Vars.reserve(C->varlist_size());
+  for (auto *VE : C->varlists()) {
+    ExprResult EVar = getDerived().TransformExpr(cast<Expr>(VE));
+    if (EVar.isInvalid())
+      return 0;
+    Vars.push_back(EVar.get());
+  }
+  return getDerived().RebuildOMPFromClause(Vars, C->getLocStart(),
+                                           C->getLParenLoc(), C->getLocEnd());
+}
+
 //===----------------------------------------------------------------------===//
 // Expression transformation
 //===----------------------------------------------------------------------===//
 template<typename Derived>
 ExprResult
 TreeTransform<Derived>::TransformPredefinedExpr(PredefinedExpr *E) {
   if (!E->isTypeDependent())
     return E;

lib/Serialization/ASTReaderStmt.cpp

case OMPC_hint:
     C = new (Context) OMPHintClause();
     break;
   case OMPC_dist_schedule:
     C = new (Context) OMPDistScheduleClause();
     break;
   case OMPC_defaultmap:
     C = new (Context) OMPDefaultmapClause();
     break;
+  case OMPC_to:
+    C = OMPToClause::CreateEmpty(Context, Record[Idx++]);
+    break;
+  case OMPC_from:
+    C = OMPFromClause::CreateEmpty(Context, Record[Idx++]);
+    break;
   }
   Visit(C);
   C->setLocStart(Reader->ReadSourceLocation(Record, Idx));
   C->setLocEnd(Reader->ReadSourceLocation(Record, Idx));
 
   return C;
 }
 
C->setDefaultmapKind(
        static_cast<OpenMPDefaultmapClauseKind>(Record[Idx++]));
   C->setDefaultmapModifier(
       static_cast<OpenMPDefaultmapClauseModifier>(Record[Idx++]));
   C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx));
   C->setDefaultmapModifierLoc(Reader->ReadSourceLocation(Record, Idx));
   C->setDefaultmapKindLoc(Reader->ReadSourceLocation(Record, Idx));
 }
 
+void OMPClauseReader::VisitOMPToClause(OMPToClause *C) {
+  C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx));
+  unsigned NumVars = C->varlist_size();
+  SmallVector<Expr *, 16> Vars;
+  Vars.reserve(NumVars);
+  for (unsigned i = 0; i != NumVars; ++i) {
+    Vars.push_back(Reader->Reader.ReadSubExpr());
+  }
+  C->setVarRefs(Vars);
+}
+
+void OMPClauseReader::VisitOMPFromClause(OMPFromClause *C) {
+  C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx));
+  unsigned NumVars = C->varlist_size();
+  SmallVector<Expr *, 16> Vars;
+  Vars.reserve(NumVars);
+  for (unsigned i = 0; i != NumVars; ++i) {
+    Vars.push_back(Reader->Reader.ReadSubExpr());
+  }
+  C->setVarRefs(Vars);
+}
+
 //===----------------------------------------------------------------------===//
 // OpenMP Directives.
 //===----------------------------------------------------------------------===//
 void ASTStmtReader::VisitOMPExecutableDirective(OMPExecutableDirective *E) {
   E->setLocStart(ReadSourceLocation(Record, Idx));
   E->setLocEnd(ReadSourceLocation(Record, Idx));
   OMPClauseReader ClauseReader(this, Reader.getContext(), Record, Idx);
   SmallVector<OMPClause *, 5> Clauses;

 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 OMPClauseWriter::VisitOMPDefaultmapClause(OMPDefaultmapClause *C) {
   Record.push_back(C->getDefaultmapKind());
   Record.push_back(C->getDefaultmapModifier());
   Writer->Writer.AddSourceLocation(C->getLParenLoc(), Record);
   Writer->Writer.AddSourceLocation(C->getDefaultmapModifierLoc(), Record);
   Writer->Writer.AddSourceLocation(C->getDefaultmapKindLoc(), Record);
 }
 
+void OMPClauseWriter::VisitOMPToClause(OMPToClause *C) {
+  Record.push_back(C->varlist_size());
+  Writer->Writer.AddSourceLocation(C->getLParenLoc(), Record);
+  for (auto *VE : C->varlists())
+    Writer->Writer.AddStmt(VE);
+}
+
+void OMPClauseWriter::VisitOMPFromClause(OMPFromClause *C) {
+  Record.push_back(C->varlist_size());
+  Writer->Writer.AddSourceLocation(C->getLParenLoc(), Record);
+  for (auto *VE : C->varlists())
+    Writer->Writer.AddStmt(VE);
+}
+
 //===----------------------------------------------------------------------===//
 // OpenMP Directives.
 //===----------------------------------------------------------------------===//
 void ASTStmtWriter::VisitOMPExecutableDirective(OMPExecutableDirective *E) {
   Writer.AddSourceLocation(E->getLocStart(), Record);
   Writer.AddSourceLocation(E->getLocEnd(), Record);
   OMPClauseWriter ClauseWriter(this, Record);
   for (unsigned i = 0; i < E->getNumClauses(); ++i) {
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 to(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 to(a) // 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 to(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 to(a) // 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 to(a)
+  }
 
 // 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 to(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 to(a)
+    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 to(a)
+    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 to(a)
+    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 from(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 to(a) // 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 to(a)
+  }
 
 // 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 to(a)
+    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 to(a)
+    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 to(a) // 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 to(a) // 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 from(a) // 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 from(a) // 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 to(a) // 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 to(a)
+    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 to(a) // 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 to(a)
+    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 from(a) // 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 from(a)
+    ++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 from(a) // 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 from(a)
+  }
 
 // 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 from(a)
+      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 from(a)
+    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 from(a)
+    ++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 from(a)
+    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 from(a)
+    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 from(a) // 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 from(a)
+  }
 
 // 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 from(a)
+    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 from(a) // 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 from(a) // 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 from(a) // 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 from(a) // 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 from(a)
+    ++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 from(a) // expected-error {{region cannot be nested inside 'target' region}}
+    ++a;
+  }
 }

test/OpenMP/target_map_messages.cpp

void func(int arg) {
     {}
     #pragma omp target map(arg,(c+1)[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
     {}
     #pragma omp target map(arg,a[:2],d) // expected-error {{subscripted value is not an array or pointer}}
     {}
     #pragma omp target map(arg,a,d[:2]) // expected-error {{subscripted value is not an array or pointer}}
     {}
 
-    #pragma omp target map(to:ss) // expected-error {{threadprivate variables are not allowed in map clause}}
+    #pragma omp target map(to:ss) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
     {}
 
     #pragma omp target map(to:b,e)
     {}
     #pragma omp target map(to:b,e) map(to:b) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
     {}
     #pragma omp target map(to:b[:2],e)
     {}
void SAclient(int arg) {
   #pragma omp target map(r.C) map(r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
   {}
   #pragma omp target map(r.C, r.S)  // this would be an error only caught at runtime - Sema would have to make sure there is not way for the missing data between fields to be mapped somewhere else.
   {}
   #pragma omp target map(r, r.S)  // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
   {}
   #pragma omp target map(r.C, t.C)
   {}
-  #pragma omp target map(r.A)   // expected-error {{bit fields cannot be used to specify storage in a map clause}}
+  #pragma omp target map(r.A)   // expected-error {{bit fields cannot be used to specify storage in a 'map' clause}}
   {}
   #pragma omp target map(r.Arr)
   {}
   #pragma omp target map(r.Arr[3:5])
   {}
   #pragma omp target map(r.Ptr[3:5])
   {}
   #pragma omp target map(r.ArrS[3:5].A)   // expected-error {{OpenMP array section is not allowed here}}

 #pragma omp target data map(S1) // expected-error {{'S1' does not refer to a value}}
 #pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}}
 #pragma omp target data map(ba) // expected-error 2 {{type 'S2' is not mappable to target}}
 #pragma omp target data map(ca)
 #pragma omp target data map(da)
 #pragma omp target data map(S2::S2s)
 #pragma omp target data map(S2::S2sc)
 #pragma omp target data map(e, g)
-#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in map clause}}
+#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
 #pragma omp target data map(k) map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
 #pragma omp target map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
   foo();
 #pragma omp target data map(da)
 #pragma omp target map(da[:4])
   foo();
 #pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
 #pragma omp target data map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}

 #pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}}
 #pragma omp target data map(argv[1])
 #pragma omp target data map(ba) // expected-error 2 {{type 'S2' is not mappable to target}}
 #pragma omp target data map(ca)
 #pragma omp target data map(da)
 #pragma omp target data map(S2::S2s)
 #pragma omp target data map(S2::S2sc)
 #pragma omp target data map(e, g)
-#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in map clause}}
+#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
 #pragma omp target data map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
 #pragma omp target map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
   foo();
 #pragma omp target data map(da)
 #pragma omp target map(da[:4])
   foo();
 #pragma omp target data map(k, j, l) // expected-note {{used here}}
 #pragma omp target data map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}

test/OpenMP/target_parallel_for_map_messages.cpp

 #pragma omp target parallel for map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(e, g)
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(h) // expected-error {{threadprivate variables are not allowed in map clause}}
+#pragma omp target parallel for map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(da[:4])

 #pragma omp target parallel for map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(e, g)
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(h) // expected-error {{threadprivate variables are not allowed in map clause}}
+#pragma omp target parallel for map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(da)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(da[:4])

test/OpenMP/target_parallel_map_messages.cpp

 #pragma omp target parallel map(da)
   foo();
 #pragma omp target parallel map(S2::S2s)
   foo();
 #pragma omp target parallel map(S2::S2sc)
   foo();
 #pragma omp target parallel map(e, g)
   foo();
-#pragma omp target parallel map(h) // expected-error {{threadprivate variables are not allowed in map clause}}
+#pragma omp target parallel map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
   foo();
 #pragma omp target parallel map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
   foo();
 #pragma omp target parallel map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
   foo();
 #pragma omp target parallel map(da)
   foo();
 #pragma omp target parallel map(da[:4])

 #pragma omp target parallel map(da)
   foo();
 #pragma omp target parallel map(S2::S2s)
   foo();
 #pragma omp target parallel map(S2::S2sc)
   foo();
 #pragma omp target parallel map(e, g)
   foo();
-#pragma omp target parallel map(h) // expected-error {{threadprivate variables are not allowed in map clause}}
+#pragma omp target parallel map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
   foo();
 #pragma omp target parallel map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
   foo();
 #pragma omp target parallel map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
   foo();
 #pragma omp target parallel map(da)
   foo();
 #pragma omp target parallel map(da[:4])

test/OpenMP/target_update_ast_print.cpp

+// 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
+
+void foo() {}
+
+template <class T, class U>
+T foo(T targ, U uarg) {
+  static T a;
+  U b;
+  int l;
+#pragma omp target update to(a) if(l>5) device(l)
+
+#pragma omp target update from(b) if(l<5) device(l-1)
+  return a + targ + (T)b;
+}
+// CHECK:      static int a;
+// CHECK-NEXT: float b;
+// CHECK-NEXT: int l;
+// CHECK-NEXT: #pragma omp target update to(a) if(l > 5) device(l)
+// CHECK-NEXT: #pragma omp target update from(b) if(l < 5) device(l - 1)
+// CHECK:      static char a;
+// CHECK-NEXT: float b;
+// CHECK-NEXT: int l;
+// CHECK-NEXT: #pragma omp target update to(a) if(l > 5) device(l)
+// CHECK-NEXT: #pragma omp target update from(b) if(l < 5) device(l - 1)
+// CHECK:      static T a;
+// CHECK-NEXT: U b;
+// CHECK-NEXT: int l;
+// CHECK-NEXT: #pragma omp target update to(a) if(l > 5) device(l)
+// CHECK-NEXT: #pragma omp target update from(b) if(l < 5) device(l - 1)
+
+int main(int argc, char **argv) {
+  static int a;
+  int n;
+  float f;
+
+// CHECK:      static int a;
+// CHECK-NEXT: int n;
+// CHECK-NEXT: float f;
+#pragma omp target update to(a) if(f>0.0) device(n)
+  // CHECK-NEXT: #pragma omp target update to(a) if(f > 0.) device(n)
+#pragma omp target update from(f) if(f<0.0) device(n+1)
+  // CHECK-NEXT: #pragma omp target update from(f) if(f < 0.) device(n + 1)
+  return foo(argc, f) + foo(argv[0][0], f) + a;
+}
+
+#endif

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 to(i) device // expected-error {{expected '(' after 'device'}}
+#pragma omp target update to(i) device ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(i) device () // expected-error {{expected expression}}
+#pragma omp target update to(i) device (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(i) device (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+#pragma omp target update from(i) device (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+#pragma omp target update from(i) device (argc + argc)
+#pragma omp target update from(i) device (argc), device (argc+1) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'device' clause}}
+#pragma omp target update from(i) device (S1) // expected-error {{'S1' does not refer to a value}}
+#pragma omp target update from(i) device (3.14) // expected-error 2 {{expression must have integral or unscoped enumeration type, not 'double'}}
+#pragma omp target update from(i) device (-2) // expected-error {{argument to 'device' clause must be a non-negative integer value}}
+}
+
+int main(int argc, char **argv) {
+  int j;
+#pragma omp target update to(j) device // expected-error {{expected '(' after 'device'}}
+#pragma omp target update from(j) device ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(j) device () // expected-error {{expected expression}}
+#pragma omp target update from(j) device (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(j) device (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+#pragma omp target update from(j) device (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+#pragma omp target update to(j) device (argc + argc)
+#pragma omp target update from(j) device (argc), device (argc+1) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'device' clause}}
+#pragma omp target update to(j) device (S1) // expected-error {{'S1' does not refer to a value}}
+#pragma omp target update from(j) device (-2) // expected-error {{argument to 'device' clause must be a non-negative integer value}}
+#pragma omp target update to(j) device (3.14) // expected-error {{expression must have integral or unscoped enumeration type, not 'double'}}
+
+  return tmain(argc, argv); // expected-note {{in instantiation of function template specialization 'tmain<int, char>' requested here}}
+}

test/OpenMP/target_update_from_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}}
+extern S1 a;
+class S2 {
+  mutable int a;
+public:
+  S2():a(0) { }
+  S2(S2 &s2):a(s2.a) { }
+  static float S2s; // expected-note 4 {{mappable type cannot contain static members}}
+  static const float S2sc; // expected-note 4 {{mappable type cannot contain static members}}
+};
+const float S2::S2sc = 0;
+const S2 b;
+const S2 ba[5];
+class S3 {
+  int a;
+public:
+  S3():a(0) { }
+  S3(S3 &s3):a(s3.a) { }
+};
+const S3 c;
+const S3 ca[5];
+extern const int f;
+class S4 {
+  int a;
+  S4();
+  S4(const S4 &s4);
+public:
+  S4(int v):a(v) { }
+};
+class S5 {
+  int a;
+  S5():a(0) {}
+  S5(const S5 &s5):a(s5.a) { }
+public:
+  S5(int v):a(v) { }
+};
+struct S6 {
+  int ii;
+  int aa[30];
+  float xx;
+  double *pp;
+};
+struct S7 {
+  int i;
+  int a[50];
+  float x;
+  S6 s6[5];
+  double *p;
+  unsigned bfa : 4;
+};
+
+S3 h;
+#pragma omp threadprivate(h) // expected-note 2 {{defined as threadprivate or thread local}}
+
+typedef int to;
+
+template <typename T, int I> // expected-note {{declared here}}
+T tmain(T argc) {
+  const T d = 5;
+  const T da[5] = { 0 };
+  S4 e(4);
+  S5 g(5);
+  T i, t[20];
+  T &j = i;
+  T *k = &j;
+  T x;
+  T y;
+  T from;
+  const T (&l)[5] = da;
+  T *m;
+  S7 s7;
+
+#pragma omp target update from // expected-error {{expected '(' after 'from'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from( // expected-error {{expected ')'}} expected-note {{to match this '('}} 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 from() // 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() // 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 from(alloc) // expected-error {{use of undeclared identifier 'alloc'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(x)
+#pragma omp target update from(t[:I])
+#pragma omp target update from(T) // expected-error {{'T' 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 from(I) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update from(S2::S2s)
+#pragma omp target update from(S2::S2sc)
+#pragma omp target update from(from)
+#pragma omp target update from(y x) // expected-error {{expected ',' or ')' in 'from' clause}}
+#pragma omp target update from(argc > 0 ? x : y) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} 
+#pragma omp target update from(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 from(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}}
+#pragma omp target update from(ba) // expected-error 2 {{type 'S2' is not mappable to target}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(h) // expected-error {{threadprivate variables are not allowed in 'from' clause}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(k), to(k) // expected-error 2 {{variable can appear only once in OpenMP 'target update' construct}} expected-note 2 {{used here}}
+#pragma omp target update from(t), from(t[:5]) // expected-error 2 {{variable can appear only once in OpenMP 'target update' construct}} expected-note 2 {{used here}}
+#pragma omp target update from(da)
+#pragma omp target update from(da[:4])
+
+#pragma omp target update from(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}}
+#pragma omp target update from(x, c[:]) // expected-error {{subscripted value is not an array or pointer}}
+#pragma omp target update from(x, (m+1)[2]) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update from(s7.i, s7.a[:3])
+#pragma omp target update from(s7.s6[1].aa[0:5])
+#pragma omp target update from(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target update from(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target update from(s7.p[:10])
+#pragma omp target update from(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'from' clause}}
+#pragma omp target update from(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+#pragma omp target data map(to: s7.i)
+  {
+#pragma omp target update from(s7.x)
+  }
+
+  return 0;
+}
+
+int main(int argc, char **argv) {
+  const int d = 5;
+  const int da[5] = { 0 };
+  S4 e(4);
+  S5 g(5);
+  int i, t[20];
+  int &j = i;
+  int *k = &j;
+  int x;
+  int y;
+  int from;
+  const int (&l)[5] = da;
+  int *m;
+  S7 s7;
+
+#pragma omp target update from // expected-error {{expected '(' after 'from'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from( // expected-error {{expected ')'}} expected-note {{to match this '('}} 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 from() // 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() // 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 from(alloc) // expected-error {{use of undeclared identifier 'alloc'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(x)
+#pragma omp target update from(t[:i])
+#pragma omp target update from(S2::S2s)
+#pragma omp target update from(S2::S2sc)
+#pragma omp target update from(from)
+#pragma omp target update from(y x) // expected-error {{expected ',' or ')' in 'from' clause}}
+#pragma omp target update from(argc > 0 ? x : y) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(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 from(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}}
+#pragma omp target update from(ba) // expected-error 2 {{type 'S2' is not mappable to target}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(h) // expected-error {{threadprivate variables are not allowed in 'from' clause}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(k), to(k) // expected-error {{variable can appear only once in OpenMP 'target update' construct}} expected-note {{used here}}
+#pragma omp target update from(t), from(t[:5]) // expected-error {{variable can appear only once in OpenMP 'target update' construct}} expected-note {{used here}}
+#pragma omp target update from(da)
+#pragma omp target update from(da[:4])
+
+#pragma omp target update from(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}}
+#pragma omp target update from(x, c[:]) // expected-error {{subscripted value is not an array or pointer}}
+#pragma omp target update from(x, (m+1)[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update from(s7.i, s7.a[:3])
+#pragma omp target update from(s7.s6[1].aa[0:5])
+#pragma omp target update from(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target update from(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target update from(s7.p[:10])
+#pragma omp target update from(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'from' clause}}
+#pragma omp target update from(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+#pragma omp target data map(to: s7.i)
+  {
+#pragma omp target update from(s7.x)
+  }
+
+  return tmain<int, 3>(argc)+tmain<to, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
+}
+

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 to(n) if // expected-error {{expected '(' after 'if'}}
+#pragma omp target update from(n) if ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(n) if () // expected-error {{expected expression}}
+#pragma omp target update from(n) if (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(n) if (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+#pragma omp target update from(n) if (argc > 0 ? argv[1] : argv[2])
+#pragma omp target update to(n) if (foobool(argc)), if (true) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'if' clause}}
+#pragma omp target update from(n) if (S) // expected-error {{'S' does not refer to a value}}
+#pragma omp target update to(n) if (argv[1]=2) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update from(n) if (argc argc) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(n) if(argc)
+#pragma omp target update from(n) if(target update // expected-warning {{missing ':' after directive name modifier - ignoring}} expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(n) if(target update : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update from(n) if(target update : argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(n) if(target update : argc)
+#pragma omp target update from(n) if(target update : argc) if (for:argc) // expected-error {{directive name modifier 'for' is not allowed for '#pragma omp target update'}}
+#pragma omp target update to(n) 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}}
+#pragma omp target update from(n) if(target update : argc) if (argc) // expected-error {{no more 'if' clause is allowed}} expected-note {{previous clause with directive name modifier specified here}}
+  return 0;
+}
+
+int main(int argc, char **argv) {
+  int m;
+#pragma omp target update to(m) if // expected-error {{expected '(' after 'if'}}
+#pragma omp target update from(m) if ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(m) if () // expected-error {{expected expression}}
+#pragma omp target update from(m) if (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(m) if (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+#pragma omp target update from(m) if (argc > 0 ? argv[1] : argv[2])
+#pragma omp target update to(m) if (foobool(argc)), if (true) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'if' clause}}
+#pragma omp target update from(m) if (S1) // expected-error {{'S1' does not refer to a value}}
+#pragma omp target update to(m) if (argv[1]=2) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update from(m) if (argc argc) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(m) if (1 0) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update from(m) if(if(tmain(argc, argv) // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(m) if(target update // expected-warning {{missing ':' after directive name modifier - ignoring}} expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update from(m) if(target update : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update to(m) if(target update : argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target update from(m) if(target update : argc)
+#pragma omp target update to(m) if(target update : argc) if (for:argc) // expected-error {{directive name modifier 'for' is not allowed for '#pragma omp target update'}}
+#pragma omp target update from(m) 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}}
+#pragma omp target update to(m) if(target update : argc) if (argc) // expected-error {{no more 'if' clause is allowed}} expected-note {{previous clause with directive name modifier specified here}}
+  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 to(m) { // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+  #pragma omp target update to(m) ( // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+  #pragma omp target update to(m) [ // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+  #pragma omp target update to(m) ] // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+  #pragma omp target update to(m) ) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}}
+
+  #pragma omp target update from(m) // OK
+  {
+    foo();
+  }
+  return tmain(argc, argv);
+}

test/OpenMP/target_update_to_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}}
+extern S1 a;
+class S2 {
+  mutable int a;
+public:
+  S2():a(0) { }
+  S2(S2 &s2):a(s2.a) { }
+  static float S2s; // expected-note 4 {{mappable type cannot contain static members}}
+  static const float S2sc; // expected-note 4 {{mappable type cannot contain static members}}
+};
+const float S2::S2sc = 0;
+const S2 b;
+const S2 ba[5];
+class S3 {
+  int a;
+public:
+  S3():a(0) { }
+  S3(S3 &s3):a(s3.a) { }
+};
+const S3 c;
+const S3 ca[5];
+extern const int f;
+class S4 {
+  int a;
+  S4();
+  S4(const S4 &s4);
+public:
+  S4(int v):a(v) { }
+};
+class S5 {
+  int a;
+  S5():a(0) {}
+  S5(const S5 &s5):a(s5.a) { }
+public:
+  S5(int v):a(v) { }
+};
+struct S6 {
+  int ii;
+  int aa[30];
+  float xx;
+  double *pp;
+};
+struct S7 {
+  int i;
+  int a[50];
+  float x;
+  S6 s6[5];
+  double *p;
+  unsigned bfa : 4;
+};
+
+S3 h;
+#pragma omp threadprivate(h) // expected-note 2 {{defined as threadprivate or thread local}}
+
+typedef int from;
+
+template <typename T, int I> // expected-note {{declared here}}
+T tmain(T argc) {
+  const T d = 5;
+  const T da[5] = { 0 };
+  S4 e(4);
+  S5 g(5);
+  T *m;
+  T i, t[20];
+  T &j = i;
+  T *k = &j;
+  T x;
+  T y;
+  T to;
+  const T (&l)[5] = da;
+  S7 s7;
+
+#pragma omp target update to // expected-error {{expected '(' after 'to'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to( // expected-error {{expected ')'}} expected-note {{to match this '('}} 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 to() // 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() // 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 to(alloc) // expected-error {{use of undeclared identifier 'alloc'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(x)
+#pragma omp target update to(t[:I])
+#pragma omp target update to(T) // expected-error {{'T' 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 to(I) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update to(S2::S2s)
+#pragma omp target update to(S2::S2sc)
+#pragma omp target update to(to)
+#pragma omp target update to(y x) // expected-error {{expected ',' or ')' in 'to' clause}}
+#pragma omp target update to(argc > 0 ? x : y) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} 
+#pragma omp target update to(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 to(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}}
+#pragma omp target update to(ba) // expected-error 2 {{type 'S2' is not mappable to target}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(h) // expected-error {{threadprivate variables are not allowed in 'to' clause}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(k), from(k) // expected-error 2 {{variable can appear only once in OpenMP 'target update' construct}} expected-note 2 {{used here}}
+#pragma omp target update to(t), to(t[:5]) // expected-error 2 {{variable can appear only once in OpenMP 'target update' construct}} expected-note 2 {{used here}}
+#pragma omp target update to(da)
+#pragma omp target update to(da[:4])
+
+#pragma omp target update to(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}}
+#pragma omp target update to(x, c[:]) // expected-error {{subscripted value is not an array or pointer}}
+#pragma omp target update to(x, (m+1)[2]) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update to(s7.i, s7.a[:3])
+#pragma omp target update to(s7.s6[1].aa[0:5])
+#pragma omp target update to(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target update to(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target update to(s7.p[:10])
+#pragma omp target update to(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'to' clause}}
+#pragma omp target update to(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+#pragma omp target data map(to: s7.i)
+  {
+#pragma omp target update to(s7.x)
+  }
+  return 0;
+}
+
+int main(int argc, char **argv) {
+  const int d = 5;
+  const int da[5] = { 0 };
+  S4 e(4);
+  S5 g(5);
+  int i, t[20];
+  int &j = i;
+  int *k = &j;
+  int x;
+  int y;
+  int to;
+  const int (&l)[5] = da;
+  S7 s7;
+  int *m;
+
+#pragma omp target update to // expected-error {{expected '(' after 'to'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to( // expected-error {{expected ')'}} expected-note {{to match this '('}} 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 to() // 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() // 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 to(alloc) // expected-error {{use of undeclared identifier 'alloc'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(x)
+#pragma omp target update to(t[:i])
+#pragma omp target update to(S2::S2s)
+#pragma omp target update to(S2::S2sc)
+#pragma omp target update to(to)
+#pragma omp target update to(y x) // expected-error {{expected ',' or ')' in 'to' clause}}
+#pragma omp target update to(argc > 0 ? x : y) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(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 to(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}}
+#pragma omp target update to(ba) // expected-error 2 {{type 'S2' is not mappable to target}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(h) // expected-error {{threadprivate variables are not allowed in 'to' clause}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(k), from(k) // expected-error {{variable can appear only once in OpenMP 'target update' construct}} expected-note {{used here}}
+#pragma omp target update to(t), to(t[:5]) // expected-error {{variable can appear only once in OpenMP 'target update' construct}} expected-note {{used here}}
+#pragma omp target update to(da)
+#pragma omp target update to(da[:4])
+
+#pragma omp target update to(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}}
+#pragma omp target update to(x, c[:]) // expected-error {{subscripted value is not an array or pointer}}
+#pragma omp target update to(x, (m+1)[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update to(s7.i, s7.a[:3])
+#pragma omp target update to(s7.s6[1].aa[0:5])
+#pragma omp target update to(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target update to(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}}
+#pragma omp target update to(s7.p[:10])
+#pragma omp target update to(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'to' clause}}
+#pragma omp target update to(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+#pragma omp target data map(to: s7.i)
+  {
+#pragma omp target update to(s7.x)
+  }
+
+  return tmain<int, 3>(argc)+tmain<from, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
+}
+

tools/libclang/CIndex.cpp

 }
 void OMPClauseEnqueue::VisitOMPDistScheduleClause(
     const OMPDistScheduleClause *C) {
   VisitOMPClauseWithPreInit(C);
   Visitor->AddStmt(C->getChunkSize());
 }
 void OMPClauseEnqueue::VisitOMPDefaultmapClause(
     const OMPDefaultmapClause * /*C*/) {}
+void OMPClauseEnqueue::VisitOMPFromClause(const OMPFromClause *C) {
+  VisitOMPClauseList(C);
+}
+void OMPClauseEnqueue::VisitOMPToClause(const OMPToClause *C) {
+  VisitOMPClauseList(C);
+}
 }
 
 void EnqueueVisitor::EnqueueChildren(const OMPClause *S) {
   unsigned size = WL.size();
   OMPClauseEnqueue Visitor(this);
   Visitor.Visit(S);
   if (size == WL.size())
     return;
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;