[OpenMP] Parsing and sema support for map clause

[kkwli0]

Summary

This patch is to add parsing and sema support for map clause. This includes the new map types and the map type modifier added in OpenMP 4.5.

[kkwli0]

Ping

[ABataev]

include/clang/AST/OpenMPClause.h

2628

I don't like the idea of friend Sema. If Sema needs some interfaces, they must be public.

include/clang/Basic/OpenMPKinds.h

83

This must be the last element in enum.

87

Remove this one, it is not required

include/clang/Sema/Sema.h

8035–8036

I think DepLinLoc and MapLoc can joined into one argument, like DepLinMapLoc

lib/AST/StmtPrinter.cpp

853–854

I think MapType can be optional, just like MapTypeModifier

856–858

Better to use ranged for, if possible

lib/Parse/ParseOpenMP.cpp

776–781

These two can be joined

lib/Sema/SemaOpenMP.cpp

7481–7483

Use ranged loop and do not use noload version.

7502–7504

Ranged loop

7506–7508

Braces are not needed

7521–7523

Extra braces

7635–7637

Extra braces

test/OpenMP/target_map_messages.cpp

2

Tests for templates, serialization/deserialization, AST printing, region nesting

[kkwli0]

Update the patch after the first review.

[ABataev]

include/clang/AST/OpenMPClause.h

2660–2662

Not all important arguments are described

2683–2686

Not all important arguments are described

lib/Basic/OpenMPKinds.cpp

199

Extra spaces between arg and ')'

lib/Sema/SemaOpenMP.cpp

318–342

Tmp is bad name.

test/OpenMP/target_data_ast_print.cpp

1

Also add tests for templates, serialization/deserialization

test/OpenMP/target_map_messages.cpp

3

Still no tests for templates

[kkwli0]

Address the comments and will post an updated patch.

[kkwli0]

Updated patch with the 2nd review comments addressed.

[ABataev]

LG

[kkwli0]

Committed revision 253849.
Committed revision 253850.

include/clang/AST/DataRecursiveASTVisitor.h

 }
 
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::VisitOMPDeviceClause(OMPDeviceClause *C) {
   TRY_TO(TraverseStmt(C->getDevice()));
   return true;
 }
 
+template <typename Derived>
+bool RecursiveASTVisitor<Derived>::VisitOMPMapClause(OMPMapClause *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/OpenMPClause.h

static bool classof(const OMPClause *T) {
     return T->getClauseKind() == OMPC_simd;
   }
 
   child_range children() {
     return child_range(child_iterator(), child_iterator());
   }
 };
 
+/// \brief This represents clause 'map' in the '#pragma omp ...'
+/// directives.
+///
+/// \code
+/// #pragma omp target map(a,b)
+/// \endcode
+/// In this example directive '#pragma omp target' has clause 'map'
+/// with the variables 'a' and 'b'.
+///
+class OMPMapClause : public OMPVarListClause<OMPMapClause> {
+  friend class OMPClauseReader;
+
+  /// \brief Map type modifier for the 'map' clause.
+  OpenMPMapClauseKind MapTypeModifier;
+  /// \brief Map type for the 'map' clause.
+  OpenMPMapClauseKind MapType;
+  /// \brief Location of the map type.
+  SourceLocation MapLoc;
+  /// \brief Colon location.
+  SourceLocation ColonLoc;
+
+  /// \brief Set type modifier for the clause.
+  ///
+  /// \param T Type Modifier for the clause.
+  ///
+  void setMapTypeModifier(OpenMPMapClauseKind T) { MapTypeModifier = T; }
+
+  /// \brief Set type for the clause.
+  ///
+  /// \param T Type for the clause.
+  ///
+  void setMapType(OpenMPMapClauseKind T) { MapType = T; }
+
+  /// \brief Set type location.
+  ///
+  /// \param TLoc Type location.
+  ///
+  void setMapLoc(SourceLocation TLoc) { MapLoc = TLoc; }
+
+  /// \brief Set colon location.
+  void setColonLoc(SourceLocation Loc) { ColonLoc = Loc; }
+
+  /// \brief Build clause with number of variables \a N.
+  ///
+  /// \param MayTypeModifier Map type modifier.
+  /// \param MapType Map type.
+  /// \param MapLoc Location of the map type.
+  /// \param StartLoc Starting location of the clause.
+  /// \param EndLoc Ending location of the clause.
+  /// \param N Number of the variables in the clause.
+  ///
+  explicit OMPMapClause(OpenMPMapClauseKind MapTypeModifier,
+                        OpenMPMapClauseKind MapType, SourceLocation MapLoc,
+                        SourceLocation StartLoc, SourceLocation LParenLoc,
+                        SourceLocation EndLoc, unsigned N)
+    : OMPVarListClause<OMPMapClause>(OMPC_map, StartLoc, LParenLoc, EndLoc, N),
+      MapTypeModifier(MapTypeModifier), MapType(MapType), MapLoc(MapLoc) {}
+
+  /// \brief Build an empty clause.
+  ///
+  /// \param N Number of variables.
+  ///
+  explicit OMPMapClause(unsigned N)
+      : OMPVarListClause<OMPMapClause>(OMPC_map, SourceLocation(),
+                                       SourceLocation(), SourceLocation(), N),
+        MapTypeModifier(OMPC_MAP_unknown), MapType(OMPC_MAP_unknown), MapLoc() {}
+
+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.
+  /// \param TypeModifier Map type modifier.
+  /// \param Type Map type.
+  /// \param TypeLoc Location of the map type.
+  ///
+  static OMPMapClause *Create(const ASTContext &C, SourceLocation StartLoc,
+                              SourceLocation LParenLoc,
+                              SourceLocation EndLoc, ArrayRef<Expr *> VL,
+                              OpenMPMapClauseKind TypeModifier,
+                              OpenMPMapClauseKind Type, SourceLocation TypeLoc);
+  /// \brief Creates an empty clause with the place for \a N variables.
+  ///
+  /// \param C AST context.
+  /// \param N The number of variables.
+  ///
+  static OMPMapClause *CreateEmpty(const ASTContext &C, unsigned N);
+
+  /// \brief Fetches mapping kind for the clause.
+  OpenMPMapClauseKind getMapType() const LLVM_READONLY { return MapType; }
+
+  /// \brief Fetches the map type modifier for the clause.
+  OpenMPMapClauseKind getMapTypeModifier() const LLVM_READONLY {
+    return MapTypeModifier;
+  }
+
+  /// \brief Fetches location of clause mapping kind.
+  SourceLocation getMapLoc() const LLVM_READONLY { return MapLoc; }
+
+  /// \brief Get colon location.
+  SourceLocation getColonLoc() const { return ColonLoc; }
+
+  static bool classof(const OMPClause *T) {
+    return T->getClauseKind() == OMPC_map;
+  }
+
+  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>::VisitOMPDeviceClause(OMPDeviceClause *C) {
   TRY_TO(TraverseStmt(C->getDevice()));
   return true;
 }
 
+template <typename Derived>
+bool RecursiveASTVisitor<Derived>::VisitOMPMapClause(OMPMapClause *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/DiagnosticParseKinds.td

 def err_omp_expected_punc : Error<
   "expected ',' or ')' in '%0' %select{clause|directive}1">;
 def err_omp_unexpected_clause : Error<
   "unexpected OpenMP clause '%0' in directive '#pragma omp %1'">;
 def err_omp_immediate_directive : Error<
   "'#pragma omp %0' cannot be an immediate substatement">;
 def err_omp_expected_identifier_for_critical : Error<
   "expected identifier specifying the name of the 'omp critical' directive">;
+def err_omp_unknown_map_type : Error<
+  "incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'">;
+def err_omp_unknown_map_type_modifier : Error<
+  "incorrect map type modifier, expected 'always'">;
+def err_omp_map_type_missing : Error<
+  "missing map type">;
 
 // Pragma loop support.
 def err_pragma_loop_missing_argument : Error<
   "missing argument; expected %select{an integer value|"
   "'enable', %select{'assume_safety'|'full'}1 or 'disable'}0">;
 def err_pragma_loop_invalid_option : Error<
   "%select{invalid|missing}0 option%select{ %1|}0; expected vectorize, "
   "vectorize_width, interleave, interleave_count, unroll, or unroll_count">;

include/clang/Basic/DiagnosticSemaKinds.td

 def note_omp_previous_named_if_clause : Note<
   "previous clause with directive name modifier specified here">;
 def err_omp_ordered_directive_with_param : Error<
   "'ordered' directive %select{without any clauses|with 'threads' clause}0 cannot be closely nested inside ordered region with specified parameter">;
 def note_omp_ordered_param : Note<
   "'ordered' clause with specified parameter">;
 def err_omp_expected_base_var_name : Error<
   "expected variable name as a base of the array %select{subscript|section}0">;
+def err_omp_map_shared_storage : Error<
+  "variable already marked as mapped in current construct">;
+def err_omp_not_mappable_type : Error<
+  "type %0 is not mappable to target">;
+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">;
 } // 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.h

 /// \brief OpenMP attributes for 'linear' clause.
 enum OpenMPLinearClauseKind {
 #define OPENMP_LINEAR_KIND(Name) \
   OMPC_LINEAR_##Name,
 #include "clang/Basic/OpenMPKinds.def"
   OMPC_LINEAR_unknown
 };
 
+/// \brief OpenMP mapping kind for 'map' clause.
+enum OpenMPMapClauseKind {
+#define OPENMP_MAP_KIND(Name) \
+  OMPC_MAP_##Name,
+#include "clang/Basic/OpenMPKinds.def"
+  OMPC_MAP_unknown
+};
+
 OpenMPDirectiveKind getOpenMPDirectiveKind(llvm::StringRef Str);
 const char *getOpenMPDirectiveName(OpenMPDirectiveKind Kind);
 
 OpenMPClauseKind getOpenMPClauseKind(llvm::StringRef Str);
 const char *getOpenMPClauseName(OpenMPClauseKind Kind);
 
 unsigned getOpenMPSimpleClauseType(OpenMPClauseKind Kind, llvm::StringRef Str);
 const char *getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind, unsigned Type);

