[OpenMP] Parsing and sema support for the to clause

[sfantao]

Summary

The patch contains the parsing and sema support for the to clause.

Patch based on the original post by Kelvin Li.

[ABataev]

lib/Sema/SemaOpenMP.cpp

9765–9766

I believe this check is not required anymore, because later you check that VE is not dependent.

9823

Why did you remove this line and inserted an empty one?

lib/Serialization/ASTReaderStmt.cpp

2294–2296

Remove extra braces

[kkwli0]

Thanks for the review.

lib/Sema/SemaOpenMP.cpp

9823

Will put it back.

[kkwli0]

Update the patch based on the review comments.

[ABataev]

lib/Sema/SemaOpenMP.cpp

9758–9764

General problem of all map/to/from clauses - they don't know anything about captured decls.
If we have the next code, it will crash the compiler or lead to incorrect code:

struct S {
  int a;
  S() {
   #pragma omp parallel private(a)
   {
    #pragma omp target map(a:....)
    ...
    }
};

The same problem wil be with the to/from clauses. Try to write codegen/dump tests for this and check that private copy of a member decl is used

[sfantao]

Hi Alexey,

I have inlined comment.

Thanks!
Samuel

lib/Sema/SemaOpenMP.cpp

9758–9764

The example you just gave will work just fine for the same reason I gave in http://reviews.llvm.org/D16749. So, this does not pose any problem to map/to/from. Privatization is now changing the outline region signature (we now have the extra privatized fields), and the map codegen (as well as to/from) will attempt to emit the reference to the variable, whatever its local address currently is. So, it is is a private address that's the one that will be used, if it is something in this the right field would be used. Given that this is also implicitly mapped, the rules of default map will be used and the this pointee mapped as a to/from-aggregate along with all the individual declarations that have been privatized.

Map clause does not change the code emission inside the outlined region - it only matters to get the right information to pass the runtime library (and the eventual capture by copy). So, if something works in the generation of the region without the map clause (implicit map), there is no reason why it shouldn't when a map clause is used.

For 'to'/'from' is even more irrelevant because no outlined regions are created. Again, only the information passed to the runtime library is different.

In any case, I will add your example to the regression tests in http://reviews.llvm.org/D16749 just to add some peace of mind.

[ABataev]

lib/Sema/SemaOpenMP.cpp

9758–9764

Add tests, please, for all possible situations with the data members

[kkwli0]

lib/Sema/SemaOpenMP.cpp

9758–9764

I will add the test cases under http://reviews.llvm.org/D15944 since the clauses rely on the support of target update construct. Is it okay?

[sfantao]

Add implementation and tests only for the to clause. The implementation is based on the existing infrastructure used by the map clause already available upstream.

[ABataev]

lib/Sema/SemaOpenMP.cpp

9764–9766

1. Use SmallVectorImpl<Expr *> instead of SmallVector<Expr *, n>.
2. Is it possible to reduce number of arguments of this function by gathering them into a record?

9765–9766

Still think that this check is not required.

9877

You already get DKind few lines above, why need to update it?

test/OpenMP/nesting_of_regions.cpp

136 ↗

(On Diff #58370)

Test 'nesting_of_regions.cpp' should be updated only when adding a new directive, not a clause.

[sfantao]

• Address comments from the last review by Alexey.

[sfantao]

Hi Alexey,

Thanks for the review.

lib/Sema/SemaOpenMP.cpp

9758–9765

Just noticed these things were unanswered. This comment was prior to the map clause patch, so now all data members cases are tested in the upstreamed regression tests.

9764–9766

Ok, I created a new container that I named MappableVarListInfo to keep all the required lists. I now pass that container instead of the four lists. Hope this is more or less what you had in mind.

9765–9766

Correct, this is not needed. Sorry, forgot to remove that in the previous diff.

9877

Right, I don't. I removed the latest diff.

test/OpenMP/nesting_of_regions.cpp

136 ↗

(On Diff #58370)

Ok, I added the nesting tests to the dependence patch http://reviews.llvm.org/D15944.

I only do a few changes in this patch to use a clause because that allows me to remove the expected errors "need to/from clause" that are unrelated with nesting. These errors are already checked by the other tests.

[ABataev]

LG, with a small nit.

lib/Sema/SemaOpenMP.cpp

9760

Mark it as a 'final' struct

[sfantao]

• Mark MappableVarListInfo as final.

include/clang/AST/OpenMPClause.h

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 '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

 }
 
 template <typename Derived>
 bool
 RecursiveASTVisitor<Derived>::VisitOMPDefaultmapClause(OMPDefaultmapClause *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/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_array_section_does_not_specify_contiguous_storage : Error<
   "array section does not specify contiguous storage">;
 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_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_variable_in_map_and_dsa : Error<
   "%0 variable cannot be in a map clause in '#pragma omp %1' 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

 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)
 
 // 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)

include/clang/Sema/Sema.h

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 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.

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

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::VisitOMPDistScheduleClause(OMPDistScheduleClause *Node) {
   OS << "dist_schedule(" << getOpenMPSimpleClauseTypeName(
                            OMPC_dist_schedule, Node->getDistScheduleKind());
   if (auto *E = Node->getChunkSize()) {
     OS << ", ";
     E->printPretty(OS, nullptr, Policy);
   }
   OS << ")";

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

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:
     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:
     break;
   }
   llvm_unreachable("Invalid OpenMP simple clause kind");
 }
 
