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

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.

Patch by Ahsan Saghir (saghir)

Differential Revision: https://reviews.llvm.org/D55719


git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@349551 91177308-0d34-0410-b5e6-96231b3b80d8

include/clang/AST/OpenMPClause.h

@@ -4061,8 +4061,19 @@ class OMPMapClause final : public OMPMappableExprListClause<OMPMapClause>,
     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;
@@ -4080,7 +4091,8 @@ class OMPMapClause final : public OMPMappableExprListClause<OMPMapClause>,
   /// 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.
@@ -4091,7 +4103,8 @@ class OMPMapClause final : public OMPMappableExprListClause<OMPMapClause>,
   /// 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,
@@ -4100,8 +4113,17 @@ class OMPMapClause final : public OMPMappableExprListClause<OMPMapClause>,
       : 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.
   ///
@@ -4116,10 +4138,25 @@ class OMPMapClause final : public OMPMappableExprListClause<OMPMapClause>,
             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.
   ///
@@ -4143,7 +4180,8 @@ public:
   /// \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.
@@ -4152,7 +4190,8 @@ public:
                               ArrayRef<Expr *> Vars,
                               ArrayRef<ValueDecl *> Declarations,
                               MappableExprComponentListsRef ComponentLists,
-                              OpenMPMapClauseKind TypeModifier,
+                              ArrayRef<OpenMPMapModifierKind> MapModifiers,
+                              ArrayRef<SourceLocation> MapModifiersLoc,
                               OpenMPMapClauseKind Type, bool TypeIsImplicit,
                               SourceLocation TypeLoc);
 
@@ -4182,9 +4221,33 @@ public:
   /// 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.

include/clang/Basic/DiagnosticParseKinds.td

@@ -1157,9 +1157,11 @@ def err_omp_decl_in_declare_simd : Error<
 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<

include/clang/Basic/DiagnosticSemaKinds.td

@@ -8725,6 +8725,8 @@ 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<

include/clang/Basic/OpenMPKinds.def

@@ -120,6 +120,9 @@
 #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
@@ -559,14 +562,17 @@ 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)
@@ -919,6 +925,7 @@ OPENMP_TASKGROUP_CLAUSE(task_reduction)
 #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

include/clang/Basic/OpenMPKinds.h

@@ -96,6 +96,15 @@ enum OpenMPMapClauseKind {
   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/Parse/Parser.h

@@ -14,6 +14,7 @@
 #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"
@@ -2876,7 +2877,10 @@ public:
     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;

include/clang/Sema/Sema.h

@@ -9233,7 +9233,9 @@ public:
       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.
@@ -9317,7 +9319,8 @@ public:
                                      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,

lib/AST/OpenMPClause.cpp

@@ -796,8 +796,10 @@ 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);
@@ -820,12 +822,12 @@ OMPMapClause::Create(const ASTContext &C, SourceLocation StartLoc,
           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;
@@ -1426,10 +1428,12 @@ void OMPClausePrinter::VisitOMPMapClause(OMPMapClause *Node) {
   if (!Node->varlist_empty()) {
     OS << "map(";
     if (Node->getMapType() != OMPC_MAP_unknown) {
-      if (Node->getMapTypeModifier() != OMPC_MAP_unknown) {
-        OS << getOpenMPSimpleClauseTypeName(OMPC_map,
-                                            Node->getMapTypeModifier());
-        OS << ',';
+      for (unsigned I = 0; I < OMPMapClause::NumberOfModifiers; ++I) {
+        if (Node->getMapTypeModifier(I) != OMPC_MAP_MODIFIER_unknown) {
+          OS << getOpenMPSimpleClauseTypeName(OMPC_map,
+                                              Node->getMapTypeModifier(I));
+          OS << ',';
+        }
       }
       OS << getOpenMPSimpleClauseTypeName(OMPC_map, Node->getMapType());
       OS << ':';

lib/Basic/OpenMPKinds.cpp

@@ -108,8 +108,11 @@ unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind,
 #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:
@@ -243,10 +246,14 @@ const char *clang::getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind,
   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;

lib/CodeGen/CGOpenMPRuntime.cpp

@@ -6645,17 +6645,17 @@ private:
   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) {}
   };
 
@@ -6732,10 +6732,9 @@ private:
   /// 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) {
@@ -6758,7 +6757,6 @@ private:
     case OMPC_MAP_delete:
       Bits |= OMP_MAP_DELETE;
       break;
-    case OMPC_MAP_always:
     case OMPC_MAP_unknown:
       llvm_unreachable("Unexpected map type!");
     }
@@ -6766,7 +6764,8 @@ private:
       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;
   }
@@ -6815,7 +6814,8 @@ private:
   /// \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,
@@ -7125,7 +7125,7 @@ private:
           // 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;
@@ -7175,7 +7175,7 @@ private:
           // 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) {
@@ -7395,28 +7395,29 @@ public:
     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());
       }
 