include/clang/Basic/OpenMPKinds.def

 #define OPENMP_SCHEDULE_KIND(Name)
 #endif
 #ifndef OPENMP_DEPEND_KIND
 #define OPENMP_DEPEND_KIND(Name)
 #endif
 #ifndef OPENMP_LINEAR_KIND
 #define OPENMP_LINEAR_KIND(Name)
 #endif
+#ifndef OPENMP_MAP_KIND
+#define OPENMP_MAP_KIND(Name)
+#endif
 
 // OpenMP directives.
 OPENMP_DIRECTIVE(threadprivate)
 OPENMP_DIRECTIVE(parallel)
 OPENMP_DIRECTIVE(task)
 OPENMP_DIRECTIVE(simd)
 OPENMP_DIRECTIVE(for)
 OPENMP_DIRECTIVE(sections)

 OPENMP_CLAUSE(write, OMPWriteClause)
 OPENMP_CLAUSE(update, OMPUpdateClause)
 OPENMP_CLAUSE(capture, OMPCaptureClause)
 OPENMP_CLAUSE(seq_cst, OMPSeqCstClause)
 OPENMP_CLAUSE(depend, OMPDependClause)
 OPENMP_CLAUSE(device, OMPDeviceClause)
 OPENMP_CLAUSE(threads, OMPThreadsClause)
 OPENMP_CLAUSE(simd, OMPSIMDClause)
+OPENMP_CLAUSE(map, OMPMapClause)
 
 // 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_ATOMIC_CLAUSE(update)
 OPENMP_ATOMIC_CLAUSE(capture)
 OPENMP_ATOMIC_CLAUSE(seq_cst)
 
 // Clauses allowed for OpenMP directive 'target'.
 // TODO More clauses for 'target' directive.
 OPENMP_TARGET_CLAUSE(if)
 OPENMP_TARGET_CLAUSE(device)
+OPENMP_TARGET_CLAUSE(map)
 
 // Clauses allowed for OpenMP directive 'target data'.
 // TODO More clauses for 'target data' directive.
 OPENMP_TARGET_DATA_CLAUSE(if)
 OPENMP_TARGET_DATA_CLAUSE(device)
+OPENMP_TARGET_DATA_CLAUSE(map)
 
 // 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)
 
 // Clauses allowed for OpenMP directive 'ordered'.
 // TODO More clauses for 'ordered' directive.
 OPENMP_ORDERED_CLAUSE(threads)
 OPENMP_ORDERED_CLAUSE(simd)
 
+// Map types and map type modifier for 'map' clause.
+OPENMP_MAP_KIND(alloc)
+OPENMP_MAP_KIND(to)
+OPENMP_MAP_KIND(from)
+OPENMP_MAP_KIND(tofrom)
+OPENMP_MAP_KIND(delete)
+OPENMP_MAP_KIND(release)
+OPENMP_MAP_KIND(always)
+
 #undef OPENMP_LINEAR_KIND
 #undef OPENMP_DEPEND_KIND
 #undef OPENMP_SCHEDULE_KIND
 #undef OPENMP_PROC_BIND_KIND
 #undef OPENMP_DEFAULT_KIND
 #undef OPENMP_DIRECTIVE
 #undef OPENMP_DIRECTIVE_EXT
 #undef OPENMP_CLAUSE
 #undef OPENMP_ORDERED_CLAUSE
 #undef OPENMP_CANCEL_CLAUSE
 #undef OPENMP_SINGLE_CLAUSE
 #undef OPENMP_SECTIONS_CLAUSE
 #undef OPENMP_PARALLEL_CLAUSE
 #undef OPENMP_PARALLEL_FOR_CLAUSE
 #undef OPENMP_PARALLEL_FOR_SIMD_CLAUSE
 #undef OPENMP_PARALLEL_SECTIONS_CLAUSE
 #undef OPENMP_TASK_CLAUSE
 #undef OPENMP_ATOMIC_CLAUSE
 #undef OPENMP_TARGET_CLAUSE
 #undef OPENMP_TARGET_DATA_CLAUSE
 #undef OPENMP_TEAMS_CLAUSE
 #undef OPENMP_SIMD_CLAUSE
 #undef OPENMP_FOR_CLAUSE
 #undef OPENMP_FOR_SIMD_CLAUSE
-
+#undef OPENMP_MAP_KIND

include/clang/Sema/Sema.h

OMPClause *ActOnOpenMPSIMDClause(SourceLocation StartLoc,
                                    SourceLocation EndLoc);
 
   OMPClause *ActOnOpenMPVarListClause(
       OpenMPClauseKind Kind, ArrayRef<Expr *> Vars, Expr *TailExpr,
       SourceLocation StartLoc, SourceLocation LParenLoc,
       SourceLocation ColonLoc, SourceLocation EndLoc,
       CXXScopeSpec &ReductionIdScopeSpec,
       const DeclarationNameInfo &ReductionId, OpenMPDependClauseKind DepKind,
-      OpenMPLinearClauseKind LinKind, SourceLocation DepLinLoc);
+      OpenMPLinearClauseKind LinKind, OpenMPMapClauseKind MapTypeModifier,
+      OpenMPMapClauseKind MapType, SourceLocation DepLinMapLoc);
   /// \brief Called on well-formed 'private' clause.
   OMPClause *ActOnOpenMPPrivateClause(ArrayRef<Expr *> VarList,
                                       SourceLocation StartLoc,
                                       SourceLocation LParenLoc,
                                       SourceLocation EndLoc);
   /// \brief Called on well-formed 'firstprivate' clause.
   OMPClause *ActOnOpenMPFirstprivateClause(ArrayRef<Expr *> VarList,
                                            SourceLocation StartLoc,
                                            SourceLocation LParenLoc,
public:
   ActOnOpenMPDependClause(OpenMPDependClauseKind DepKind, SourceLocation DepLoc,
                           SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
                           SourceLocation StartLoc, SourceLocation LParenLoc,
                           SourceLocation EndLoc);
   /// \brief Called on well-formed 'device' clause.
   OMPClause *ActOnOpenMPDeviceClause(Expr *Device, SourceLocation StartLoc,
                                      SourceLocation LParenLoc,
                                      SourceLocation EndLoc);
- 
+   /// \brief Called on well-formed 'map' clause.
+  OMPClause *ActOnOpenMPMapClause(
+      OpenMPMapClauseKind MapTypeModifier, OpenMPMapClauseKind MapType,
+      SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
+      SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc);
   /// \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.
     CCK_FunctionalCast,

lib/AST/OpenMPClause.cpp

 }
 
 OMPDependClause *OMPDependClause::CreateEmpty(const ASTContext &C, unsigned N) {
   void *Mem = C.Allocate(llvm::RoundUpToAlignment(sizeof(OMPDependClause),
                                                   llvm::alignOf<Expr *>()) +
                          sizeof(Expr *) * N);
   return new (Mem) OMPDependClause(N);
 }
