[OpenMP] parsing and sema support for 'close' map-type-modifier

[saghir]

Summary

A map clause with the close map-type-modifier is a hint to prefer that the variables are mapped using a copy into faster memory.

Reference: TR7 Section 2.15.8 and 2.22.7.1

[saghir]

Made following changes:

• added more tests for checking diagnostic error messages.
• added initialization of modifiers' location in OMPMapClause.
• improved variable names to depict multiple modifiers as opposed to a single modifier.
• improved comments.

Also updated to get latest changes.

[ABataev]

LG

include/clang/AST/OpenMPClause.h

class OMPMapClause final : public OMPMappableExprListClause<OMPMapClause>,
   }
   size_t numTrailingObjects(OverloadToken<ValueDecl *>) const {
     return getUniqueDeclarationsNum();
   }
   size_t numTrailingObjects(OverloadToken<unsigned>) const {
     return getUniqueDeclarationsNum() + getTotalComponentListNum();
   }
 
-  /// Map type modifier for the 'map' clause.
-  OpenMPMapClauseKind MapTypeModifier = OMPC_MAP_unknown;
+public:
+  /// Number of allowed map-type-modifiers.
+  static constexpr unsigned NumberOfModifiers =
+      OMPC_MAP_MODIFIER_last - OMPC_MAP_MODIFIER_unknown - 1;
+
+private:
+  /// Map-type-modifiers for the 'map' clause.
+  OpenMPMapModifierKind MapTypeModifiers[NumberOfModifiers] = {
+    OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown
+  };
+  
+  /// Location of map-type-modifiers for the 'map' clause.
+  SourceLocation MapTypeModifiersLoc[NumberOfModifiers];
 
   /// Map type for the 'map' clause.
   OpenMPMapClauseKind MapType = OMPC_MAP_unknown;
 
   /// Is this an implicit map type or not.
   bool MapTypeIsImplicit = false;
 
   /// Location of the map type.
   SourceLocation MapLoc;
 
   /// Colon location.
   SourceLocation ColonLoc;
 
   /// Build a clause for \a NumVars listed expressions, \a
   /// NumUniqueDeclarations declarations, \a NumComponentLists total component
   /// lists, and \a NumComponents total expression components.
   ///
-  /// \param MapTypeModifier Map type modifier.
+  /// \param MapModifiers Map-type-modifiers.
+  /// \param MapModifiersLoc Locations of map-type-modifiers.
   /// \param MapType Map type.
   /// \param MapTypeIsImplicit Map type is inferred implicitly.
   /// \param MapLoc Location of the map type.
   /// \param StartLoc Starting location of the clause.
   /// \param EndLoc Ending location of the clause.
   /// \param NumVars Number of expressions listed in this clause.
   /// \param NumUniqueDeclarations Number of unique base declarations in this
   /// clause.
   /// \param NumComponentLists Number of component lists in this clause.
   /// \param NumComponents Total number of expression components in the clause.