 bool clang::isAllowedClauseForDirective(OpenMPDirectiveKind DKind,
                                         OpenMPClauseKind CKind) {
   assert(DKind <= OMPD_unknown);

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:
     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()) {

lib/Parse/ParseOpenMP.cpp

 ///       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
 ///
 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:
     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 clause
 ///
 ///    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 ')'
 ///
 /// 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

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_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_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_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_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_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 (!Type->isAnyPointerType() && Type->isVariablyModifiedType() &&
           << D;
       continue;
     }
 
     // OpenMP 4.5 [2.15.5.1, Restrictions, p.3]
     // A list item cannot appear in both a map clause and a data-sharing
     // attribute clause on the same construct
     if (DSAStack->getCurrentDirective() == OMPD_target) {
-      if(DSAStack->checkMapInfoForVar(VD, /* CurrentRegionOnly = */ true,
-                                      [&](Expr *RE) -> bool {return true;})) {
+      if (DSAStack->checkMapInfoForVar(
+              VD, /* CurrentRegionOnly = */ true,
         Diag(ELoc, diag::err_omp_variable_in_map_and_dsa)
             << getOpenMPClauseName(OMPC_private)
             << getOpenMPDirectiveName(DSAStack->getCurrentDirective());
         ReportOriginalDSA(*this, DSAStack, D, DVar);
         continue;
       }
     }
 
if (!IsImplicitClause) {
           ReportOriginalDSA(*this, DSAStack, D, DVar);
           continue;
         }
       }
       // OpenMP 4.5 [2.15.5.1, Restrictions, p.3]
       // A list item cannot appear in both a map clause and a data-sharing
       // attribute clause on the same construct
       if (CurrDir == OMPD_target) {
-        if(DSAStack->checkMapInfoForVar(VD, /* CurrentRegionOnly = */ true,
-                                        [&](Expr *RE) -> bool {return true;})) {
+        if (DSAStack->checkMapInfoForVar(
+                VD, /* CurrentRegionOnly = */ true,
           Diag(ELoc, diag::err_omp_variable_in_map_and_dsa)
               << getOpenMPClauseName(OMPC_firstprivate)
               << getOpenMPDirectiveName(DSAStack->getCurrentDirective());
           ReportOriginalDSA(*this, DSAStack, D, DVar);
           continue;
         }
       }
     }
if (!Length->EvaluateAsInt(ConstLength, SemaRef.getASTContext()))
     return false; // Can't get the integer value as a constant.
 
   return ConstLength.getSExtValue() != 1;
 }
 
 // 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().getNonReferenceType();
 
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 {
+              assert(CKind == OMPC_to);
+              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 {
+            assert(CKind == OMPC_to);
+            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)
if (EnclosingExpr && !IsEnclosedByDataEnvironmentExpr) {
     SemaRef.Diag(EnclosingExpr->getExprLoc(), diag::note_used_here)
         << EnclosingExpr->getSourceRange();
     return true;
   }
 
   return FoundError;
 }
 
-OMPClause *
-Sema::ActOnOpenMPMapClause(OpenMPMapClauseKind MapTypeModifier,
-                           OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
-                           SourceLocation MapLoc, SourceLocation ColonLoc,
-                           ArrayRef<Expr *> VarList, SourceLocation StartLoc,
-                           SourceLocation LParenLoc, SourceLocation EndLoc) {
-  SmallVector<Expr *, 4> Vars;
-
+static void checkOpenMPToFromMapVars(Sema &SemaRef, DSAStackTy *DSAS,
+                                     ArrayRef<Expr *> VarList,
+                                     SmallVector<Expr *, 4> &Vars,
+                                     SourceLocation StartLoc,
+                                     OpenMPClauseKind CKind,
+                                     OpenMPMapClauseKind MapType,
+                                     bool IsMapTypeImplicit) {
   for (auto &RE : VarList) {
-    assert(RE && "Null expr in omp map");
-    if (isa<DependentScopeDeclRefExpr>(RE)) {
-      // It will be analyzed later.
-      Vars.push_back(RE);
-      continue;
-    }
-    SourceLocation ELoc = RE->getExprLoc();
-
+    assert(RE && "Null expr in omp to/map clause");
+    auto 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)
+      SemaRef.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(SemaRef, SimpleExpr, CKind);
     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.
+    // 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 map clause.");
+    assert(D && "Null decl on to/map clause.");
 
     auto *VD = dyn_cast<VarDecl>(D);
     auto *FD = dyn_cast<FieldDecl>(D);
 
     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);
-      ReportOriginalDSA(*this, DSAStack, VD, DVar);
+    // threadprivate variables cannot appear in a map clause.
+    // OpenMP 4.5 [2.10.5, target update Construct]
+    // threadprivate variables cannot appear in a to clause.
+    if (VD && DSAS->isThreadPrivate(VD)) {
+      auto DVar = DSAS->getTopDSA(VD, false);
+      SemaRef.Diag(ELoc, diag::err_omp_threadprivate_in_clause)
+          << getOpenMPClauseName(CKind);
+      ReportOriginalDSA(SemaRef, DSAS, 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))
-      break;
-    if (CheckMapConflicts(*this, DSAStack, D, SimpleExpr,
-                          /*CurrentRegionOnly=*/false))
+    // Check conflicts with other map or to clause expressions. We check
+    // the conflicts with the current construct separately from the enclosing
+    // data environment, because the restrictions are different.
+    if (CheckMapConflicts(SemaRef, DSAS, D, SimpleExpr,
+                          /*CurrentRegionOnly=*/true, CKind))
+      break;
+    if (CKind == OMPC_map &&
+        CheckMapConflicts(SemaRef, DSAS, D, SimpleExpr,
+                          /*CurrentRegionOnly=*/false, CKind))
       break;
 
+    // OpenMP 4.5 [2.10.5, target update Construct]
     // 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 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.
     // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.9]
-    //  A list item must have a mappable type.
-    if (!CheckTypeMappable(VE->getExprLoc(), VE->getSourceRange(), *this,
-                           DSAStack, Type))
+    // A list item must have a mappable type.
+    if (!CheckTypeMappable(VE->getExprLoc(), VE->getSourceRange(), SemaRef,
+                           DSAS, Type))
       continue;
 
+    if (CKind == OMPC_map) {
       // target enter data
       // OpenMP [2.10.2, Restrictions, p. 99]
       // A map-type must be specified in all map clauses and must be either
       // to or alloc.
-    OpenMPDirectiveKind DKind = DSAStack->getCurrentDirective();
+      OpenMPDirectiveKind DKind = DSAS->getCurrentDirective();
       if (DKind == OMPD_target_enter_data &&
           !(MapType == OMPC_MAP_to || MapType == OMPC_MAP_alloc)) {
-      Diag(StartLoc, diag::err_omp_invalid_map_type_for_directive)
+        SemaRef.Diag(StartLoc, diag::err_omp_invalid_map_type_for_directive)
             << (IsMapTypeImplicit ? 1 : 0)
             << getOpenMPSimpleClauseTypeName(OMPC_map, MapType)
             << getOpenMPDirectiveName(DKind);
         continue;
       }
 
       // target exit_data
       // OpenMP [2.10.3, Restrictions, p. 102]
       // A map-type must be specified in all map clauses and must be either
       // from, release, or delete.
-    DKind = DSAStack->getCurrentDirective();
+      DKind = DSAS->getCurrentDirective();
       if (DKind == OMPD_target_exit_data &&
           !(MapType == OMPC_MAP_from || MapType == OMPC_MAP_release ||
             MapType == OMPC_MAP_delete)) {
-      Diag(StartLoc, diag::err_omp_invalid_map_type_for_directive)
+        SemaRef.Diag(StartLoc, diag::err_omp_invalid_map_type_for_directive)
             << (IsMapTypeImplicit ? 1 : 0)
             << getOpenMPSimpleClauseTypeName(OMPC_map, MapType)
             << getOpenMPDirectiveName(DKind);
         continue;
       }
 
       // OpenMP 4.5 [2.15.5.1, Restrictions, p.3]
       // A list item cannot appear in both a map clause and a data-sharing
       // attribute clause on the same construct
       if (DKind == OMPD_target && VD) {
-      auto DVar = DSAStack->getTopDSA(VD, false);
+        auto DVar = DSAS->getTopDSA(VD, false);
         if (isOpenMPPrivate(DVar.CKind)) {
-        Diag(ELoc, diag::err_omp_variable_in_map_and_dsa)
+          SemaRef.Diag(ELoc, diag::err_omp_variable_in_map_and_dsa)
               << getOpenMPClauseName(DVar.CKind)
-            << getOpenMPDirectiveName(DSAStack->getCurrentDirective());
-        ReportOriginalDSA(*this, DSAStack, D, DVar);
+              << getOpenMPDirectiveName(DSAS->getCurrentDirective());
+          ReportOriginalDSA(SemaRef, DSAS, D, DVar);
           continue;
         }
       }
+    }
 
     Vars.push_back(RE);
-    DSAStack->addExprToVarMapInfo(D, RE);
-  }
+    DSAS->addExprToVarMapInfo(D, RE);
+  }
 
+OMPClause *
+Sema::ActOnOpenMPMapClause(OpenMPMapClauseKind MapTypeModifier,
+                           OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
+                           SourceLocation MapLoc, SourceLocation ColonLoc,
+                           ArrayRef<Expr *> VarList, SourceLocation StartLoc,
+                           SourceLocation LParenLoc, SourceLocation EndLoc) {
+  SmallVector<Expr *, 4> Vars;
+  checkOpenMPToFromMapVars(*this, DSAStack, VarList, Vars, StartLoc, OMPC_map,
+                           MapType, IsMapTypeImplicit);
+
   // We need to produce a map clause even if we don't have variables so that
   // other diagnostics related with non-existing map clauses are accurate.
   return OMPMapClause::Create(Context, StartLoc, LParenLoc, EndLoc, Vars,
                               MapTypeModifier, MapType, IsMapTypeImplicit,
                               MapLoc);
 }
 
 QualType Sema::ActOnOpenMPDeclareReductionType(SourceLocation TyLoc,
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;
+  checkOpenMPToFromMapVars(*this, DSAStack, VarList, Vars, StartLoc, OMPC_to,
+                           OMPC_MAP_unknown, /*IsMapTypeImplicit=*/false);
+  if (Vars.empty())
+    return nullptr;
+
+  return OMPToClause::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 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);
   }

 }
 
 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());