+
+OMPMapClause *OMPMapClause::Create(const ASTContext &C, SourceLocation StartLoc,
+                                   SourceLocation LParenLoc,
+                                   SourceLocation EndLoc, ArrayRef<Expr *> VL,
+                                   OpenMPMapClauseKind TypeModifier,
+                                   OpenMPMapClauseKind Type,
+                                   SourceLocation TypeLoc) {
+  void *Mem = C.Allocate(
+      llvm::RoundUpToAlignment(sizeof(OMPMapClause), llvm::alignOf<Expr *>()) +
+      sizeof(Expr *) * VL.size());
+  OMPMapClause *Clause = new (Mem) OMPMapClause(
+      TypeModifier, Type, TypeLoc, StartLoc, LParenLoc, EndLoc, VL.size());
+  Clause->setVarRefs(VL);
+  Clause->setMapTypeModifier(TypeModifier);
+  Clause->setMapType(Type);
+  Clause->setMapLoc(TypeLoc);
+  return Clause;
+}
+
+OMPMapClause *OMPMapClause::CreateEmpty(const ASTContext &C, unsigned N) {
+  void *Mem = C.Allocate(
+      llvm::RoundUpToAlignment(sizeof(OMPMapClause), llvm::alignOf<Expr *>()) +
+      sizeof(Expr *) * N);
+  return new (Mem) OMPMapClause(N);
+}

lib/AST/StmtPrinter.cpp

if (!Node->varlist_empty()) {
     OS << "depend(";
     OS << getOpenMPSimpleClauseTypeName(Node->getClauseKind(),
                                         Node->getDependencyKind())
        << " :";
     VisitOMPClauseList(Node, ' ');
     OS << ")";
   }
 }