-  explicit OMPMapClause(OpenMPMapClauseKind MapTypeModifier,
+  explicit OMPMapClause(ArrayRef<OpenMPMapModifierKind> MapModifiers,
+                        ArrayRef<SourceLocation> MapModifiersLoc,
                         OpenMPMapClauseKind MapType, bool MapTypeIsImplicit,
                         SourceLocation MapLoc, SourceLocation StartLoc,
                         SourceLocation LParenLoc, SourceLocation EndLoc,
                         unsigned NumVars, unsigned NumUniqueDeclarations,
                         unsigned NumComponentLists, unsigned NumComponents)
       : OMPMappableExprListClause(OMPC_map, StartLoc, LParenLoc, EndLoc,
                                   NumVars, NumUniqueDeclarations,
                                   NumComponentLists, NumComponents),
-        MapTypeModifier(MapTypeModifier), MapType(MapType),
-        MapTypeIsImplicit(MapTypeIsImplicit), MapLoc(MapLoc) {}
+        MapType(MapType), MapTypeIsImplicit(MapTypeIsImplicit),
+        MapLoc(MapLoc) {
+          assert(llvm::array_lengthof(MapTypeModifiers) == MapModifiers.size()
+                 && "Unexpected number of map type modifiers.");
+          llvm::copy(MapModifiers, std::begin(MapTypeModifiers));
+
+          assert(llvm::array_lengthof(MapTypeModifiersLoc) ==
+                     MapModifiersLoc.size() &&
+                 "Unexpected number of map type modifier locations.");
+          llvm::copy(MapModifiersLoc, std::begin(MapTypeModifiersLoc));
+  }
 
   /// Build an empty clause.
   ///
   /// \param NumVars Number of expressions listed in this clause.
   /// \param NumUniqueDeclarations Number of unique base declarations in this
   /// clause.
   /// \param NumComponentLists Number of component lists in this clause.
   /// \param NumComponents Total number of expression components in the clause.
   explicit OMPMapClause(unsigned NumVars, unsigned NumUniqueDeclarations,
                         unsigned NumComponentLists, unsigned NumComponents)
       : OMPMappableExprListClause(
             OMPC_map, SourceLocation(), SourceLocation(), SourceLocation(),
             NumVars, NumUniqueDeclarations, NumComponentLists, NumComponents) {}
 
-  /// Set type modifier for the clause.
+  /// Set map-type-modifier for the clause.
   ///
-  /// \param T Type Modifier for the clause.
-  void setMapTypeModifier(OpenMPMapClauseKind T) { MapTypeModifier = T; }
+  /// \param I index for map-type-modifier.
+  /// \param T map-type-modifier for the clause.
+  void setMapTypeModifier(unsigned I, OpenMPMapModifierKind T) {
+    assert(I < NumberOfModifiers &&
+           "Unexpected index to store map type modifier, exceeds array size.");
+    MapTypeModifiers[I] = T;
+  }
+
+  /// Set location for the map-type-modifier.
+  ///
+  /// \param I index for map-type-modifier location.
+  /// \param TLoc map-type-modifier location.
+  void setMapTypeModifierLoc(unsigned I, SourceLocation TLoc) {
+    assert(I < NumberOfModifiers &&
+           "Index to store map type modifier location exceeds array size.");
+    MapTypeModifiersLoc[I] = TLoc;
+  }
 
   /// Set type for the clause.
   ///
   /// \param T Type for the clause.
   void setMapType(OpenMPMapClauseKind T) { MapType = T; }
 
   /// Set type location.
   ///
   /// \param TLoc Type location.
   void setMapLoc(SourceLocation TLoc) { MapLoc = TLoc; }
 
   /// Set colon location.
   void setColonLoc(SourceLocation Loc) { ColonLoc = Loc; }
 
 public:
   /// Creates clause with a list of variables \a VL.
   ///
   /// \param C AST context.
   /// \param StartLoc Starting location of the clause.
   /// \param EndLoc Ending location of the clause.
   /// \param Vars The original expression used in the clause.
   /// \param Declarations Declarations used in the clause.
   /// \param ComponentLists Component lists used in the clause.
-  /// \param TypeModifier Map type modifier.
+  /// \param MapModifiers Map-type-modifiers.
+  /// \param MapModifiersLoc Location of map-type-modifiers.
   /// \param Type Map type.
   /// \param TypeIsImplicit Map type is inferred implicitly.
   /// \param TypeLoc Location of the map type.
   static OMPMapClause *Create(const ASTContext &C, SourceLocation StartLoc,
                               SourceLocation LParenLoc, SourceLocation EndLoc,
                               ArrayRef<Expr *> Vars,
                               ArrayRef<ValueDecl *> Declarations,
                               MappableExprComponentListsRef ComponentLists,
-                              OpenMPMapClauseKind TypeModifier,
+                              ArrayRef<OpenMPMapModifierKind> MapModifiers,
+                              ArrayRef<SourceLocation> MapModifiersLoc,
                               OpenMPMapClauseKind Type, bool TypeIsImplicit,
                               SourceLocation TypeLoc);
 
   /// Creates an empty clause with the place for \a NumVars original
   /// expressions, \a NumUniqueDeclarations declarations, \NumComponentLists
   /// lists, and \a NumComponents expression components.
   ///
   /// \param C AST context.
public:
 
   /// Is this an implicit map type?
   /// We have to capture 'IsMapTypeImplicit' from the parser for more
   /// informative error messages.  It helps distinguish map(r) from
   /// map(tofrom: r), which is important to print more helpful error
   /// messages for some target directives.
   bool isImplicitMapType() const LLVM_READONLY { return MapTypeIsImplicit; }
 
-  /// Fetches the map type modifier for the clause.
-  OpenMPMapClauseKind getMapTypeModifier() const LLVM_READONLY {
-    return MapTypeModifier;
+  /// Fetches the map-type-modifier at 'Cnt' index of array of modifiers.
+  ///
+  /// \param Cnt index for map-type-modifier.
+  OpenMPMapModifierKind getMapTypeModifier(unsigned Cnt) const LLVM_READONLY {
+    assert(Cnt < NumberOfModifiers &&
+           "Requested modifier exceeds the total number of modifiers.");
+    return MapTypeModifiers[Cnt];
+  }
+
+  /// Fetches the map-type-modifier location at 'Cnt' index of array of
+  /// modifiers' locations.
+  ///
+  /// \param Cnt index for map-type-modifier location.  
+  SourceLocation getMapTypeModifierLoc(unsigned Cnt) const LLVM_READONLY {
+    assert(Cnt < NumberOfModifiers &&
+           "Requested modifier location exceeds total number of modifiers.");
+    return MapTypeModifiersLoc[Cnt];
+  }
+
+  /// Fetches ArrayRef of map-type-modifiers.
+  ArrayRef<OpenMPMapModifierKind> getMapTypeModifiers() const LLVM_READONLY {
+    return llvm::makeArrayRef(MapTypeModifiers);
+  }
+  
+  /// Fetches ArrayRef of location of map-type-modifiers.
+  ArrayRef<SourceLocation> getMapTypeModifiersLoc() const LLVM_READONLY {
+    return llvm::makeArrayRef(MapTypeModifiersLoc);
   }
 
   /// Fetches location of clause mapping kind.
   SourceLocation getMapLoc() const LLVM_READONLY { return MapLoc; }
 
   /// Get colon location.
   SourceLocation getColonLoc() const { return ColonLoc; }
 

include/clang/Basic/DiagnosticParseKinds.td

def err_omp_expected_identifier_for_critical : Error<
   "expected identifier specifying the name of the 'omp critical' directive">;
 def err_omp_expected_reduction_identifier : Error<
   "expected identifier or one of the following operators: '+', '-', '*', '&', '|', '^', '&&', or '||'">;
 def err_omp_decl_in_declare_simd : Error<
   "function declaration is expected after 'declare simd' 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'">;
+  "incorrect map type modifier, expected 'always' or 'close'">;
 def err_omp_map_type_missing : Error<
   "missing map type">;
+def err_omp_map_type_modifier_missing : Error<
+  "missing map type modifier">;
 def err_omp_declare_simd_inbranch_notinbranch : Error<
   "unexpected '%0' clause, '%1' is specified already">;
 def err_expected_end_declare_target : Error<
   "expected '#pragma omp end declare target'">;
 def err_omp_declare_target_unexpected_clause: Error<
   "unexpected '%0' clause, only 'to' or 'link' clauses expected">;
 def err_omp_expected_clause: Error<
   "expected at least one clause on '#pragma omp %0' directive">;

include/clang/Basic/DiagnosticSemaKinds.td

 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<
   "mapping of union members is not allowed">;
 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_duplicate_map_type_modifier : Error<
+  "same map type modifier has been specified more than once">;
 def err_omp_pointer_mapped_along_with_derived_section : Error<
   "pointer cannot be mapped along with a section derived from itself">;
 def err_omp_original_storage_is_shared_and_does_not_contain : Error<
   "original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage">;
 def err_omp_same_pointer_dereferenced : Error<
   "same pointer dereferenced in multiple different ways in map clause expressions">;
 def note_omp_task_predetermined_firstprivate_here : Note<
   "predetermined as a firstprivate in a task construct here">;

include/clang/Basic/OpenMPKinds.h

 /// OpenMP mapping kind for 'map' clause.
 enum OpenMPMapClauseKind {
 #define OPENMP_MAP_KIND(Name) \
   OMPC_MAP_##Name,
 #include "clang/Basic/OpenMPKinds.def"
   OMPC_MAP_unknown
 };
 
+/// OpenMP modifier kind for 'map' clause.
+enum OpenMPMapModifierKind {
+  OMPC_MAP_MODIFIER_unknown = OMPC_MAP_unknown,
+#define OPENMP_MAP_MODIFIER_KIND(Name) \
+  OMPC_MAP_MODIFIER_##Name,
+#include "clang/Basic/OpenMPKinds.def"
+  OMPC_MAP_MODIFIER_last
+};
+
 /// OpenMP attributes for 'dist_schedule' clause.
 enum OpenMPDistScheduleClauseKind {
 #define OPENMP_DIST_SCHEDULE_KIND(Name) OMPC_DIST_SCHEDULE_##Name,
 #include "clang/Basic/OpenMPKinds.def"
   OMPC_DIST_SCHEDULE_unknown
 };
 
 /// OpenMP attributes for 'defaultmap' clause.

include/clang/Basic/OpenMPKinds.def

 #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
+#ifndef OPENMP_MAP_MODIFIER_KIND
+#define OPENMP_MAP_MODIFIER_KIND(Name)
+#endif
 #ifndef OPENMP_DIST_SCHEDULE_KIND
 #define OPENMP_DIST_SCHEDULE_KIND(Name)
 #endif
 #ifndef OPENMP_DEFAULTMAP_KIND
 #define OPENMP_DEFAULTMAP_KIND(Name)
 #endif
 #ifndef OPENMP_ATOMIC_DEFAULT_MEM_ORDER_KIND
 #define OPENMP_ATOMIC_DEFAULT_MEM_ORDER_KIND(Name)

 OPENMP_TEAMS_CLAUSE(num_teams)
 OPENMP_TEAMS_CLAUSE(thread_limit)
 
 // Clauses allowed for OpenMP directive 'ordered'.
 OPENMP_ORDERED_CLAUSE(threads)
 OPENMP_ORDERED_CLAUSE(simd)
 OPENMP_ORDERED_CLAUSE(depend)
 
-// Map types and map type modifier for 'map' clause.
+// Map types 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)
+
+// Map-type-modifiers for 'map' clause.
+OPENMP_MAP_MODIFIER_KIND(always)
+OPENMP_MAP_MODIFIER_KIND(close)
 
 // Clauses allowed for OpenMP directive 'taskloop'.
 OPENMP_TASKLOOP_CLAUSE(if)
 OPENMP_TASKLOOP_CLAUSE(shared)
 OPENMP_TASKLOOP_CLAUSE(private)
 OPENMP_TASKLOOP_CLAUSE(firstprivate)
 OPENMP_TASKLOOP_CLAUSE(lastprivate)
 OPENMP_TASKLOOP_CLAUSE(default)

 #undef OPENMP_TARGET_EXIT_DATA_CLAUSE
 #undef OPENMP_TARGET_PARALLEL_CLAUSE
 #undef OPENMP_TARGET_PARALLEL_FOR_CLAUSE
 #undef OPENMP_TEAMS_CLAUSE
 #undef OPENMP_SIMD_CLAUSE
 #undef OPENMP_FOR_CLAUSE
 #undef OPENMP_FOR_SIMD_CLAUSE
 #undef OPENMP_MAP_KIND
+#undef OPENMP_MAP_MODIFIER_KIND
 #undef OPENMP_DISTRIBUTE_CLAUSE
 #undef OPENMP_DIST_SCHEDULE_KIND
 #undef OPENMP_DEFAULTMAP_KIND
 #undef OPENMP_DEFAULTMAP_MODIFIER
 #undef OPENMP_TARGET_UPDATE_CLAUSE
 #undef OPENMP_DISTRIBUTE_PARALLEL_FOR_CLAUSE
 #undef OPENMP_DISTRIBUTE_PARALLEL_FOR_SIMD_CLAUSE
 #undef OPENMP_DISTRIBUTE_SIMD_CLAUSE

include/clang/Parse/Parser.h

 //===--- Parser.h - C Language Parser ---------------------------*- C++ -*-===//
 //
 //                     The LLVM Compiler Infrastructure
 //
 // This file is distributed under the University of Illinois Open Source
 // License. See LICENSE.TXT for details.
 //
 //===----------------------------------------------------------------------===//
 //
 //  This file defines the Parser interface.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_CLANG_PARSE_PARSER_H
 #define LLVM_CLANG_PARSE_PARSER_H
 
+#include "clang/AST/OpenMPClause.h"
 #include "clang/AST/Availability.h"
 #include "clang/Basic/BitmaskEnum.h"
 #include "clang/Basic/OpenMPKinds.h"
 #include "clang/Basic/OperatorPrecedence.h"
 #include "clang/Basic/Specifiers.h"
 #include "clang/Lex/CodeCompletionHandler.h"
 #include "clang/Lex/Preprocessor.h"
 #include "clang/Sema/DeclSpec.h"
public:
   struct OpenMPVarListDataTy {
     Expr *TailExpr = nullptr;
     SourceLocation ColonLoc;
     SourceLocation RLoc;
     CXXScopeSpec ReductionIdScopeSpec;
     DeclarationNameInfo ReductionId;
     OpenMPDependClauseKind DepKind = OMPC_DEPEND_unknown;
     OpenMPLinearClauseKind LinKind = OMPC_LINEAR_val;
-    OpenMPMapClauseKind MapTypeModifier = OMPC_MAP_unknown;
+    SmallVector<OpenMPMapModifierKind, OMPMapClause::NumberOfModifiers>
+    MapTypeModifiers;
+    SmallVector<SourceLocation, OMPMapClause::NumberOfModifiers>
+    MapTypeModifiersLoc;
     OpenMPMapClauseKind MapType = OMPC_MAP_unknown;
     bool IsMapTypeImplicit = false;
     SourceLocation DepLinMapLoc;
   };
 
   /// Parses clauses with list.
   bool ParseOpenMPVarList(OpenMPDirectiveKind DKind, OpenMPClauseKind Kind,
                           SmallVectorImpl<Expr *> &Vars,

include/clang/Sema/Sema.h

OMPClause *ActOnOpenMPAtomicDefaultMemOrderClause(
       SourceLocation StartLoc, SourceLocation LParenLoc, 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, OpenMPMapClauseKind MapTypeModifier,
+      OpenMPLinearClauseKind LinKind,
+      ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
+      ArrayRef<SourceLocation> MapTypeModifiersLoc,
       OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
       SourceLocation DepLinMapLoc);
   /// Called on well-formed 'private' clause.
   OMPClause *ActOnOpenMPPrivateClause(ArrayRef<Expr *> VarList,
                                       SourceLocation StartLoc,
                                       SourceLocation LParenLoc,
                                       SourceLocation EndLoc);
   /// Called on well-formed 'firstprivate' clause.
ActOnOpenMPDependClause(OpenMPDependClauseKind DepKind, SourceLocation DepLoc,
                           SourceLocation StartLoc, SourceLocation LParenLoc,
                           SourceLocation EndLoc);
   /// Called on well-formed 'device' clause.
   OMPClause *ActOnOpenMPDeviceClause(Expr *Device, SourceLocation StartLoc,
                                      SourceLocation LParenLoc,
                                      SourceLocation EndLoc);
   /// Called on well-formed 'map' clause.
   OMPClause *
-  ActOnOpenMPMapClause(OpenMPMapClauseKind MapTypeModifier,
+  ActOnOpenMPMapClause(ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
+                       ArrayRef<SourceLocation> MapTypeModifiersLoc,
                        OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
                        SourceLocation MapLoc, SourceLocation ColonLoc,
                        ArrayRef<Expr *> VarList, SourceLocation StartLoc,
                        SourceLocation LParenLoc, SourceLocation EndLoc);
   /// Called on well-formed 'num_teams' clause.
   OMPClause *ActOnOpenMPNumTeamsClause(Expr *NumTeams, SourceLocation StartLoc,
                                        SourceLocation LParenLoc,
                                        SourceLocation EndLoc);

lib/AST/OpenMPClause.cpp

unsigned OMPClauseMappableExprCommon::getUniqueDeclarationsTotalNumber(
   return TotalNum;
 }
 
 OMPMapClause *
 OMPMapClause::Create(const ASTContext &C, SourceLocation StartLoc,
                      SourceLocation LParenLoc, SourceLocation EndLoc,
                      ArrayRef<Expr *> Vars, ArrayRef<ValueDecl *> Declarations,
                      MappableExprComponentListsRef ComponentLists,
-                     OpenMPMapClauseKind TypeModifier, OpenMPMapClauseKind Type,
-                     bool TypeIsImplicit, SourceLocation TypeLoc) {
+                     ArrayRef<OpenMPMapModifierKind> MapModifiers,
+                     ArrayRef<SourceLocation> MapModifiersLoc,
+                     OpenMPMapClauseKind Type, bool TypeIsImplicit,
+                     SourceLocation TypeLoc) {
   unsigned NumVars = Vars.size();
   unsigned NumUniqueDeclarations =
       getUniqueDeclarationsTotalNumber(Declarations);
   unsigned NumComponentLists = ComponentLists.size();
   unsigned NumComponents = getComponentsTotalNumber(ComponentLists);
 
   // We need to allocate:
   // NumVars x Expr* - we have an original list expression for each clause list
   // entry.
   // NumUniqueDeclarations x ValueDecl* - unique base declarations associated
   // with each component list.
   // (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the
   // number of lists for each unique declaration and the size of each component
   // list.
   // NumComponents x MappableComponent - the total of all the components in all
   // the lists.
   void *Mem = C.Allocate(
       totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
           NumVars, NumUniqueDeclarations,
           NumUniqueDeclarations + NumComponentLists, NumComponents));
   OMPMapClause *Clause = new (Mem) OMPMapClause(
-      TypeModifier, Type, TypeIsImplicit, TypeLoc, StartLoc, LParenLoc, EndLoc,
-      NumVars, NumUniqueDeclarations, NumComponentLists, NumComponents);
+      MapModifiers, MapModifiersLoc, Type, TypeIsImplicit, TypeLoc, StartLoc,
+      LParenLoc, EndLoc, NumVars, NumUniqueDeclarations, NumComponentLists,
+      NumComponents);
 
   Clause->setVarRefs(Vars);
   Clause->setClauseInfo(Declarations, ComponentLists);
-  Clause->setMapTypeModifier(TypeModifier);
   Clause->setMapType(Type);
   Clause->setMapLoc(TypeLoc);
   return Clause;
 }
 
 OMPMapClause *OMPMapClause::CreateEmpty(const ASTContext &C, unsigned NumVars,
                                         unsigned NumUniqueDeclarations,
                                         unsigned NumComponentLists,
void OMPClausePrinter::VisitOMPDependClause(OMPDependClause *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) {
+      for (unsigned I = 0; I < OMPMapClause::NumberOfModifiers; ++I) {
+        if (Node->getMapTypeModifier(I) != OMPC_MAP_MODIFIER_unknown) {
           OS << getOpenMPSimpleClauseTypeName(OMPC_map,
-                                            Node->getMapTypeModifier());
+                                              Node->getMapTypeModifier(I));
           OS << ',';
         }
+      }
       OS << getOpenMPSimpleClauseTypeName(OMPC_map, Node->getMapType());
       OS << ':';
     }
     VisitOMPClauseList(Node, ' ');
     OS << ")";
   }
 }
 