+}
+
 //===----------------------------------------------------------------------===//
 // 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;
   }
   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);
+}
+
 //===----------------------------------------------------------------------===//
 // 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;

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);
+}
+
 //===----------------------------------------------------------------------===//
 // 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) {

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

 #pragma omp target data map(always, tofrom: x)
 #pragma omp target data map(always: x) // expected-error {{missing map type}}
 #pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
 #pragma omp target data map(always, tofrom: always, tofrom, x)
 #pragma omp target map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
 #pragma omp target private(j) map(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target' directive}}  expected-note {{defined as private}}
   {}
-#pragma omp target firstprivate(j) map(j)  // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target' directive}} expected-note {{defined as firstprivate}}
+#pragma omp target map(j) private(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target' directive}}
+  {}
+#pragma omp target firstprivate(j) map(j) // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target' directive}} expected-note {{defined as firstprivate}}
   {}
+#pragma omp target map(j) firstprivate(j) // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target' directive}}
+  {}
   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}}
 }
 #endif

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])

tools/libclang/CIndex.cpp

 }
 void OMPClauseEnqueue::VisitOMPDistScheduleClause(
     const OMPDistScheduleClause *C) {
   VisitOMPClauseWithPreInit(C);
   Visitor->AddStmt(C->getChunkSize());
 }
 void OMPClauseEnqueue::VisitOMPDefaultmapClause(
     const OMPDefaultmapClause * /*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;