+
+void OMPClausePrinter::VisitOMPMapClause(OMPMapClause *Node) {
+  if (!Node->varlist_empty()) {
+    OS << "map(";
+    if (Node->getMapType() != OMPC_MAP_unknown) {
+      if (Node->getMapTypeModifier() != OMPC_MAP_unknown) {
+        OS << getOpenMPSimpleClauseTypeName(OMPC_map, 
+                                            Node->getMapTypeModifier());
+        OS << ',';
+      }
+      OS << getOpenMPSimpleClauseTypeName(OMPC_map, Node->getMapType());
+      OS << ':';
+    }
+    VisitOMPClauseList(Node, ' ');
+    OS << ")";
+  }
+}
 }
 
 //===----------------------------------------------------------------------===//
 //  OpenMP directives printing methods
 //===----------------------------------------------------------------------===//
 
 void StmtPrinter::PrintOMPExecutableDirective(OMPExecutableDirective *S) {
   OMPClausePrinter Printer(OS, Policy);

lib/AST/StmtProfile.cpp

void OMPClauseProfiler::VisitOMPFlushClause(const OMPFlushClause *C) {
   VisitOMPClauseList(C);
 }
 void OMPClauseProfiler::VisitOMPDependClause(const OMPDependClause *C) {
   VisitOMPClauseList(C);
 }
 void OMPClauseProfiler::VisitOMPDeviceClause(const OMPDeviceClause *C) {
   Profiler->VisitStmt(C->getDevice());
 }
+void OMPClauseProfiler::VisitOMPMapClause(const OMPMapClause *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

 #define OPENMP_DEPEND_KIND(Name) .Case(#Name, OMPC_DEPEND_##Name)
 #include "clang/Basic/OpenMPKinds.def"
         .Default(OMPC_DEPEND_unknown);
   case OMPC_linear:
     return llvm::StringSwitch<OpenMPLinearClauseKind>(Str)
 #define OPENMP_LINEAR_KIND(Name) .Case(#Name, OMPC_LINEAR_##Name)
 #include "clang/Basic/OpenMPKinds.def"
         .Default(OMPC_LINEAR_unknown);
+  case OMPC_map:
+    return llvm::StringSwitch<OpenMPMapClauseKind>(Str)
+#define OPENMP_MAP_KIND(Name) .Case(#Name, OMPC_MAP_##Name)
+#include "clang/Basic/OpenMPKinds.def"
+        .Default(OMPC_MAP_unknown);
   case OMPC_unknown:
   case OMPC_threadprivate:
   case OMPC_if:
   case OMPC_final:
   case OMPC_num_threads:
   case OMPC_safelen:
   case OMPC_simdlen:
   case OMPC_collapse:
case OMPC_linear:
     case OMPC_LINEAR_unknown:
       return "unknown";
 #define OPENMP_LINEAR_KIND(Name)                                             \
   case OMPC_LINEAR_##Name:                                                   \
     return #Name;
 #include "clang/Basic/OpenMPKinds.def"
     }
     llvm_unreachable("Invalid OpenMP 'linear' clause type");
+  case OMPC_map:
+    switch (Type) {
+    case OMPC_MAP_unknown:
+      return "unknown";
+#define OPENMP_MAP_KIND(Name)                                                \
+  case OMPC_MAP_##Name:                                                      \
+    return #Name;
+#include "clang/Basic/OpenMPKinds.def"
+    default:
+      break;
+    }
+    llvm_unreachable("Invalid OpenMP 'map' clause type");
   case OMPC_unknown:
   case OMPC_threadprivate:
   case OMPC_if:
   case OMPC_final:
   case OMPC_num_threads:
   case OMPC_safelen:
   case OMPC_simdlen:
   case OMPC_collapse:

lib/CodeGen/CGStmtOpenMP.cpp

static void EmitOMPAtomicExpr(CodeGenFunction &CGF, OpenMPClauseKind Kind,
   case OMPC_nowait:
   case OMPC_untied:
   case OMPC_threadprivate:
   case OMPC_depend:
   case OMPC_mergeable:
   case OMPC_device:
   case OMPC_threads:
   case OMPC_simd:
+  case OMPC_map:
     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

OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
   case OMPC_shared:
   case OMPC_reduction:
   case OMPC_linear:
   case OMPC_aligned:
   case OMPC_copyin:
   case OMPC_copyprivate:
   case OMPC_flush:
   case OMPC_depend:
+  case OMPC_map:
     Clause = ParseOpenMPVarListClause(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:

 ///    reduction-clause:
 ///       'reduction' '(' reduction-identifier ':' list ')'
 ///    copyprivate-clause:
 ///       'copyprivate' '(' list ')'
 ///    flush-clause:
 ///       'flush' '(' list ')'
 ///    depend-clause:
 ///       'depend' '(' in | out | inout : list ')'
+///    map-clause:
+///       'map' '(' [ [ always , ]
+///          to | from | tofrom | alloc | release | delete ':' ] 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(OpenMPClauseKind Kind) {
   SourceLocation Loc = Tok.getLocation();
   SourceLocation LOpen = ConsumeToken();
   SourceLocation ColonLoc = SourceLocation();
   // Optional scope specifier and unqualified id for reduction identifier.
   CXXScopeSpec ReductionIdScopeSpec;
   UnqualifiedId ReductionId;
   bool InvalidReductionId = false;
   OpenMPDependClauseKind DepKind = OMPC_DEPEND_unknown;
   // OpenMP 4.1 [2.15.3.7, linear Clause]
   //  If no modifier is specified it is assumed to be val.
   OpenMPLinearClauseKind LinearModifier = OMPC_LINEAR_val;
-  SourceLocation DepLinLoc;
+  OpenMPMapClauseKind MapType = OMPC_MAP_unknown;
+  OpenMPMapClauseKind MapTypeModifier = OMPC_MAP_unknown;
+  bool MapTypeModifierSpecified = false;
+  bool UnexpectedId = false;
+  SourceLocation DepLinMapLoc;
 
   // Parse '('.
   BalancedDelimiterTracker T(*this, tok::l_paren, tok::annot_pragma_openmp_end);
   if (T.expectAndConsume(diag::err_expected_lparen_after,
                          getOpenMPClauseName(Kind)))
     return nullptr;
 
   bool NeedRParenForLinear = false;
   BalancedDelimiterTracker LinearT(*this, tok::l_paren,
if (Kind == OMPC_reduction) {
     } else {
       Diag(Tok, diag::warn_pragma_expected_colon) << "reduction identifier";
     }
   } else if (Kind == OMPC_depend) {
   // Handle dependency type for depend clause.
     ColonProtectionRAIIObject ColonRAII(*this);
     DepKind = static_cast<OpenMPDependClauseKind>(getOpenMPSimpleClauseType(
         Kind, Tok.is(tok::identifier) ? PP.getSpelling(Tok) : ""));
-    DepLinLoc = Tok.getLocation();
+    DepLinMapLoc = Tok.getLocation();
 
     if (DepKind == OMPC_DEPEND_unknown) {
       SkipUntil(tok::colon, tok::r_paren, tok::annot_pragma_openmp_end,
                 StopBeforeMatch);
     } else {
       ConsumeToken();
     }
     if (Tok.is(tok::colon)) {
       ColonLoc = ConsumeToken();
     } else {
       Diag(Tok, diag::warn_pragma_expected_colon) << "dependency type";
     }
   } else if (Kind == OMPC_linear) {
     // Try to parse modifier if any.
     if (Tok.is(tok::identifier) && PP.LookAhead(0).is(tok::l_paren)) {
       LinearModifier = static_cast<OpenMPLinearClauseKind>(
           getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok)));
-      DepLinLoc = ConsumeToken();
+      DepLinMapLoc = ConsumeToken();
       LinearT.consumeOpen();
       NeedRParenForLinear = true;
     }
+  } else if (Kind == OMPC_map) {
+    // Handle map type for map clause.
+    ColonProtectionRAIIObject ColonRAII(*this);
+
+    // the first identifier may be a list item, a map-type or
+    //   a map-type-modifier
+    MapType = static_cast<OpenMPMapClauseKind>(getOpenMPSimpleClauseType(
+        Kind, Tok.is(tok::identifier) ? PP.getSpelling(Tok) : ""));
+    DepLinMapLoc = Tok.getLocation();
+    bool ColonExpected = false;
+
+    if (Tok.is(tok::identifier)) {
+      if (PP.LookAhead(0).is(tok::colon)) {
+        MapType = static_cast<OpenMPMapClauseKind>(getOpenMPSimpleClauseType(
+            Kind, Tok.is(tok::identifier) ? PP.getSpelling(Tok) : ""));
+        if (MapType == OMPC_MAP_unknown) {
+          Diag(Tok, diag::err_omp_unknown_map_type);
+        } else if (MapType == OMPC_MAP_always) {
+          Diag(Tok, diag::err_omp_map_type_missing);
+        }
+        ConsumeToken();
+      } else if (PP.LookAhead(0).is(tok::comma)) {
+        if (PP.LookAhead(1).is(tok::identifier) &&
+            PP.LookAhead(2).is(tok::colon)) {
+          MapTypeModifier =
+              static_cast<OpenMPMapClauseKind>(getOpenMPSimpleClauseType(
+                   Kind, Tok.is(tok::identifier) ? PP.getSpelling(Tok) : ""));
+          if (MapTypeModifier != OMPC_MAP_always) {
+            Diag(Tok, diag::err_omp_unknown_map_type_modifier);
+            MapTypeModifier = OMPC_MAP_unknown;
+          } else {
+            MapTypeModifierSpecified = true;
+          }
+
+          ConsumeToken();
+          ConsumeToken();
+
+          MapType = static_cast<OpenMPMapClauseKind>(getOpenMPSimpleClauseType(
+              Kind, Tok.is(tok::identifier) ? PP.getSpelling(Tok) : ""));
+          if (MapType == OMPC_MAP_unknown || MapType == OMPC_MAP_always) {
+            Diag(Tok, diag::err_omp_unknown_map_type);
+          }
+          ConsumeToken();
+        } else {
+          MapType = OMPC_MAP_tofrom;
+        }
+      } else {
+        MapType = OMPC_MAP_tofrom;
+      }
+    } else {
+      UnexpectedId = true;
+    }
+
+    if (Tok.is(tok::colon)) {
+      ColonLoc = ConsumeToken();
+    } else if (ColonExpected) {
+      Diag(Tok, diag::warn_pragma_expected_colon) << "map type";
+    }
   }
 
   SmallVector<Expr *, 5> Vars;
-  bool IsComma = ((Kind != OMPC_reduction) && (Kind != OMPC_depend)) ||
-                 ((Kind == OMPC_reduction) && !InvalidReductionId) ||
-                 ((Kind == OMPC_depend) && DepKind != OMPC_DEPEND_unknown);
+  bool IsComma =
+      ((Kind != OMPC_reduction) && (Kind != OMPC_depend) &&
+       (Kind != OMPC_map)) ||
+      ((Kind == OMPC_map) && (UnexpectedId || MapType != OMPC_MAP_unknown) &&
+       (!MapTypeModifierSpecified ||
+        (MapTypeModifierSpecified && MapTypeModifier == OMPC_MAP_always))) ||
   const bool MayHaveTail = (Kind == OMPC_linear || Kind == OMPC_aligned);
   while (IsComma || (Tok.isNot(tok::r_paren) && Tok.isNot(tok::colon) &&
                      Tok.isNot(tok::annot_pragma_openmp_end))) {
     ColonProtectionRAIIObject ColonRAII(*this, MayHaveTail);
     // Parse variable
     ExprResult VarExpr =
         Actions.CorrectDelayedTyposInExpr(ParseAssignmentExpression());
     if (VarExpr.isUsable()) {
else
       SkipUntil(tok::comma, tok::r_paren, tok::annot_pragma_openmp_end,
                 StopBeforeMatch);
   }
 
   // Parse ')'.
   T.consumeClose();
   if ((Kind == OMPC_depend && DepKind != OMPC_DEPEND_unknown && Vars.empty()) ||
       (Kind != OMPC_depend && Vars.empty()) || (MustHaveTail && !TailExpr) ||
-      InvalidReductionId)
+      (Kind == OMPC_map && MapType == OMPC_MAP_unknown) ||
+      InvalidReductionId) {
     return nullptr;
+  }
 
   return Actions.ActOnOpenMPVarListClause(
       Kind, Vars, TailExpr, Loc, LOpen, ColonLoc, Tok.getLocation(),
       ReductionIdScopeSpec,
       ReductionId.isValid() ? Actions.GetNameFromUnqualifiedId(ReductionId)
                             : DeclarationNameInfo(),
-      DepKind, LinearModifier, DepLinLoc);
+      DepKind, LinearModifier, MapTypeModifier, MapType, DepLinMapLoc);
 }
 

lib/Sema/SemaOpenMP.cpp

struct DSAVarData {
     OpenMPClauseKind CKind;
     DeclRefExpr *RefExpr;
     SourceLocation ImplicitDSALoc;
     DSAVarData()
         : DKind(OMPD_unknown), CKind(OMPC_unknown), RefExpr(nullptr),
           ImplicitDSALoc() {}
   };
 
+public:
+  struct MapInfo {
+    Expr *RefExpr;
+  };
+
 private:
   struct DSAInfo {
     OpenMPClauseKind Attributes;
     DeclRefExpr *RefExpr;
   };
   typedef llvm::SmallDenseMap<VarDecl *, DSAInfo, 64> DeclSAMapTy;
   typedef llvm::SmallDenseMap<VarDecl *, DeclRefExpr *, 64> AlignedMapTy;
   typedef llvm::DenseSet<VarDecl *> LoopControlVariablesSetTy;