lib/Basic/OpenMPKinds.cpp

 #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)
+    return llvm::StringSwitch<unsigned>(Str)
+#define OPENMP_MAP_KIND(Name)                                                  \
+  .Case(#Name, static_cast<unsigned>(OMPC_MAP_##Name))
+#define OPENMP_MAP_MODIFIER_KIND(Name)                                         \
+  .Case(#Name, static_cast<unsigned>(OMPC_MAP_MODIFIER_##Name))
 #include "clang/Basic/OpenMPKinds.def"
         .Default(OMPC_MAP_unknown);
   case OMPC_dist_schedule:
     return llvm::StringSwitch<OpenMPDistScheduleClauseKind>(Str)
 #define OPENMP_DIST_SCHEDULE_KIND(Name) .Case(#Name, OMPC_DIST_SCHEDULE_##Name)
 #include "clang/Basic/OpenMPKinds.def"
         .Default(OMPC_DIST_SCHEDULE_unknown);
   case OMPC_defaultmap:
#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:
+    case OMPC_MAP_MODIFIER_last:
       return "unknown";
 #define OPENMP_MAP_KIND(Name)                                                \
   case OMPC_MAP_##Name:                                                      \
     return #Name;
+#define OPENMP_MAP_MODIFIER_KIND(Name)                                       \
+  case OMPC_MAP_MODIFIER_##Name:                                             \
+    return #Name;
 #include "clang/Basic/OpenMPKinds.def"
     default:
       break;
     }
     llvm_unreachable("Invalid OpenMP 'map' clause type");
   case OMPC_dist_schedule:
     switch (Type) {
     case OMPC_DIST_SCHEDULE_unknown:

lib/CodeGen/CGOpenMPRuntime.cpp

struct StructRangeInfoTy {
     Address Base = Address::invalid();
   };
 
 private:
   /// Kind that defines how a device pointer has to be returned.
   struct MapInfo {
     OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
     OpenMPMapClauseKind MapType = OMPC_MAP_unknown;
-    OpenMPMapClauseKind MapTypeModifier = OMPC_MAP_unknown;
+    ArrayRef<OpenMPMapModifierKind> MapModifiers;
     bool ReturnDevicePointer = false;
     bool IsImplicit = false;
 
     MapInfo() = default;
     MapInfo(
         OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
-        OpenMPMapClauseKind MapType, OpenMPMapClauseKind MapTypeModifier,
+        OpenMPMapClauseKind MapType,
+        ArrayRef<OpenMPMapModifierKind> MapModifiers,
         bool ReturnDevicePointer, bool IsImplicit)
-        : Components(Components), MapType(MapType),
-          MapTypeModifier(MapTypeModifier),
+        : Components(Components), MapType(MapType), MapModifiers(MapModifiers),
           ReturnDevicePointer(ReturnDevicePointer), IsImplicit(IsImplicit) {}
   };
 
   /// If use_device_ptr is used on a pointer which is a struct member and there
   /// is no map information about it, then emission of that entry is deferred
   /// until the whole struct has been processed.
   struct DeferredDevicePtrEntryTy {
     const Expr *IE = nullptr;
llvm::Value *getExprTypeSize(const Expr *E) const {
     }
     return CGF.getTypeSize(ExprTy);
   }
 
   /// Return the corresponding bits for a given map clause modifier. Add
   /// a flag marking the map as a pointer if requested. Add a flag marking the
   /// map as the first one of a series of maps that relate to the same map
   /// expression.
-  OpenMPOffloadMappingFlags getMapTypeBits(OpenMPMapClauseKind MapType,
-                                           OpenMPMapClauseKind MapTypeModifier,
-                                           bool IsImplicit, bool AddPtrFlag,
-                                           bool AddIsTargetParamFlag) const {
+  OpenMPOffloadMappingFlags getMapTypeBits(
+      OpenMPMapClauseKind MapType, ArrayRef<OpenMPMapModifierKind> MapModifiers,
+      bool IsImplicit, bool AddPtrFlag, bool AddIsTargetParamFlag) const {
     OpenMPOffloadMappingFlags Bits =
         IsImplicit ? OMP_MAP_IMPLICIT : OMP_MAP_NONE;
     switch (MapType) {
     case OMPC_MAP_alloc:
     case OMPC_MAP_release:
       // alloc and release is the default behavior in the runtime library,  i.e.
       // if we don't pass any bits alloc/release that is what the runtime is
       // going to do. Therefore, we don't need to signal anything for these two
       // type modifiers.
       break;
     case OMPC_MAP_to:
       Bits |= OMP_MAP_TO;
       break;
     case OMPC_MAP_from:
       Bits |= OMP_MAP_FROM;
       break;
     case OMPC_MAP_tofrom:
       Bits |= OMP_MAP_TO | OMP_MAP_FROM;
       break;
     case OMPC_MAP_delete:
       Bits |= OMP_MAP_DELETE;
       break;
-    case OMPC_MAP_always:
     case OMPC_MAP_unknown:
       llvm_unreachable("Unexpected map type!");
     }
     if (AddPtrFlag)
       Bits |= OMP_MAP_PTR_AND_OBJ;
     if (AddIsTargetParamFlag)
       Bits |= OMP_MAP_TARGET_PARAM;
-    if (MapTypeModifier == OMPC_MAP_always)
+    if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_always)
+        != MapModifiers.end())
       Bits |= OMP_MAP_ALWAYS;
     return Bits;
   }
 
   /// Return true if the provided expression is a final array section. A
   /// final array section, is one whose length can't be proved to be one.
   bool isFinalArraySectionExpression(const Expr *E) const {
     const auto *OASE = dyn_cast<OMPArraySectionExpr>(E);
bool isFinalArraySectionExpression(const Expr *E) const {
     return ConstLength.getSExtValue() != 1;
   }
 
   /// Generate the base pointers, section pointers, sizes and map type
   /// bits for the provided map type, map modifier, and expression components.
   /// \a IsFirstComponent should be set to true if the provided set of
   /// components is the first associated with a capture.
   void generateInfoForComponentList(
-      OpenMPMapClauseKind MapType, OpenMPMapClauseKind MapTypeModifier,
+      OpenMPMapClauseKind MapType,
+      ArrayRef<OpenMPMapModifierKind> MapModifiers,
       OMPClauseMappableExprCommon::MappableExprComponentListRef Components,
       MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers,
       MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types,
       StructRangeInfoTy &PartialStruct, bool IsFirstComponentList,
       bool IsImplicit,
       ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
           OverlappedElements = llvm::None) const {
     // The following summarizes what has to be generated for each map and the
for (; I != CE; ++I) {
           PartialStruct.HighestElem = {
               std::numeric_limits<decltype(
                   PartialStruct.HighestElem.first)>::max(),
               HB};
           PartialStruct.Base = BP;
           // Emit data for non-overlapped data.
           OpenMPOffloadMappingFlags Flags =
               OMP_MAP_MEMBER_OF |
-              getMapTypeBits(MapType, MapTypeModifier, IsImplicit,
+              getMapTypeBits(MapType, MapModifiers, IsImplicit,
                              /*AddPtrFlag=*/false,
                              /*AddIsTargetParamFlag=*/false);
           LB = BP;
           llvm::Value *Size = nullptr;
           // Do bitcopy of all non-overlapped structure elements.
           for (OMPClauseMappableExprCommon::MappableExprComponentListRef
                    Component : OverlappedElements) {
             Address ComponentLB = Address::invalid();
for (; I != CE; ++I) {
           Pointers.push_back(LB.getPointer());
           Sizes.push_back(Size);
 
           // We need to add a pointer flag for each map that comes from the
           // same expression except for the first one. We also need to signal
           // this map is the first one that relates with the current capture
           // (there is a set of entries for each capture).
           OpenMPOffloadMappingFlags Flags = getMapTypeBits(
-              MapType, MapTypeModifier, IsImplicit,
+              MapType, MapModifiers, IsImplicit,
               !IsExpressionFirstInfo || IsLink, IsCaptureFirstInfo && !IsLink);
 
           if (!IsExpressionFirstInfo) {
             // If we have a PTR_AND_OBJ pair where the OBJ is a pointer as well,
             // then we reset the TO/FROM/ALWAYS/DELETE flags.
             if (IsPointer)
               Flags &= ~(OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_ALWAYS |
                          OMP_MAP_DELETE);
void generateAllInfo(MapBaseValuesArrayTy &BasePointers,
     // correctly. Therefore, we organize all lists in a map.
     llvm::MapVector<const ValueDecl *, SmallVector<MapInfo, 8>> Info;
 
     // Helper function to fill the information map for the different supported
     // clauses.
     auto &&InfoGen = [&Info](
         const ValueDecl *D,
         OMPClauseMappableExprCommon::MappableExprComponentListRef L,
-        OpenMPMapClauseKind MapType, OpenMPMapClauseKind MapModifier,
+        OpenMPMapClauseKind MapType,
+        ArrayRef<OpenMPMapModifierKind> MapModifiers,
         bool ReturnDevicePointer, bool IsImplicit) {
       const ValueDecl *VD =
           D ? cast<ValueDecl>(D->getCanonicalDecl()) : nullptr;
-      Info[VD].emplace_back(L, MapType, MapModifier, ReturnDevicePointer,
+      Info[VD].emplace_back(L, MapType, MapModifiers, ReturnDevicePointer,
                             IsImplicit);
     };
 
     // FIXME: MSVC 2013 seems to require this-> to find member CurDir.
     for (const auto *C : this->CurDir.getClausesOfKind<OMPMapClause>())
       for (const auto &L : C->component_lists()) {
-        InfoGen(L.first, L.second, C->getMapType(), C->getMapTypeModifier(),
+        InfoGen(L.first, L.second, C->getMapType(), C->getMapTypeModifiers(),
             /*ReturnDevicePointer=*/false, C->isImplicit());
       }
     for (const auto *C : this->CurDir.getClausesOfKind<OMPToClause>())
       for (const auto &L : C->component_lists()) {
-        InfoGen(L.first, L.second, OMPC_MAP_to, OMPC_MAP_unknown,
+        InfoGen(L.first, L.second, OMPC_MAP_to, llvm::None,
             /*ReturnDevicePointer=*/false, C->isImplicit());
       }
     for (const auto *C : this->CurDir.getClausesOfKind<OMPFromClause>())
       for (const auto &L : C->component_lists()) {
-        InfoGen(L.first, L.second, OMPC_MAP_from, OMPC_MAP_unknown,
+        InfoGen(L.first, L.second, OMPC_MAP_from, llvm::None,
             /*ReturnDevicePointer=*/false, C->isImplicit());
       }
 
     // Look at the use_device_ptr clause information and mark the existing map
     // entries as such. If there is no map information for an entry in the
     // use_device_ptr list, we create one with map type 'alloc' and zero size
     // section. It is the user fault if that was not mapped before. If there is
     // no map information and the pointer is a struct member, then we defer the
for (const auto *C :
         if (isa<MemberExpr>(IE)) {
           // Insert the pointer into Info to be processed by
           // generateInfoForComponentList. Because it is a member pointer
           // without a pointee, no entry will be generated for it, therefore
           // we need to generate one after the whole struct has been processed.
           // Nonetheless, generateInfoForComponentList must be called to take
           // the pointer into account for the calculation of the range of the
           // partial struct.
-          InfoGen(nullptr, L.second, OMPC_MAP_unknown, OMPC_MAP_unknown,
+          InfoGen(nullptr, L.second, OMPC_MAP_unknown, llvm::None,
                   /*ReturnDevicePointer=*/false, C->isImplicit());
           DeferredInfo[nullptr].emplace_back(IE, VD);
         } else {
           llvm::Value *Ptr = this->CGF.EmitLoadOfScalar(
               this->CGF.EmitLValue(IE), IE->getExprLoc());
           BasePointers.emplace_back(Ptr, VD);
           Pointers.push_back(Ptr);
           Sizes.push_back(llvm::Constant::getNullValue(this->CGF.SizeTy));
for (const auto &M : Info) {
       for (const MapInfo &L : M.second) {
         assert(!L.Components.empty() &&
                "Not expecting declaration with no component lists.");
 
         // Remember the current base pointer index.
         unsigned CurrentBasePointersIdx = CurBasePointers.size();
         // FIXME: MSVC 2013 seems to require this-> to find the member method.
         this->generateInfoForComponentList(
-            L.MapType, L.MapTypeModifier, L.Components, CurBasePointers,
+            L.MapType, L.MapModifiers, L.Components, CurBasePointers,
             CurPointers, CurSizes, CurTypes, PartialStruct,
             IsFirstComponentList, L.IsImplicit);
 
         // If this entry relates with a device pointer, set the relevant
         // declaration and add the 'return pointer' flag.
         if (L.ReturnDevicePointer) {
           assert(CurBasePointers.size() > CurrentBasePointersIdx &&
                  "Unexpected number of mapped base pointers.");
if (DevPointersMap.count(VD)) {
       Pointers.push_back(Arg);
       Sizes.push_back(CGF.getTypeSize(CGF.getContext().VoidPtrTy));
       Types.push_back(OMP_MAP_LITERAL | OMP_MAP_TARGET_PARAM);
       return;
     }
 
     using MapData =
         std::tuple<OMPClauseMappableExprCommon::MappableExprComponentListRef,
-                   OpenMPMapClauseKind, OpenMPMapClauseKind, bool>;
+                   OpenMPMapClauseKind, ArrayRef<OpenMPMapModifierKind>, bool>;
     SmallVector<MapData, 4> DeclComponentLists;
     // FIXME: MSVC 2013 seems to require this-> to find member CurDir.
     for (const auto *C : this->CurDir.getClausesOfKind<OMPMapClause>()) {
       for (const auto &L : C->decl_component_lists(VD)) {
         assert(L.first == VD &&
                "We got information for the wrong declaration??");
         assert(!L.second.empty() &&
                "Not expecting declaration with no component lists.");
         DeclComponentLists.emplace_back(L.second, C->getMapType(),
-                                        C->getMapTypeModifier(),
+                                        C->getMapTypeModifiers(),
                                         C->isImplicit());
       }
     }
 
     // Find overlapping elements (including the offset from the base element).
     llvm::SmallDenseMap<
         const MapData *,
         llvm::SmallVector<
             OMPClauseMappableExprCommon::MappableExprComponentListRef, 4>,
         4>
         OverlappedData;
     size_t Count = 0;
     for (const MapData &L : DeclComponentLists) {
       OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
       OpenMPMapClauseKind MapType;
-      OpenMPMapClauseKind MapTypeModifier;
+      ArrayRef<OpenMPMapModifierKind> MapModifiers;
       bool IsImplicit;
-      std::tie(Components, MapType, MapTypeModifier, IsImplicit) = L;
+      std::tie(Components, MapType, MapModifiers, IsImplicit) = L;
       ++Count;
       for (const MapData &L1 : makeArrayRef(DeclComponentLists).slice(Count)) {
         OMPClauseMappableExprCommon::MappableExprComponentListRef Components1;
-        std::tie(Components1, MapType, MapTypeModifier, IsImplicit) = L1;
+        std::tie(Components1, MapType, MapModifiers, IsImplicit) = L1;
         auto CI = Components.rbegin();
         auto CE = Components.rend();
         auto SI = Components1.rbegin();
         auto SE = Components1.rend();
         for (; CI != CE && SI != SE; ++CI, ++SI) {
           if (CI->getAssociatedExpression()->getStmtClass() !=
               SI->getAssociatedExpression()->getStmtClass())
             break;
void generateInfoForCapture(const CapturedStmt::Capture *Cap,
     }
 
     // Associated with a capture, because the mapping flags depend on it.
     // Go through all of the elements with the overlapped elements.
     for (const auto &Pair : OverlappedData) {
       const MapData &L = *Pair.getFirst();
       OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
       OpenMPMapClauseKind MapType;
-      OpenMPMapClauseKind MapTypeModifier;
+      ArrayRef<OpenMPMapModifierKind> MapModifiers;
       bool IsImplicit;
-      std::tie(Components, MapType, MapTypeModifier, IsImplicit) = L;
+      std::tie(Components, MapType, MapModifiers, IsImplicit) = L;
       ArrayRef<OMPClauseMappableExprCommon::MappableExprComponentListRef>
           OverlappedComponents = Pair.getSecond();
       bool IsFirstComponentList = true;
-      generateInfoForComponentList(MapType, MapTypeModifier, Components,
+      generateInfoForComponentList(MapType, MapModifiers, Components,
                                    BasePointers, Pointers, Sizes, Types,
                                    PartialStruct, IsFirstComponentList,
                                    IsImplicit, OverlappedComponents);
     }
     // Go through other elements without overlapped elements.
     bool IsFirstComponentList = OverlappedData.empty();
     for (const MapData &L : DeclComponentLists) {
       OMPClauseMappableExprCommon::MappableExprComponentListRef Components;
       OpenMPMapClauseKind MapType;
-      OpenMPMapClauseKind MapTypeModifier;
+      ArrayRef<OpenMPMapModifierKind> MapModifiers;
       bool IsImplicit;
-      std::tie(Components, MapType, MapTypeModifier, IsImplicit) = L;
+      std::tie(Components, MapType, MapModifiers, IsImplicit) = L;
       auto It = OverlappedData.find(&L);
       if (It == OverlappedData.end())
-        generateInfoForComponentList(MapType, MapTypeModifier, Components,
+        generateInfoForComponentList(MapType, MapModifiers, Components,
                                      BasePointers, Pointers, Sizes, Types,
                                      PartialStruct, IsFirstComponentList,
                                      IsImplicit);
       IsFirstComponentList = false;
     }
   }
 
   /// Generate the base pointers, section pointers, sizes and map types
for (const auto *C : this->CurDir.getClausesOfKind<OMPMapClause>()) {
         if (!VD)
           continue;
         llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
             OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD);
         if (!Res || *Res != OMPDeclareTargetDeclAttr::MT_Link)
           continue;
         StructRangeInfoTy PartialStruct;
         generateInfoForComponentList(
-            C->getMapType(), C->getMapTypeModifier(), L.second, BasePointers,
+            C->getMapType(), C->getMapTypeModifiers(), L.second, BasePointers,
             Pointers, Sizes, Types, PartialStruct,
             /*IsFirstComponentList=*/true, C->isImplicit());
         assert(!PartialStruct.Base.isValid() &&
                "No partial structs for declare target link expected.");
       }
     }
   }
 

lib/Parse/ParseOpenMP.cpp

static bool ParseReductionId(Parser &P, CXXScopeSpec &ReductionIdScopeSpec,
   }
   return P.ParseUnqualifiedId(ReductionIdScopeSpec, /*EnteringContext*/ false,
                               /*AllowDestructorName*/ false,
                               /*AllowConstructorName*/ false,
                               /*AllowDeductionGuide*/ false,
                               nullptr, nullptr, ReductionId);
 }
 
+/// Checks if the token is a valid map-type-modifier.
+static OpenMPMapModifierKind isMapModifier(Parser &P) {
+  Token Tok = P.getCurToken();
+  if (!Tok.is(tok::identifier))
+    return OMPC_MAP_MODIFIER_unknown;
+
+  Preprocessor &PP = P.getPreprocessor();
+  OpenMPMapModifierKind TypeModifier = static_cast<OpenMPMapModifierKind>(
+      getOpenMPSimpleClauseType(OMPC_map, PP.getSpelling(Tok)));
+  return TypeModifier;
+}
+
+/// Parse map-type-modifiers in map clause.
+/// map([ [map-type-modifier[,] [map-type-modifier[,] ...] map-type : ] list)
+/// where, map-type-modifier ::= always | close 
+static void parseMapTypeModifiers(Parser &P,
+                                  Parser::OpenMPVarListDataTy &Data) {
+  Preprocessor &PP = P.getPreprocessor();
+  while (P.getCurToken().isNot(tok::colon)) {
+    Token Tok = P.getCurToken();
+    OpenMPMapModifierKind TypeModifier = isMapModifier(P);
+    if (TypeModifier == OMPC_MAP_MODIFIER_always ||
+        TypeModifier == OMPC_MAP_MODIFIER_close) {
+      Data.MapTypeModifiers.push_back(TypeModifier);
+      Data.MapTypeModifiersLoc.push_back(Tok.getLocation());
+      P.ConsumeToken();
+    } else {
+      // For the case of unknown map-type-modifier or a map-type.
+      // Map-type is followed by a colon; the function returns when it
+      // encounters a token followed by a colon.
+      if (Tok.is(tok::comma)) {
+        P.Diag(Tok, diag::err_omp_map_type_modifier_missing);
+        P.ConsumeToken();
+        continue;
+      }
+      // Potential map-type token as it is followed by a colon.
+      if (PP.LookAhead(0).is(tok::colon))
+        return;
+      P.Diag(Tok, diag::err_omp_unknown_map_type_modifier);
+      P.ConsumeToken();
+    }
+    if (P.getCurToken().is(tok::comma))
+      P.ConsumeToken();
+  }
+}
+
+/// Checks if the token is a valid map-type.
+static OpenMPMapClauseKind isMapType(Parser &P) {
+  Token Tok = P.getCurToken();
+  // The map-type token can be either an identifier or the C++ delete keyword.
+  if (!Tok.isOneOf(tok::identifier, tok::kw_delete))
+    return OMPC_MAP_unknown;
+  Preprocessor &PP = P.getPreprocessor();
+  OpenMPMapClauseKind MapType = static_cast<OpenMPMapClauseKind>(
+      getOpenMPSimpleClauseType(OMPC_map, PP.getSpelling(Tok)));
+  return MapType;
+}
+
+/// Parse map-type in map clause.
+/// map([ [map-type-modifier[,] [map-type-modifier[,] ...] map-type : ] list)
+/// where, map-type ::= to | from | tofrom | alloc | release | delete 
+static void parseMapType(Parser &P, Parser::OpenMPVarListDataTy &Data) {
+  Token Tok = P.getCurToken();
+  if (Tok.is(tok::colon)) {
+    P.Diag(Tok, diag::err_omp_map_type_missing);
+    return;
+  }
+  Data.MapType = isMapType(P);
+  if (Data.MapType == OMPC_MAP_unknown)
+    P.Diag(Tok, diag::err_omp_unknown_map_type);
+  P.ConsumeToken();
+}
+
 /// Parses clauses with list.
 bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
                                 OpenMPClauseKind Kind,
                                 SmallVectorImpl<Expr *> &Vars,
                                 OpenMPVarListDataTy &Data) {
   UnqualifiedId UnqualifiedReductionId;
   bool InvalidReductionId = false;
 
if (Tok.is(tok::identifier) && PP.LookAhead(0).is(tok::l_paren)) {
       Data.DepLinMapLoc = ConsumeToken();
       LinearT.consumeOpen();
       NeedRParenForLinear = true;
     }
   } else if (Kind == OMPC_map) {
     // Handle map type for map clause.
     ColonProtectionRAIIObject ColonRAII(*this);
 
-    /// The map clause modifier token can be either a identifier or the C++
-    /// delete keyword.
-    auto &&IsMapClauseModifierToken = [](const Token &Tok) -> bool {
-      return Tok.isOneOf(tok::identifier, tok::kw_delete);
-    };
-
     // The first identifier may be a list item, a map-type or a
-    // map-type-modifier. The map modifier can also be delete which has the same
+    // map-type-modifier. The map-type can also be delete which has the same
     // spelling of the C++ delete keyword.
-    Data.MapType =
-        IsMapClauseModifierToken(Tok)
-            ? static_cast<OpenMPMapClauseKind>(
-                  getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok)))
-            : OMPC_MAP_unknown;
     Data.DepLinMapLoc = Tok.getLocation();
 
-    if (IsMapClauseModifierToken(Tok)) {
-      if (PP.LookAhead(0).is(tok::colon)) {
-        if (Data.MapType == OMPC_MAP_unknown)
-          Diag(Tok, diag::err_omp_unknown_map_type);
-        else if (Data.MapType == OMPC_MAP_always)
-          Diag(Tok, diag::err_omp_map_type_missing);
-        ConsumeToken();
-      } else if (PP.LookAhead(0).is(tok::comma)) {
-        if (IsMapClauseModifierToken(PP.LookAhead(1)) &&
-            PP.LookAhead(2).is(tok::colon)) {
-          Data.MapTypeModifier = Data.MapType;
-          if (Data.MapTypeModifier != OMPC_MAP_always) {
-            Diag(Tok, diag::err_omp_unknown_map_type_modifier);
-            Data.MapTypeModifier = OMPC_MAP_unknown;
-          }
-
-          ConsumeToken();
-          ConsumeToken();
-
-          Data.MapType =
-              IsMapClauseModifierToken(Tok)
-                  ? static_cast<OpenMPMapClauseKind>(
-                        getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok)))
-                  : OMPC_MAP_unknown;
-          if (Data.MapType == OMPC_MAP_unknown ||
-              Data.MapType == OMPC_MAP_always)
-            Diag(Tok, diag::err_omp_unknown_map_type);
-          ConsumeToken();
-        } else {
-          Data.MapType = OMPC_MAP_tofrom;
-          Data.IsMapTypeImplicit = true;
-        }
-      } else if (IsMapClauseModifierToken(PP.LookAhead(0))) {
-        if (PP.LookAhead(1).is(tok::colon)) {
-          Data.MapTypeModifier = Data.MapType;
-          if (Data.MapTypeModifier != OMPC_MAP_always) {
-            Diag(Tok, diag::err_omp_unknown_map_type_modifier);
-            Data.MapTypeModifier = OMPC_MAP_unknown;
-          }
-
-          ConsumeToken();
-
-          Data.MapType =
-              IsMapClauseModifierToken(Tok)
-                  ? static_cast<OpenMPMapClauseKind>(
-                        getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok)))
-                  : OMPC_MAP_unknown;
-          if (Data.MapType == OMPC_MAP_unknown ||
-              Data.MapType == OMPC_MAP_always)
-            Diag(Tok, diag::err_omp_unknown_map_type);
-          ConsumeToken();
-        } else {
-          Data.MapType = OMPC_MAP_tofrom;
-          Data.IsMapTypeImplicit = true;
-        }
-      } else {
-        Data.MapType = OMPC_MAP_tofrom;
-        Data.IsMapTypeImplicit = true;
-      }
-    } else {
+    // Check for presence of a colon in the map clause.
+    TentativeParsingAction TPA(*this);
+    bool ColonPresent = false;
+    if (SkipUntil(tok::colon, tok::r_paren, tok::annot_pragma_openmp_end,
+        StopBeforeMatch)) {
+      if (Tok.is(tok::colon))
+        ColonPresent = true;
+    TPA.Revert();
+    // Only parse map-type-modifier[s] and map-type if a colon is present in
+    // the map clause.
+    if (ColonPresent) {
+      parseMapTypeModifiers(*this, Data);
+      parseMapType(*this, Data);
+    if (Data.MapType == OMPC_MAP_unknown) {
       Data.MapType = OMPC_MAP_tofrom;
       Data.IsMapTypeImplicit = true;
     }
 
     if (Tok.is(tok::colon))
       Data.ColonLoc = ConsumeToken();
   }
 

 ///       'in_reduction' '(' reduction-identifier ':' list ')'
 ///    copyprivate-clause:
 ///       'copyprivate' '(' list ')'
 ///    flush-clause:
 ///       'flush' '(' list ')'
 ///    depend-clause:
 ///       'depend' '(' in | out | inout : list | source ')'
 ///    map-clause:
-///       'map' '(' [ [ always , ]
+///       'map' '(' [ [ always [,] ] [ close [,] ]
 ///          to | from | tofrom | alloc | release | delete ':' ] list ')';
 ///    to-clause:
 ///       'to' '(' list ')'
 ///    from-clause:
 ///       'from' '(' list ')'
 ///    use_device_ptr-clause:
 ///       'use_device_ptr' '(' list ')'
 ///    is_device_ptr-clause:
OMPClause *Parser::ParseOpenMPVarListClause(OpenMPDirectiveKind DKind,
   if (ParseOpenMPVarList(DKind, Kind, Vars, Data))
     return nullptr;
 
   if (ParseOnly)
     return nullptr;
   return Actions.ActOnOpenMPVarListClause(
       Kind, Vars, Data.TailExpr, Loc, LOpen, Data.ColonLoc, Data.RLoc,
       Data.ReductionIdScopeSpec, Data.ReductionId, Data.DepKind, Data.LinKind,
-      Data.MapTypeModifier, Data.MapType, Data.IsMapTypeImplicit,
-      Data.DepLinMapLoc);
+      Data.MapTypeModifiers, Data.MapTypeModifiersLoc, Data.MapType,
+      Data.IsMapTypeImplicit, Data.DepLinMapLoc);
 }
 

lib/Sema/SemaOpenMP.cpp

if (!ImplicitFirstprivates.empty()) {
         ErrorFound = cast<OMPFirstprivateClause>(Implicit)->varlist_size() !=
                      ImplicitFirstprivates.size();
       } else {
         ErrorFound = true;
       }
     }
     if (!ImplicitMaps.empty()) {
       if (OMPClause *Implicit = ActOnOpenMPMapClause(
-              OMPC_MAP_unknown, OMPC_MAP_tofrom, /*IsMapTypeImplicit=*/true,
-              SourceLocation(), SourceLocation(), ImplicitMaps,
-              SourceLocation(), SourceLocation(), SourceLocation())) {
+              llvm::None, llvm::None, OMPC_MAP_tofrom,
+              /*IsMapTypeImplicit=*/true, SourceLocation(), SourceLocation(),
+              ImplicitMaps, SourceLocation(), SourceLocation(),
+              SourceLocation())) {
         ClausesWithImplicit.emplace_back(Implicit);
         ErrorFound |=
             cast<OMPMapClause>(Implicit)->varlist_size() != ImplicitMaps.size();
       } else {
         ErrorFound = true;
       }
     }
   }
OMPClause *Sema::ActOnOpenMPDynamicAllocatorsClause(SourceLocation StartLoc,
   return new (Context) OMPDynamicAllocatorsClause(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, OpenMPMapClauseKind MapTypeModifier,
+    OpenMPLinearClauseKind LinKind,
+    ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
+    ArrayRef<SourceLocation> MapTypeModifiersLoc,
     OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
     SourceLocation DepLinMapLoc) {
   OMPClause *Res = nullptr;
   switch (Kind) {
   case OMPC_private:
     Res = ActOnOpenMPPrivateClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_firstprivate:
OMPClause *Sema::ActOnOpenMPVarListClause(
   case OMPC_flush:
     Res = ActOnOpenMPFlushClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   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);
+    Res = ActOnOpenMPMapClause(MapTypeModifiers, MapTypeModifiersLoc, MapType,
+                               IsMapTypeImplicit, DepLinMapLoc, ColonLoc,
+                               VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_to:
     Res = ActOnOpenMPToClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_from:
     Res = ActOnOpenMPFromClause(VarList, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_use_device_ptr:
for (Expr *RE : MVLI.VarList) {
     MVLI.VarComponents.back().append(CurComponents.begin(),
                                      CurComponents.end());
     MVLI.VarBaseDeclarations.push_back(isa<MemberExpr>(BE) ? nullptr
                                                            : CurDeclaration);
   }
 }
 
 OMPClause *
-Sema::ActOnOpenMPMapClause(OpenMPMapClauseKind MapTypeModifier,
+Sema::ActOnOpenMPMapClause(ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
+                           ArrayRef<SourceLocation> MapTypeModifiersLoc,
                            OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
                            SourceLocation MapLoc, SourceLocation ColonLoc,
                            ArrayRef<Expr *> VarList, SourceLocation StartLoc,
                            SourceLocation LParenLoc, SourceLocation EndLoc) {
   MappableVarListInfo MVLI(VarList);
   checkMappableExpressionList(*this, DSAStack, OMPC_map, MVLI, StartLoc,
                               MapType, IsMapTypeImplicit);
 
+  OpenMPMapModifierKind Modifiers[] = { OMPC_MAP_MODIFIER_unknown,
+                                        OMPC_MAP_MODIFIER_unknown };
+  SourceLocation ModifiersLoc[OMPMapClause::NumberOfModifiers];
+
+  // Process map-type-modifiers, flag errors for duplicate modifiers.
+  unsigned Count = 0;
+  for (unsigned I = 0, E = MapTypeModifiers.size(); I < E; ++I) {
+    if (MapTypeModifiers[I] != OMPC_MAP_MODIFIER_unknown &&
+        llvm::find(Modifiers, MapTypeModifiers[I]) != std::end(Modifiers)) {
+      Diag(MapTypeModifiersLoc[I], diag::err_omp_duplicate_map_type_modifier);
+      continue;
+    }
+    assert(Count < OMPMapClause::NumberOfModifiers &&
+           "Modifiers exceed the allowed number of map type modifiers"); 
+    Modifiers[Count] = MapTypeModifiers[I];
+    ModifiersLoc[Count] = MapTypeModifiersLoc[I];
+    ++Count;
+  }
+
   // 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,
                               MVLI.ProcessedVarList, MVLI.VarBaseDeclarations,
-                              MVLI.VarComponents, MapTypeModifier, MapType,
-                              IsMapTypeImplicit, MapLoc);
+                              MVLI.VarComponents, Modifiers, ModifiersLoc,
+                              MapType, IsMapTypeImplicit, MapLoc);
 }
 
 QualType Sema::ActOnOpenMPDeclareReductionType(SourceLocation TyLoc,
                                                TypeResult ParsedType) {
   assert(ParsedType.isUsable());
 
   QualType ReductionType = GetTypeFromParser(ParsedType.get());
   if (ReductionType.isNull())

lib/Sema/TreeTransform.h

return getSema().ActOnOpenMPDeviceClause(Device, StartLoc, LParenLoc,
                                              EndLoc);
   }
 
   /// 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,
+  RebuildOMPMapClause(ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
+                      ArrayRef<SourceLocation> MapTypeModifiersLoc,
                       OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
                       SourceLocation MapLoc, SourceLocation ColonLoc,
                       ArrayRef<Expr *> VarList, SourceLocation StartLoc,
                       SourceLocation LParenLoc, SourceLocation EndLoc) {
-    return getSema().ActOnOpenMPMapClause(MapTypeModifier, MapType,
-                                          IsMapTypeImplicit, MapLoc, ColonLoc,
-                                          VarList, StartLoc, LParenLoc, EndLoc);
+    return getSema().ActOnOpenMPMapClause(MapTypeModifiers, MapTypeModifiersLoc,
+                                          MapType, IsMapTypeImplicit, MapLoc,
+                                          ColonLoc, VarList, StartLoc,
+                                          LParenLoc, EndLoc);
   }
 
   /// Build a new OpenMP 'num_teams' clause.
   ///
   /// By default, performs semantic analysis to build the new statement.
   /// Subclasses may override this routine to provide different behavior.
   OMPClause *RebuildOMPNumTeamsClause(Expr *NumTeams, SourceLocation StartLoc,
                                       SourceLocation LParenLoc,
OMPClause *TreeTransform<Derived>::TransformOMPMapClause(OMPMapClause *C) {
   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->isImplicitMapType(),
-      C->getMapLoc(), C->getColonLoc(), Vars, C->getBeginLoc(),
-      C->getLParenLoc(), C->getEndLoc());
+      C->getMapTypeModifiers(), C->getMapTypeModifiersLoc(), C->getMapType(),
+      C->isImplicitMapType(), C->getMapLoc(), C->getColonLoc(), Vars,
+      C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc());
 }
 
 template <typename Derived>
 OMPClause *
 TreeTransform<Derived>::TransformOMPNumTeamsClause(OMPNumTeamsClause *C) {
   ExprResult E = getDerived().TransformExpr(C->getNumTeams());
   if (E.isInvalid())
     return nullptr;

lib/Serialization/ASTReader.cpp

 void OMPClauseReader::VisitOMPDeviceClause(OMPDeviceClause *C) {
   VisitOMPClauseWithPreInit(C);
   C->setDevice(Record.readSubExpr());
   C->setLParenLoc(Record.readSourceLocation());
 }
 
 void OMPClauseReader::VisitOMPMapClause(OMPMapClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
+  for (unsigned I = 0; I < OMPMapClause::NumberOfModifiers; ++I) {
     C->setMapTypeModifier(
-     static_cast<OpenMPMapClauseKind>(Record.readInt()));
+        I, static_cast<OpenMPMapModifierKind>(Record.readInt()));
+    C->setMapTypeModifierLoc(I, Record.readSourceLocation());
+  }
   C->setMapType(
      static_cast<OpenMPMapClauseKind>(Record.readInt()));
   C->setMapLoc(Record.readSourceLocation());
   C->setColonLoc(Record.readSourceLocation());
   auto NumVars = C->varlist_size();
   auto UniqueDecls = C->getUniqueDeclarationsNum();
   auto TotalLists = C->getTotalComponentListNum();
   auto TotalComponents = C->getTotalComponentsNum();

lib/Serialization/ASTWriter.cpp

 }
 
 void OMPClauseWriter::VisitOMPMapClause(OMPMapClause *C) {
   Record.push_back(C->varlist_size());
   Record.push_back(C->getUniqueDeclarationsNum());
   Record.push_back(C->getTotalComponentListNum());
   Record.push_back(C->getTotalComponentsNum());
   Record.AddSourceLocation(C->getLParenLoc());
-  Record.push_back(C->getMapTypeModifier());
+  for (unsigned I = 0; I < OMPMapClause::NumberOfModifiers; ++I) {
+    Record.push_back(C->getMapTypeModifier(I));
+    Record.AddSourceLocation(C->getMapTypeModifierLoc(I));
+  }
   Record.push_back(C->getMapType());
   Record.AddSourceLocation(C->getMapLoc());
   Record.AddSourceLocation(C->getColonLoc());
   for (auto *E : C->varlists())
     Record.AddStmt(E);
   for (auto *D : C->all_decls())
     Record.AddDeclRef(D);
   for (auto N : C->all_num_lists())

test/OpenMP/target_ast_print.cpp

 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -ast-print %s | FileCheck %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
 
 // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=45 -ast-print %s | FileCheck %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -x c++ -std=c++11 -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=45 -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], always;
+  T i, j, a[20], always, close;
 #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();
 #pragma omp target map(always from: i)
   foo();
 #pragma omp target map(always)
   {always++;}
 #pragma omp target map(always,i)
   {always++;i++;}
+#pragma omp target map(close,alloc: i)
+  foo();
+#pragma omp target map(close from: i)
+  foo();
+#pragma omp target map(close)
+  {close++;}
+#pragma omp target map(close,i)
+  {close++;i++;}
 #pragma omp target nowait
   foo();
 #pragma omp target depend(in : argc, argv[i:argc], a[:])
   foo();
 #pragma omp target defaultmap(tofrom: scalar)
   foo();
   return 0;
 }

 // CHECK-NEXT: {
 // CHECK-NEXT: always++;
 // CHECK-NEXT: }
 // CHECK-NEXT: #pragma omp target map(tofrom: always,i)
 // CHECK-NEXT: {
 // CHECK-NEXT: always++;
 // CHECK-NEXT: i++;
 // CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target map(close,alloc: i)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(close,from: i)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(tofrom: close)
+// CHECK-NEXT: {
+// CHECK-NEXT: close++;
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target map(tofrom: close,i)
+// CHECK-NEXT: {
+// CHECK-NEXT: close++;
+// CHECK-NEXT: i++;
+// CHECK-NEXT: }
 // CHECK-NEXT: #pragma omp target nowait
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target defaultmap(tofrom: scalar)
 // CHECK-NEXT: foo()
 // CHECK: template<> int tmain<int, 5>(int argc, int *argv) {
 // CHECK-NEXT: int i, j, a[20]

 // CHECK-NEXT: {
 // CHECK-NEXT: always++;
 // CHECK-NEXT: }
 // CHECK-NEXT: #pragma omp target map(tofrom: always,i)
 // CHECK-NEXT: {
 // CHECK-NEXT: always++;
 // CHECK-NEXT: i++;
 // CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target map(close,alloc: i)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(close,from: i)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(tofrom: close)
+// CHECK-NEXT: {
+// CHECK-NEXT: close++;
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target map(tofrom: close,i)
+// CHECK-NEXT: {
+// CHECK-NEXT: close++;
+// CHECK-NEXT: i++;
+// CHECK-NEXT: }
 // CHECK-NEXT: #pragma omp target nowait
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target defaultmap(tofrom: scalar)
 // CHECK-NEXT: foo()
 // CHECK: template<> char tmain<char, 1>(char argc, char *argv) {
 // CHECK-NEXT: char i, j, a[20]

 // CHECK-NEXT: {
 // CHECK-NEXT: always++;
 // CHECK-NEXT: }
 // CHECK-NEXT: #pragma omp target map(tofrom: always,i)
 // CHECK-NEXT: {
 // CHECK-NEXT: always++;
 // CHECK-NEXT: i++;
 // CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target map(close,alloc: i)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(close,from: i)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target map(tofrom: close)
+// CHECK-NEXT: {
+// CHECK-NEXT: close++;
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target map(tofrom: close,i)
+// CHECK-NEXT: {
+// CHECK-NEXT: close++;
+// CHECK-NEXT: i++;
+// CHECK-NEXT: }
 // CHECK-NEXT: #pragma omp target nowait
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
 // CHECK-NEXT: foo()
 // CHECK-NEXT: #pragma omp target defaultmap(tofrom: scalar)
 // CHECK-NEXT: foo()
 
 // CHECK-LABEL: int main(int argc, char **argv) {
 int main (int argc, char **argv) {
-  int i, j, a[20], always;
+  int i, j, a[20], always, close;
 // 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();

 #pragma omp target map(always,i)
 // CHECK-NEXT: #pragma omp target map(tofrom: always,i)
   {always++;i++;}
 // CHECK-NEXT: {
 // CHECK-NEXT: always++;
 // CHECK-NEXT: i++;
 // CHECK-NEXT: }
 
+#pragma omp target map(close,alloc: i)
+// CHECK-NEXT: #pragma omp target map(close,alloc: i)
+  foo();
+// CHECK-NEXT: foo();
+
+#pragma omp target map(close from: i)
+// CHECK-NEXT: #pragma omp target map(close,from: i)
+  foo();
+// CHECK-NEXT: foo();
+
+#pragma omp target map(close)
+// CHECK-NEXT: #pragma omp target map(tofrom: close)
+  {close++;}
+// CHECK-NEXT: {
+// CHECK-NEXT: close++;
+// CHECK-NEXT: }
+
+#pragma omp target map(close,i)
+// CHECK-NEXT: #pragma omp target map(tofrom: close,i)
+  {close++;i++;}
+// CHECK-NEXT: {
+// CHECK-NEXT: close++;
+// CHECK-NEXT: i++;
+// CHECK-NEXT: }
+
 #pragma omp target nowait
 // CHECK-NEXT: #pragma omp target nowait
   foo();
 // CHECK-NEXT: foo();
 
 #pragma omp target depend(in : argc, argv[i:argc], a[:])
 // CHECK-NEXT: #pragma omp target depend(in : argc,argv[i:argc],a[:])
   foo();

test/OpenMP/target_data_ast_print.cpp

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

test/OpenMP/target_map_messages.cpp

void func(int arg) {
     {}
     #pragma omp target map(to:b,e[:])
     {}
     #pragma omp target map(b[-1:]) // expected-error {{array section must be a subset of the original array}}
     {}
     #pragma omp target map(b[:-1]) // expected-error {{section length is evaluated to a negative value -1}}
     {}
 
+    #pragma omp target map(: c,f) // expected-error {{missing map type}}
+    {}
     #pragma omp target map(always, tofrom: c,f)
     {}
     #pragma omp target map(always, tofrom: c[1:2],f)
     {}
     #pragma omp target map(always, tofrom: c,f[1:2])
     {}
     #pragma omp target map(always, tofrom: c[:],f)   // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
     {}
     #pragma omp target map(always, tofrom: c,f[:])   // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
     {}
     #pragma omp target map(always)   // expected-error {{use of undeclared identifier 'always'}}
     {}
+    #pragma omp target map(close, tofrom: c,f)
+    {}
+    #pragma omp target map(close, tofrom: c[1:2],f)
+    {}
+    #pragma omp target map(close, tofrom: c,f[1:2])
+    {}
+    #pragma omp target map(close, tofrom: c[:],f)   // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+    {}
+    #pragma omp target map(close, tofrom: c,f[:])   // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+    {}
+    #pragma omp target map(close)   // expected-error {{use of undeclared identifier 'close'}}
+    {}
+    #pragma omp target map(close, close, tofrom: a)   // expected-error {{same map type modifier has been specified more than once}}
+    {}
+    #pragma omp target map(always, close, always, close, tofrom: a)   // expected-error {{same map type modifier has been specified more than once}} expected-error {{same map type modifier has been specified more than once}}
+    {}
+    #pragma omp target map( , tofrom: a)   // expected-error {{missing map type modifier}}
+    {}
+    #pragma omp target map( , , tofrom: a)   // expected-error {{missing map type modifier}} expected-error {{missing map type modifier}}
+    {}
+    #pragma omp target map( , , : a)   // expected-error {{missing map type modifier}} expected-error {{missing map type modifier}} expected-error {{missing map type}}
+    {}
+    #pragma omp target map( d, f, bf: a)   // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+    {}
+    #pragma omp target map( , f, : a)   // expected-error {{missing map type modifier}} expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+    {}
+    #pragma omp target map(always close: a)   // expected-error {{missing map type}}
+    {}
+    #pragma omp target map(always close bf: a)   // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+    {}
+    #pragma omp target map(always tofrom close: a)   // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+    {}
+    #pragma omp target map(tofrom from: a)   // expected-error {{incorrect map type modifier, expected 'always' or 'close'}}
+    {}
+    #pragma omp target map(close bf: a)   // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+    {}
     return;
   }
 };
 
 struct SB {
   unsigned A;
   unsigned B;
   float Arr[100];
T tmain(T argc) {
   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;
+  T to, tofrom, always, close;
   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 data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
 #pragma omp target data map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
 #pragma omp target data map(j)
 #pragma omp target map(l)
   foo();
 
 #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(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
 #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 data map(close, tofrom: x)
+#pragma omp target data map(close: x) // expected-error {{missing map type}}
+#pragma omp target data map(tofrom, close: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+#pragma omp target data map(close, tofrom: close, tofrom, x)
+  foo();
   return 0;
 }
 
 struct SA1{
   int a;
   struct SA1 *p;
   int b[10];
 };
int main(int argc, char **argv) {
   S4 e(4);
   S5 g(5);
   int i;
   int &j = i;
   int *k = &j;
   S6<int> m;
   int x;
   int y;
-  int to, tofrom, always;
+  int to, tofrom, always, close;
   const int (&l)[5] = da;
   SC1 s;
   SC1 *p;
 #pragma omp target data map // expected-error {{expected '(' after 'map'}} expected-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}}
 #pragma omp target data map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
 #pragma omp target data map() // expected-error {{expected expression}}
 #pragma omp target data map(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
 #pragma omp target data map(to argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected ',' or ')' in 'map' clause}}

 #pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
 #pragma omp target data map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
 #pragma omp target data map(j)
 #pragma omp target map(l)
   foo();
 
 #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(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
 #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 data map(close, tofrom: x)
+#pragma omp target data map(close: x) // expected-error {{missing map type}}
+#pragma omp target data map(tofrom, close: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
+  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(m)
   {}
 // expected-note@+1 {{used here}}
 #pragma omp target map(s.s.s)

test/OpenMP/target_parallel_for_map_messages.cpp

 #pragma omp target parallel for map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target parallel for map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for 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 parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }

 #pragma omp target parallel for map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target parallel for map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for 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 parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) 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}}
 }
 

test/OpenMP/target_parallel_for_simd_map_messages.cpp

 #pragma omp target parallel for simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target parallel for simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd 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 parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }

 #pragma omp target parallel for simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target parallel for simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd 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 parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) 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}}
 }
 

test/OpenMP/target_parallel_map_messages.cpp

 #pragma omp target parallel map(l)
   foo();
 }
 
 #pragma omp target parallel map(always, tofrom: x)
   foo();
 #pragma omp target parallel map(always: x) // expected-error {{missing map type}}
   foo();
-#pragma omp target parallel 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 parallel map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
   foo();
 #pragma omp target parallel map(always, tofrom: always, tofrom, x)
   foo();
 #pragma omp target parallel map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
 
   return 0;
 }

 #pragma omp target parallel map(l)
   foo();
 }
 
 #pragma omp target parallel map(always, tofrom: x)
   foo();
 #pragma omp target parallel map(always: x) // expected-error {{missing map type}}
   foo();
-#pragma omp target parallel 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 parallel map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
   foo();
 #pragma omp target parallel map(always, tofrom: always, tofrom, x)
   foo();
 #pragma omp target parallel 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}}
 }
 

test/OpenMP/target_simd_map_messages.cpp

 #pragma omp target simd map(l) // OK
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd 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 simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }

 #pragma omp target simd map(l) //
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd 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 simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) 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}}
 }
 

test/OpenMP/target_teams_distribute_map_messages.cpp

 #pragma omp target teams distribute map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute 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 teams distribute map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }

 #pragma omp target teams distribute map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute 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 teams distribute map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) 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}}
 }
 

test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp

 #pragma omp target teams distribute parallel for map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute parallel for map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for 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 teams distribute parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }

 #pragma omp target teams distribute parallel for map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute parallel for map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for 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 teams distribute parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) 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}}
 }
 

test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp

 #pragma omp target teams distribute parallel for simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute parallel for simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd 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 teams distribute parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }

 #pragma omp target teams distribute parallel for simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute parallel for simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd 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 teams distribute parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) 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}}
 }
 

test/OpenMP/target_teams_distribute_simd_map_messages.cpp

 #pragma omp target teams distribute simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd 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 teams distribute simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }

 #pragma omp target teams distribute simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd 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 teams distribute simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) 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}}
 }
 

test/OpenMP/target_teams_map_messages.cpp

 #pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
 #pragma omp target data map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
 #pragma omp target data map(j)
 #pragma omp target teams map(l)
   foo();
 
 #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(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
 #pragma omp target data map(always, tofrom: always, tofrom, x)
 #pragma omp target teams map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
   return 0;
 }
 
 int main(int argc, char **argv) {
   const int d = 5;

 #pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
 #pragma omp target data map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
 #pragma omp target data map(j)
 #pragma omp target teams map(l)
   foo();
 
 #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(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always' or 'close'}} expected-error {{missing map type}}
 #pragma omp target data map(always, tofrom: always, tofrom, x)
 #pragma omp target teams map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
 
 #pragma omp target teams private(j) map(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target teams' directive}}  expected-note {{defined as private}}
   {}
 
 #pragma omp target teams firstprivate(j) map(j)  // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams' directive}} expected-note {{defined as firstprivate}}
   {}
 
 #pragma omp target teams map(m)
   {}
 
   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