@@ -7469,7 +7470,7 @@ public:
           // 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 {
@@ -7503,7 +7504,7 @@ public:
         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);
 
@@ -7662,7 +7663,7 @@ public:
 
     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>()) {
@@ -7672,7 +7673,7 @@ public:
         assert(!L.second.empty() &&
                "Not expecting declaration with no component lists.");
         DeclComponentLists.emplace_back(L.second, C->getMapType(),
-                                        C->getMapTypeModifier(),
+                                        C->getMapTypeModifiers(),
                                         C->isImplicit());
       }
     }
@@ -7688,13 +7689,13 @@ public:
     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();
@@ -7778,13 +7779,13 @@ public:
       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);
@@ -7794,12 +7795,12 @@ public:
     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);
@@ -7828,7 +7829,7 @@ public:
           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() &&

lib/Parse/ParseOpenMP.cpp

@@ -1774,6 +1774,79 @@ static bool ParseReductionId(Parser &P, CXXScopeSpec &ReductionIdScopeSpec,
                               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,
@@ -1852,82 +1925,27 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
     // 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;
     }
@@ -2025,7 +2043,7 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
 ///    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 ')'
@@ -2056,7 +2074,7 @@ OMPClause *Parser::ParseOpenMPVarListClause(OpenMPDirectiveKind DKind,
   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

@@ -3336,9 +3336,10 @@ StmtResult Sema::ActOnOpenMPExecutableDirective(
     }
     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();
@@ -9535,7 +9536,9 @@ OMPClause *Sema::ActOnOpenMPVarListClause(
     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;
@@ -9588,9 +9591,9 @@ OMPClause *Sema::ActOnOpenMPVarListClause(
                                   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);
@@ -12957,7 +12960,8 @@ checkMappableExpressionList(Sema &SemaRef, DSAStackTy *DSAS,
 }
 
 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,
@@ -12966,12 +12970,31 @@ Sema::ActOnOpenMPMapClause(OpenMPMapClauseKind MapTypeModifier,
   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,

lib/Sema/TreeTransform.h

@@ -1796,14 +1796,16 @@ public:
   /// 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.
@@ -8803,9 +8805,9 @@ OMPClause *TreeTransform<Derived>::TransformOMPMapClause(OMPMapClause *C) {
     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>

lib/Serialization/ASTReader.cpp

@@ -12284,8 +12284,11 @@ void OMPClauseReader::VisitOMPDeviceClause(OMPDeviceClause *C) {
 
 void OMPClauseReader::VisitOMPMapClause(OMPMapClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
-  C->setMapTypeModifier(
-     static_cast<OpenMPMapClauseKind>(Record.readInt()));
+  for (unsigned I = 0; I < OMPMapClause::NumberOfModifiers; ++I) {
+    C->setMapTypeModifier(
+        I, static_cast<OpenMPMapModifierKind>(Record.readInt()));
+    C->setMapTypeModifierLoc(I, Record.readSourceLocation());
+  }
   C->setMapType(
      static_cast<OpenMPMapClauseKind>(Record.readInt()));
   C->setMapLoc(Record.readSourceLocation());

lib/Serialization/ASTWriter.cpp

@@ -6782,7 +6782,10 @@ void OMPClauseWriter::VisitOMPMapClause(OMPMapClause *C) {
   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());

test/OpenMP/target_ast_print.cpp

@@ -14,7 +14,7 @@ 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)
@@ -35,6 +35,14 @@ T tmain(T argc, T *argv) {
   {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[:])
@@ -71,6 +79,19 @@ T tmain(T argc, T *argv) {
 // 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[:])
@@ -104,6 +125,19 @@ T tmain(T argc, T *argv) {
 // 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[:])
@@ -137,6 +171,19 @@ T tmain(T argc, T *argv) {
 // 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[:])
@@ -146,7 +193,7 @@ T tmain(T argc, T *argv) {
 
 // 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
@@ -202,6 +249,31 @@ int main (int argc, char **argv) {
 // 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();

test/OpenMP/target_data_ast_print.cpp

@@ -40,11 +40,16 @@ T tmain(T argc, T *argv) {
 #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;
@@ -68,10 +73,14 @@ T tmain(T argc, T *argv) {
 // 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)
@@ -90,10 +99,14 @@ T tmain(T argc, T *argv) {
 // 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)
@@ -112,10 +125,14 @@ T tmain(T argc, T *argv) {
 // 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];
@@ -161,6 +178,11 @@ int main (int argc, char **argv) {
   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)
@@ -170,7 +192,11 @@ int main (int argc, char **argv) {
 // 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]);
 }
 

test/OpenMP/target_map_messages.cpp

@@ -74,6 +74,8 @@ struct SA {
     #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)
@@ -86,6 +88,42 @@ struct SA {
     {}
     #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;
   }
 };
@@ -405,7 +443,7 @@ T tmain(T argc) {
   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'}}
   {}
@@ -478,10 +516,16 @@ T tmain(T argc) {
 
 #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;
 }
 
@@ -515,7 +559,7 @@ int main(int argc, char **argv) {
   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;
@@ -569,10 +613,14 @@ int main(int argc, char **argv) {
 
 #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}}

test/OpenMP/target_parallel_for_map_messages.cpp

@@ -163,7 +163,7 @@ T tmain(T argc) {
   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();
@@ -271,7 +271,7 @@ int main(int argc, char **argv) {
   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();

test/OpenMP/target_parallel_for_simd_map_messages.cpp

@@ -163,7 +163,7 @@ T tmain(T argc) {
   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();
@@ -271,7 +271,7 @@ int main(int argc, char **argv) {
   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();

test/OpenMP/target_parallel_map_messages.cpp

@@ -163,7 +163,7 @@ T tmain(T argc) {
   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();
@@ -270,7 +270,7 @@ int main(int argc, char **argv) {
   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();

test/OpenMP/target_simd_map_messages.cpp

@@ -159,7 +159,7 @@ T tmain(T argc) {
   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();
@@ -263,7 +263,7 @@ int main(int argc, char **argv) {
   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();

test/OpenMP/target_teams_distribute_map_messages.cpp

@@ -163,7 +163,7 @@ T tmain(T argc) {
   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();
@@ -271,7 +271,7 @@ int main(int argc, char **argv) {
   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();

test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp

@@ -163,7 +163,7 @@ T tmain(T argc) {
   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();
@@ -271,7 +271,7 @@ int main(int argc, char **argv) {
   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();

test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp

@@ -163,7 +163,7 @@ T tmain(T argc) {
   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();
@@ -271,7 +271,7 @@ int main(int argc, char **argv) {
   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();

test/OpenMP/target_teams_distribute_simd_map_messages.cpp

@@ -163,7 +163,7 @@ T tmain(T argc) {
   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();
@@ -271,7 +271,7 @@ int main(int argc, char **argv) {
   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();

test/OpenMP/target_teams_map_messages.cpp

@@ -454,7 +454,7 @@ T tmain(T argc) {
 
 #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();
@@ -529,7 +529,7 @@ int main(int argc, char **argv) {
 
 #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();