+  typedef llvm::SmallDenseMap<VarDecl *, MapInfo, 64> MappedDeclsTy;
 
   struct SharingMapTy {
     DeclSAMapTy SharingMap;
     AlignedMapTy AlignedMap;
+    MappedDeclsTy MappedDecls;
     LoopControlVariablesSetTy LCVSet;
     DefaultDataSharingAttributes DefaultAttr;
     SourceLocation DefaultAttrLoc;
     OpenMPDirectiveKind Directive;
     DeclarationNameInfo DirectiveName;
     Scope *CurScope;
     SourceLocation ConstructLoc;
     /// \brief first argument (Expr *) contains optional argument of the
SourceLocation getInnerTeamsRegionLoc() const {
     if (Stack.size() > 1)
       return Stack.back().InnerTeamsRegionLoc;
     return SourceLocation();
   }
 
   Scope *getCurScope() const { return Stack.back().CurScope; }
   Scope *getCurScope() { return Stack.back().CurScope; }
   SourceLocation getConstructLoc() { return Stack.back().ConstructLoc; }
+
+  MapInfo getMapInfoForVar(VarDecl *VD) {
+    MapInfo VarMI = {0};
+    for (auto Cnt = Stack.size() - 1; Cnt > 0; --Cnt) {
+      if (Stack[Cnt].MappedDecls.count(VD)) {
+        VarMI = Stack[Cnt].MappedDecls[VD];
+        break;
+      }
+    }
+    return VarMI;
+  }
+
+  void addMapInfoForVar(VarDecl *VD, MapInfo MI) {
+    if (Stack.size() > 1) {
+      Stack.back().MappedDecls[VD] = MI;
+    }
+  }
+
+  MapInfo IsMappedInCurrentRegion(VarDecl *VD) {
+    assert(Stack.size() > 1 && "Target level is 0");
+    MapInfo VarMI = {0};
+    if (Stack.size() > 1 && Stack.back().MappedDecls.count(VD)) {
+      VarMI = Stack.back().MappedDecls[VD];
+    }
+    return VarMI;
+  }
 };
 bool isParallelOrTaskRegion(OpenMPDirectiveKind DKind) {
   return isOpenMPParallelDirective(DKind) || DKind == OMPD_task ||
          isOpenMPTeamsDirective(DKind) || DKind == OMPD_unknown;
 }
 } // namespace
 
 DSAStackTy::DSAVarData DSAStackTy::getDSA(StackTy::reverse_iterator Iter,
OMPClause *Sema::ActOnOpenMPSingleExprClause(OpenMPClauseKind Kind, Expr *Expr,
   case OMPC_read:
   case OMPC_write:
   case OMPC_update:
   case OMPC_capture:
   case OMPC_seq_cst:
   case OMPC_depend:
   case OMPC_threads:
   case OMPC_simd:
+  case OMPC_map:
   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_write:
   case OMPC_update:
   case OMPC_capture:
   case OMPC_seq_cst:
   case OMPC_depend:
   case OMPC_device:
   case OMPC_threads:
   case OMPC_simd:
+  case OMPC_map:
   case OMPC_unknown:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;
 }
 
 OMPClause *Sema::ActOnOpenMPDefaultClause(OpenMPDefaultClauseKind Kind,
                                           SourceLocation KindKwLoc,
OMPClause *Sema::ActOnOpenMPSingleExprWithArgClause(
   case OMPC_write:
   case OMPC_update:
   case OMPC_capture:
   case OMPC_seq_cst:
   case OMPC_depend:
   case OMPC_device:
   case OMPC_threads:
   case OMPC_simd:
+  case OMPC_map:
   case OMPC_unknown:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;
 }
 
 OMPClause *Sema::ActOnOpenMPScheduleClause(
     OpenMPScheduleClauseKind Kind, Expr *ChunkSize, SourceLocation StartLoc,
OMPClause *Sema::ActOnOpenMPClause(OpenMPClauseKind Kind,
   case OMPC_copyin:
   case OMPC_copyprivate:
   case OMPC_default:
   case OMPC_proc_bind:
   case OMPC_threadprivate:
   case OMPC_flush:
   case OMPC_depend:
   case OMPC_device:
+  case OMPC_map:
   case OMPC_unknown:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;
 }
 
 OMPClause *Sema::ActOnOpenMPNowaitClause(SourceLocation StartLoc,
                                          SourceLocation EndLoc) {
OMPClause *Sema::ActOnOpenMPSIMDClause(SourceLocation StartLoc,
   return new (Context) OMPSIMDClause(StartLoc, EndLoc);
 }
 
 OMPClause *Sema::ActOnOpenMPVarListClause(
     OpenMPClauseKind Kind, ArrayRef<Expr *> VarList, Expr *TailExpr,
     SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation ColonLoc,
     SourceLocation EndLoc, CXXScopeSpec &ReductionIdScopeSpec,
     const DeclarationNameInfo &ReductionId, OpenMPDependClauseKind DepKind,
-    OpenMPLinearClauseKind LinKind, SourceLocation DepLinLoc) {
+    OpenMPLinearClauseKind LinKind, OpenMPMapClauseKind MapTypeModifier, 
+    OpenMPMapClauseKind MapType, SourceLocation DepLinMapLoc) {
   OMPClause *Res = nullptr;
   switch (Kind) {
   case OMPC_private:
     Res = ActOnOpenMPPrivateClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_firstprivate:
     Res = ActOnOpenMPFirstprivateClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_lastprivate:
     Res = ActOnOpenMPLastprivateClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_shared:
     Res = ActOnOpenMPSharedClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_reduction:
     Res = ActOnOpenMPReductionClause(VarList, StartLoc, LParenLoc, ColonLoc,
                                      EndLoc, ReductionIdScopeSpec, ReductionId);
     break;
   case OMPC_linear:
     Res = ActOnOpenMPLinearClause(VarList, TailExpr, StartLoc, LParenLoc,
-                                  LinKind, DepLinLoc, ColonLoc, EndLoc);
+                                  LinKind, DepLinMapLoc, ColonLoc, EndLoc);
     break;
   case OMPC_aligned:
     Res = ActOnOpenMPAlignedClause(VarList, TailExpr, StartLoc, LParenLoc,
                                    ColonLoc, EndLoc);
     break;
   case OMPC_copyin:
     Res = ActOnOpenMPCopyinClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_copyprivate:
     Res = ActOnOpenMPCopyprivateClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_flush:
     Res = ActOnOpenMPFlushClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_depend:
-    Res = ActOnOpenMPDependClause(DepKind, DepLinLoc, ColonLoc, VarList, StartLoc,
-                                  LParenLoc, EndLoc);
+    Res = ActOnOpenMPDependClause(DepKind, DepLinMapLoc, ColonLoc, VarList, 
+                                  StartLoc, LParenLoc, EndLoc);
+    break;
+  case OMPC_map:
+    Res = ActOnOpenMPMapClause(MapTypeModifier, MapType, DepLinMapLoc, ColonLoc,
+                               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:
if (Value.get()->isIntegerConstantExpr(Result, Context) &&
         Result.isSigned() && !Result.isStrictlyPositive()) {
       Diag(Loc, diag::err_omp_negative_expression_in_clause)
           << "device" << ValExpr->getSourceRange();
       return nullptr;
     }
   }
   return new (Context) OMPDeviceClause(ValExpr, StartLoc, LParenLoc, EndLoc);
 }
+
+static bool IsCXXRecordForMappable(Sema &SemaRef, SourceLocation Loc,
+                                   DSAStackTy *Stack, CXXRecordDecl *RD) {
+  if (!RD || RD->isInvalidDecl())
+    return true;
+
+  auto QTy = SemaRef.Context.getRecordType(RD);
+  if (RD->isDynamicClass()) {
+    SemaRef.Diag(Loc, diag::err_omp_not_mappable_type) << QTy;
+    SemaRef.Diag(RD->getLocation(), diag::note_omp_polymorphic_in_target);
+    return false;
+  }
+  auto *DC = RD;
+  bool IsCorrect = true;
+  for (auto *I : DC->decls()) {
+    if (I) {
+      if (auto *MD = dyn_cast<CXXMethodDecl>(I)) {
+        if (MD->isStatic()) {
+          SemaRef.Diag(Loc, diag::err_omp_not_mappable_type) << QTy;
+          SemaRef.Diag(MD->getLocation(),
+                       diag::note_omp_static_member_in_target);
+          IsCorrect = false;
+        }
+      } else if (auto *VD = dyn_cast<VarDecl>(I)) {
+        if (VD->isStaticDataMember()) {
+          SemaRef.Diag(Loc, diag::err_omp_not_mappable_type) << QTy;
+          SemaRef.Diag(VD->getLocation(),
+                       diag::note_omp_static_member_in_target);
+          IsCorrect = false;
+        }
+      }
+    }
+  }
+
+  for (auto &I : RD->bases()) {
+    if (!IsCXXRecordForMappable(SemaRef, I.getLocStart(), Stack,
+                                I.getType()->getAsCXXRecordDecl()))
+      IsCorrect = false;
+  }
+  return IsCorrect;
+}
+
+static bool CheckTypeMappable(SourceLocation SL, SourceRange SR, Sema &SemaRef,
+                              DSAStackTy *Stack, QualType QTy) {
+  NamedDecl *ND;
+  if (QTy->isIncompleteType(&ND)) {
+    SemaRef.Diag(SL, diag::err_incomplete_type) << QTy << SR;
+    return false;
+  } else if (CXXRecordDecl *RD = dyn_cast_or_null<CXXRecordDecl>(ND)) {
+    if (!RD->isInvalidDecl() &&
+        !IsCXXRecordForMappable(SemaRef, SL, Stack, RD))
+      return false;
+  }
+  return true;
+}
+
+OMPClause *Sema::ActOnOpenMPMapClause(
+    OpenMPMapClauseKind MapTypeModifier, OpenMPMapClauseKind MapType,
+    SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
+    SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc) {
+  SmallVector<Expr *, 4> Vars;
+
+  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();
+
+    // OpenMP [2.14.5, Restrictions]
+    //  A variable that is part of another variable (such as field of a
+    //  structure) but is not an array element or an array section cannot appear
+    //  in a map clause.
+    auto *VE = RE->IgnoreParenLValueCasts();
+
+    if (VE->isValueDependent() || VE->isTypeDependent() ||
+        VE->isInstantiationDependent() ||
+        VE->containsUnexpandedParameterPack()) {
+      // It will be analyzed later.
+      Vars.push_back(RE);
+      continue;
+    }
+
+    auto *SimpleExpr = RE->IgnoreParenCasts();
+    auto *DE = dyn_cast<DeclRefExpr>(SimpleExpr);
+    auto *ASE = dyn_cast<ArraySubscriptExpr>(SimpleExpr);
+    auto *OASE = dyn_cast<OMPArraySectionExpr>(SimpleExpr);
+
+    if (!RE->IgnoreParenImpCasts()->isLValue() ||
+        (!OASE && !ASE && !DE) ||
+        (DE && !isa<VarDecl>(DE->getDecl())) ||
+        (ASE && !ASE->getBase()->getType()->isAnyPointerType() &&
+         !ASE->getBase()->getType()->isArrayType())) {
+      Diag(ELoc, diag::err_omp_expected_var_name_or_array_item)
+        << RE->getSourceRange();
+      continue;
+    }
+
+    Decl *D = nullptr;
+    if (DE) {
+      D = DE->getDecl();
+    } else if (ASE) {
+      auto *B = ASE->getBase()->IgnoreParenCasts();
+      D = dyn_cast<DeclRefExpr>(B)->getDecl();
+    } else if (OASE) {
+      auto *B = OASE->getBase();
+      D = dyn_cast<DeclRefExpr>(B)->getDecl();
+    }
+    assert(D && "Null decl on map clause.");
+    auto *VD = cast<VarDecl>(D);
+
+    // OpenMP [2.14.5, Restrictions, p.8]
+    // threadprivate variables cannot appear in a map clause.
+    if (DSAStack->isThreadPrivate(VD)) {
+      auto DVar = DSAStack->getTopDSA(VD, false);
+      Diag(ELoc, diag::err_omp_threadprivate_in_map);
+      ReportOriginalDSA(*this, DSAStack, VD, DVar);
+      continue;
+    }
+
+    // OpenMP [2.14.5, Restrictions, p.2]
+    //  At most one list item can be an array item derived from a given variable
+    //  in map clauses of the same construct.
+    // OpenMP [2.14.5, Restrictions, p.3]
+    //  List items of map clauses in the same construct must not share original
+    //  storage.
+    // OpenMP [2.14.5, Restrictions, C/C++, p.2]
+    //  A variable for which the type is pointer, reference to array, or
+    //  reference to pointer and an array section derived from that variable
+    //  must not appear as list items of map clauses of the same construct.
+    DSAStackTy::MapInfo MI = DSAStack->IsMappedInCurrentRegion(VD);
+    if (MI.RefExpr) {
+      Diag(ELoc, diag::err_omp_map_shared_storage) << ELoc;
+      Diag(MI.RefExpr->getExprLoc(), diag::note_used_here)
+          << MI.RefExpr->getSourceRange();
+      continue;
+    }
+
+    // OpenMP [2.14.5, Restrictions, C/C++, p.3,4]
+    //  A variable for which the type is pointer, reference to array, or
+    //  reference to pointer must not appear as a list item if the enclosing
+    //  device data environment already contains an array section derived from
+    //  that variable.
+    //  An array section derived from a variable for which the type is pointer,
+    //  reference to array, or reference to pointer must not appear as a list
+    //  item if the enclosing device data environment already contains that
+    //  variable.
+    QualType Type = VD->getType();
+    MI = DSAStack->getMapInfoForVar(VD);
+    if (MI.RefExpr && (isa<DeclRefExpr>(MI.RefExpr->IgnoreParenLValueCasts()) !=
+                       isa<DeclRefExpr>(VE)) &&
+        (Type->isPointerType() || Type->isReferenceType())) {
+      Diag(ELoc, diag::err_omp_map_shared_storage) << ELoc;
+      Diag(MI.RefExpr->getExprLoc(), diag::note_used_here)
+          << MI.RefExpr->getSourceRange();
+      continue;
+    }
+
+    // OpenMP [2.14.5, Restrictions, C/C++, p.7]
+    //  A list item must have a mappable type.
+    if (!CheckTypeMappable(VE->getExprLoc(), VE->getSourceRange(), *this,
+                           DSAStack, Type))
+      continue;
+
+    Vars.push_back(RE);
+    MI.RefExpr = RE;
+    DSAStack->addMapInfoForVar(VD, MI);
+  }
+  if (Vars.empty())
+    return nullptr;
+
+  return OMPMapClause::Create(Context, StartLoc, LParenLoc, EndLoc, Vars,
+                              MapTypeModifier, MapType, MapLoc);
+}

lib/Sema/TreeTransform.h

#include "clang/Basic/OpenMPKinds.def"
   /// Subclasses may override this routine to provide different behavior.
   OMPClause *RebuildOMPDeviceClause(Expr *Device, SourceLocation StartLoc,
                                     SourceLocation LParenLoc,
                                     SourceLocation EndLoc) {
     return getSema().ActOnOpenMPDeviceClause(Device, StartLoc, LParenLoc, 
                                              EndLoc);
   }
 
+  /// \brief Build a new OpenMP 'map' clause.
+  ///
+  /// By default, performs semantic analysis to build the new OpenMP clause.
+  /// Subclasses may override this routine to provide different behavior.
+  OMPClause *RebuildOMPMapClause(
+      OpenMPMapClauseKind MapTypeModifier, OpenMPMapClauseKind MapType,
+      SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
+      SourceLocation StartLoc, SourceLocation LParenLoc,
+      SourceLocation EndLoc) {
+    return getSema().ActOnOpenMPMapClause(MapTypeModifier, MapType, MapLoc,
+                                          ColonLoc, 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);
   }

 TreeTransform<Derived>::TransformOMPDeviceClause(OMPDeviceClause *C) {
   ExprResult E = getDerived().TransformExpr(C->getDevice());
   if (E.isInvalid())
     return nullptr;
   return getDerived().RebuildOMPDeviceClause(
       E.get(), C->getLocStart(), C->getLParenLoc(), C->getLocEnd());
 }
 
+template <typename Derived>
+OMPClause *TreeTransform<Derived>::TransformOMPMapClause(OMPMapClause *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 nullptr;
+    Vars.push_back(EVar.get());
+  }
+  return getDerived().RebuildOMPMapClause(
+      C->getMapTypeModifier(), C->getMapType(), C->getMapLoc(),
+      C->getColonLoc(), 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_flush:
     C = OMPFlushClause::CreateEmpty(Context, Record[Idx++]);
     break;
   case OMPC_depend:
     C = OMPDependClause::CreateEmpty(Context, Record[Idx++]);
     break;
   case OMPC_device:
     C = new (Context) OMPDeviceClause();
     break;
+  case OMPC_map:
+    C = OMPMapClause::CreateEmpty(Context, Record[Idx++]);
+    break;
   }
   Visit(C);
   C->setLocStart(Reader->ReadSourceLocation(Record, Idx));
   C->setLocEnd(Reader->ReadSourceLocation(Record, Idx));
 
   return C;
 }
 
void OMPClauseReader::VisitOMPDependClause(OMPDependClause *C) {
   C->setVarRefs(Vars);
 }
 
 void OMPClauseReader::VisitOMPDeviceClause(OMPDeviceClause *C) {
   C->setDevice(Reader->Reader.ReadSubExpr());
   C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx));
 }
 
+void OMPClauseReader::VisitOMPMapClause(OMPMapClause *C) {
+  C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx));
+  C->setMapTypeModifier(
+     static_cast<OpenMPMapClauseKind>(Record[Idx++]));
+  C->setMapType(
+     static_cast<OpenMPMapClauseKind>(Record[Idx++]));
+  C->setMapLoc(Reader->ReadSourceLocation(Record, Idx));
+  C->setColonLoc(Reader->ReadSourceLocation(Record, Idx));
+  auto 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

for (auto *VE : C->varlists())
     Writer->Writer.AddStmt(VE);
 }
 
 void OMPClauseWriter::VisitOMPDeviceClause(OMPDeviceClause *C) {
   Writer->Writer.AddStmt(C->getDevice());
   Writer->Writer.AddSourceLocation(C->getLParenLoc(), Record);
 }
 
+void OMPClauseWriter::VisitOMPMapClause(OMPMapClause *C) {
+  Record.push_back(C->varlist_size());
+  Writer->Writer.AddSourceLocation(C->getLParenLoc(), Record);
+  Record.push_back(C->getMapTypeModifier());
+  Record.push_back(C->getMapType());
+  Writer->Writer.AddSourceLocation(C->getMapLoc(), Record);
+  Writer->Writer.AddSourceLocation(C->getColonLoc(), 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_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 <typename T, int C>
 T tmain(T argc, T *argv) {
+  T i, j, a[20];
 #pragma omp target
   foo();
 #pragma omp target if (target:argc > 0)
   foo();
 #pragma omp target if (C)
   foo();
+#pragma omp target map(i)
+  foo();
+#pragma omp target map(a[0:10], i)
+  foo();
+#pragma omp target map(to: i) map(from: j)
+  foo();
+#pragma omp target map(always,alloc: i)
+  foo();
   return 0;
 }
 
 // CHECK: template <typename T = int, int C = 5> int tmain(int argc, int *argv) {
+// CHECK-NEXT: int i, j, a[20]
 // CHECK-NEXT: #pragma omp target
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target if(target: argc > 0)
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target if(5)
 // CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(tofrom: i)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(tofrom: a[0:10],i)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(to: i) map(from: j)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(always,alloc: i)
+// CHECK-NEXT: foo()
 // CHECK: template <typename T = char, int C = 1> char tmain(char argc, char *argv) {
+// CHECK-NEXT: char i, j, a[20]
 // CHECK-NEXT: #pragma omp target
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target if(target: argc > 0)
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target if(1)
 // CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(tofrom: i)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(tofrom: a[0:10],i)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(to: i) map(from: j)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(always,alloc: i)
+// CHECK-NEXT: foo()
 // CHECK: template <typename T, int C> T tmain(T argc, T *argv) {
+// CHECK-NEXT: T i, j, a[20]
 // CHECK-NEXT: #pragma omp target
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target if(target: argc > 0)
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target if(C)
 // CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(tofrom: i)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(tofrom: a[0:10],i)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(to: i) map(from: j)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(always,alloc: i)
+// CHECK-NEXT: foo()
 
 // CHECK-LABEL: int main(int argc, char **argv) {
 int main (int argc, char **argv) {
+  int i, j, a[20];
+// CHECK-NEXT: int i, j, a[20]
 #pragma omp target
 // CHECK-NEXT: #pragma omp target
   foo();
 // CHECK-NEXT: foo();
 #pragma omp target if (argc > 0)
 // CHECK-NEXT: #pragma omp target if(argc > 0)
   foo();
 // CHECK-NEXT: foo();
+
+#pragma omp target map(i) if(argc>0)
+// CHECK-NEXT: #pragma omp target map(tofrom: i) if(argc > 0)
+  foo();
+// CHECK-NEXT: foo();
+
+#pragma omp target map(i)
+// CHECK-NEXT: #pragma omp target map(tofrom: i)
+  foo();
+// CHECK-NEXT: foo();
+
+#pragma omp target map(a[0:10], i)
+// CHECK-NEXT: #pragma omp target map(tofrom: a[0:10],i)
+  foo();
+// CHECK-NEXT: foo();
+
+#pragma omp target map(to: i) map(from: j)
+// CHECK-NEXT: #pragma omp target map(to: i) map(from: j)
+  foo();
+// CHECK-NEXT: foo();
+
+#pragma omp target map(always,alloc: i)
+// CHECK-NEXT: #pragma omp target map(always,alloc: i)
+  foo();
+// CHECK-NEXT: foo();
+
   return tmain<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);
 }
 
 #endif

test/OpenMP/target_data_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 <typename T, int C>
+T tmain(T argc, T *argv) {
+  T i, j, b, c, d, e, x[20];
+
+#pragma omp target data
+  i = argc;
+
+#pragma omp target data if (target data: j > 0)
+  foo();
+
+#pragma omp target data if (b)
+  foo();
+
+#pragma omp target data map(c)
+  foo();
+
+#pragma omp target data map(c) if(b>e)
+  foo();
+
+#pragma omp target data map(x[0:10], c)
+  foo();
+
+#pragma omp target data map(to: c) map(from: d)
+  foo();
+
+#pragma omp target data map(always,alloc: e)
+  foo();
+
+// nesting a target region
+#pragma omp target data map(e)
+{
+  #pragma omp target map(always, alloc: e)
+    foo();
+}
+
+  return 0;
+}
+
+// CHECK: template <typename T = int, int C = 5> int tmain(int argc, int *argv) {
+// CHECK-NEXT: int i, j, b, c, d, e, x[20];
+// CHECK-NEXT: #pragma omp target data
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target data if(target data: j > 0)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data if(b)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(tofrom: c)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(tofrom: c) if(b > e)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(tofrom: x[0:10],c)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(to: c) map(from: d)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(always,alloc: e)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(tofrom: e)
+// CHECK-NEXT: {
+// CHECK-NEXT: #pragma omp target map(always,alloc: e)
+// CHECK-NEXT: foo();
+// CHECK: template <typename T = char, int C = 1> char tmain(char argc, char *argv) {
+// CHECK-NEXT: char i, j, b, c, d, e, x[20];
+// CHECK-NEXT: #pragma omp target data
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target data if(target data: j > 0)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data if(b)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(tofrom: c)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(tofrom: c) if(b > e)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(tofrom: x[0:10],c)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(to: c) map(from: d)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(always,alloc: e)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(tofrom: e)
+// CHECK-NEXT: {
+// CHECK-NEXT: #pragma omp target map(always,alloc: e)
+// CHECK-NEXT: foo();
+// CHECK: template <typename T, int C> T tmain(T argc, T *argv) {
+// CHECK-NEXT: T i, j, b, c, d, e, x[20];
+// CHECK-NEXT: #pragma omp target data
+// CHECK-NEXT: i = argc;
+// CHECK-NEXT: #pragma omp target data if(target data: j > 0)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data if(b)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(tofrom: c)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(tofrom: c) if(b > e)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(tofrom: x[0:10],c)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(to: c) map(from: d)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(always,alloc: e)
+// CHECK-NEXT: foo();
+// CHECK-NEXT: #pragma omp target data map(tofrom: e)
+// CHECK-NEXT: {
+// CHECK-NEXT: #pragma omp target map(always,alloc: e)
+// CHECK-NEXT: foo();
+
 int main (int argc, char **argv) {
-  int b = argc, c, d, e, f, g;
+  int b = argc, c, d, e, f, g, x[20];
   static int a;
 // CHECK: static int a;
 
 #pragma omp target data
 // CHECK:      #pragma omp target data
   a=2;
 // CHECK-NEXT: a = 2;
 #pragma omp target data if (target data: b)
 // CHECK: #pragma omp target data if(target data: b)
   foo();
 // CHECK-NEXT: foo();
 
 #pragma omp target data if (b > g)
 // CHECK: #pragma omp target data if(b > g)
   foo();
 // CHECK-NEXT: foo();
 
-  return (0);
+#pragma omp target data map(c)
+// CHECK-NEXT: #pragma omp target data map(tofrom: c)
+  foo();
+// CHECK-NEXT: foo();
+
+#pragma omp target data map(c) if(b>g)
+// CHECK-NEXT: #pragma omp target data map(tofrom: c) if(b > g)
+  foo();
+// CHECK-NEXT: foo();
+
+#pragma omp target data map(x[0:10], c)
+// CHECK-NEXT: #pragma omp target data map(tofrom: x[0:10],c)
+  foo();
+// CHECK-NEXT: foo();
+
+#pragma omp target data map(to: c) map(from: d)
+// CHECK-NEXT: #pragma omp target data map(to: c) map(from: d)
+  foo();
+// CHECK-NEXT: foo();
+
+#pragma omp target data map(always,alloc: e)
+// CHECK-NEXT: #pragma omp target data map(always,alloc: e)
+  foo();
+// CHECK-NEXT: foo();
+
+// nesting a target region
+#pragma omp target data map(e)
+// CHECK-NEXT: #pragma omp target data map(tofrom: e)
+{
+// CHECK-NEXT: {
+  #pragma omp target map(always, alloc: e)
+// CHECK-NEXT: #pragma omp target map(always,alloc: e)
+    foo();
+// CHECK-NEXT: foo();
+}
+  return tmain<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);
 }
 
 #endif

test/OpenMP/target_map_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) { }
+};
+
+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 i, t[20];
+  T &j = i;
+  T *k = &j;
+  T x;
+  T y;
+  T to, tofrom, always;
+  const T (&l)[5] = da;
+
+
+#pragma omp target map // expected-error {{expected '(' after 'map'}}
+#pragma omp target map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+#pragma omp target map() // expected-error {{expected expression}}
+#pragma omp target map(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+#pragma omp target map(to argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected ',' or ')' in 'map' clause}}
+#pragma omp target map(to:) // expected-error {{expected expression}}
+#pragma omp target map(from: argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target map(x: y) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+#pragma omp target map(x)
+  foo();
+#pragma omp target map(tofrom: t[:I])
+  foo();
+#pragma omp target map(T: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+  foo();
+#pragma omp target map(T) // expected-error {{'T' does not refer to a value}}
+  foo();
+#pragma omp target map(I) // expected-error 2 {{expected variable name, array element or array section}}
+  foo();
+#pragma omp target map(S2::S2s)
+  foo();
+#pragma omp target map(S2::S2sc)
+  foo();
+#pragma omp target map(x)
+  foo();
+#pragma omp target map(to: x)
+  foo();
+#pragma omp target map(to: to)
+  foo();
+#pragma omp target map(to)
+  foo();
+#pragma omp target map(to, x)
+  foo();
+#pragma omp target map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
+#pragma omp target map(tofrom: argc > 0 ? x : y) // expected-error 2 {{expected variable name, array element or array section}} 
+#pragma omp target map(argc)
+#pragma omp target map(S1) // expected-error {{'S1' does not refer to a value}}
+#pragma omp target 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 map(ba) // expected-error 2 {{type 'S2' is not mappable to target}}
+#pragma omp target map(ca)
+#pragma omp target map(da)
+#pragma omp target map(S2::S2s)
+#pragma omp target map(S2::S2sc)
+#pragma omp target map(e, g)
+#pragma omp target map(h) // expected-error {{threadprivate variables are not allowed in map clause}}
+#pragma omp target 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 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+  foo();
+#pragma omp target map(da)
+#pragma omp target map(da[:4])
+  foo();
+#pragma omp target map(k, j, l) // expected-note 4 {{used here}}
+#pragma omp target map(k[:4]) // expected-error 2 {{variable already marked as mapped in current construct}}
+#pragma omp target map(j)
+#pragma omp target map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}}
+  foo();
+#pragma omp target map(k[:4], j, l[:5]) // expected-note 4 {{used here}}
+#pragma omp target map(k) // expected-error 2 {{variable already marked as mapped in current construct}}
+#pragma omp target map(j)
+#pragma omp target map(l) // expected-error 2 {{variable already marked as mapped in current construct}}
+  foo();
+
+#pragma omp target map(always, tofrom: x)
+#pragma omp target map(always: x) // expected-error {{missing map type}}
+#pragma omp target 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 map(always, tofrom: always, tofrom, x)
+#pragma omp target map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
+  foo();
+
+  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;
+  int &j = i;
+  int *k = &j;
+  int x;
+  int y;
+  int to, tofrom, always;
+  const int (&l)[5] = da;
+#pragma omp target map // expected-error {{expected '(' after 'map'}}
+#pragma omp target map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+#pragma omp target map() // expected-error {{expected expression}}
+#pragma omp target map(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+#pragma omp target map(to argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected ',' or ')' in 'map' clause}}
+#pragma omp target map(to:) // expected-error {{expected expression}}
+#pragma omp target map(from: argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target map(x: y) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+#pragma omp target map(x)
+  foo();
+#pragma omp target map(to: x)
+  foo();
+#pragma omp target map(to: to)
+  foo();
+#pragma omp target map(to)
+  foo();
+#pragma omp target map(to, x)
+  foo();
+#pragma omp target map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
+#pragma omp target map(tofrom: argc > 0 ? argv[1] : argv[2]) // expected-error {{expected variable name, array element or array section}}
+#pragma omp target map(argc)
+#pragma omp target map(S1) // expected-error {{'S1' does not refer to a value}}
+#pragma omp target 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 map(argv[1])
+#pragma omp target map(ba) // expected-error 2 {{type 'S2' is not mappable to target}}
+#pragma omp target map(ca)
+#pragma omp target map(da)
+#pragma omp target map(S2::S2s)
+#pragma omp target map(S2::S2sc)
+#pragma omp target map(e, g)
+#pragma omp target map(h) // expected-error {{threadprivate variables are not allowed in map clause}}
+#pragma omp target 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 {{variable already marked as mapped in current construct}} expected-note {{used here}}
+  foo();
+#pragma omp target map(da)
+#pragma omp target map(da[:4])
+  foo();
+#pragma omp target map(k, j, l) // expected-note 2 {{used here}}
+#pragma omp target map(k[:4]) // expected-error {{variable already marked as mapped in current construct}}
+#pragma omp target map(j)
+#pragma omp target map(l[:5]) // expected-error {{variable already marked as mapped in current construct}}
+  foo();
+#pragma omp target map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
+#pragma omp target map(k) // expected-error {{variable already marked as mapped in current construct}}
+#pragma omp target map(j)
+#pragma omp target map(l) // expected-error {{variable already marked as mapped in current construct}}
+  foo();
+
+#pragma omp target map(always, tofrom: x)
+#pragma omp target map(always: x) // expected-error {{missing map type}}
+#pragma omp target 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 map(always, tofrom: always, tofrom, x)
+#pragma omp target map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
+  foo();
+
+  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

OMPClauseEnqueue::VisitOMPCopyprivateClause(const OMPCopyprivateClause *C) {
   }
 }
 void OMPClauseEnqueue::VisitOMPFlushClause(const OMPFlushClause *C) {
   VisitOMPClauseList(C);
 }
 void OMPClauseEnqueue::VisitOMPDependClause(const OMPDependClause *C) {
   VisitOMPClauseList(C);
 }
+void OMPClauseEnqueue::VisitOMPMapClause(const OMPMapClause *C) {
+  VisitOMPClauseList(C);
+}
 }
 
 void EnqueueVisitor::EnqueueChildren(const OMPClause *S) {
   unsigned size = WL.size();
   OMPClauseEnqueue Visitor(this);
   Visitor.Visit(S);
   if (size == WL.size())
     return;