[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
diff --git a/include/clang/AST/OpenMPClause.h b/include/clang/AST/OpenMPClause.h
index e93f7e2..bdcdf74 100644
--- a/include/clang/AST/OpenMPClause.h
+++ b/include/clang/AST/OpenMPClause.h
@@ -4061,8 +4061,19 @@
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 @@
/// 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 @@
/// 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 @@
: 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 @@
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 @@
/// \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 @@
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 @@
/// 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.
diff --git a/include/clang/Basic/DiagnosticParseKinds.td b/include/clang/Basic/DiagnosticParseKinds.td
index ec07bb5..1e816bc 100644
--- a/include/clang/Basic/DiagnosticParseKinds.td
+++ b/include/clang/Basic/DiagnosticParseKinds.td
@@ -1157,9 +1157,11 @@
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<
diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td
index 2e84f5f..9ab263f 100644
--- a/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8725,6 +8725,8 @@
"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<
diff --git a/include/clang/Basic/OpenMPKinds.def b/include/clang/Basic/OpenMPKinds.def
index eb98811..f86721b 100644
--- a/include/clang/Basic/OpenMPKinds.def
+++ b/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(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 @@
#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
diff --git a/include/clang/Basic/OpenMPKinds.h b/include/clang/Basic/OpenMPKinds.h
index b0fab5a..3e03a48 100644
--- a/include/clang/Basic/OpenMPKinds.h
+++ b/include/clang/Basic/OpenMPKinds.h
@@ -96,6 +96,15 @@
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,
diff --git a/include/clang/Parse/Parser.h b/include/clang/Parse/Parser.h
index 09e81d2..46e4431 100644
--- a/include/clang/Parse/Parser.h
+++ b/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 @@
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;
diff --git a/include/clang/Sema/Sema.h b/include/clang/Sema/Sema.h
index 4ddbf6c..7685b3f 100644
--- a/include/clang/Sema/Sema.h
+++ b/include/clang/Sema/Sema.h
@@ -9233,7 +9233,9 @@
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 @@
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,
diff --git a/lib/AST/OpenMPClause.cpp b/lib/AST/OpenMPClause.cpp
index 3124b0f..76098f1 100644
--- a/lib/AST/OpenMPClause.cpp
+++ b/lib/AST/OpenMPClause.cpp
@@ -796,8 +796,10 @@
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 @@
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 @@
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 << ':';
diff --git a/lib/Basic/OpenMPKinds.cpp b/lib/Basic/OpenMPKinds.cpp
index dcd36c4..a5bfac8 100644
--- a/lib/Basic/OpenMPKinds.cpp
+++ b/lib/Basic/OpenMPKinds.cpp
@@ -108,8 +108,11 @@
#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 @@
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;
diff --git a/lib/CodeGen/CGOpenMPRuntime.cpp b/lib/CodeGen/CGOpenMPRuntime.cpp
index 66f0783..d7a0a72 100644
--- a/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -6645,17 +6645,17 @@
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 @@
/// 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 @@
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 @@
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 @@
/// \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 @@
// 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 @@
// 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 @@
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 @@
// 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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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() &&
diff --git a/lib/Parse/ParseOpenMP.cpp b/lib/Parse/ParseOpenMP.cpp
index 1fb5280..17c3fa3 100644
--- a/lib/Parse/ParseOpenMP.cpp
+++ b/lib/Parse/ParseOpenMP.cpp
@@ -1774,6 +1774,79 @@
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 @@
// 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 @@
/// 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 @@
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);
}
diff --git a/lib/Sema/SemaOpenMP.cpp b/lib/Sema/SemaOpenMP.cpp
index 2daf454..38a329a 100644
--- a/lib/Sema/SemaOpenMP.cpp
+++ b/lib/Sema/SemaOpenMP.cpp
@@ -3336,9 +3336,10 @@
}
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 @@
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 @@
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 @@
}
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 @@
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,
diff --git a/lib/Sema/TreeTransform.h b/lib/Sema/TreeTransform.h
index 3f4b21e..50ecd6d 100644
--- a/lib/Sema/TreeTransform.h
+++ b/lib/Sema/TreeTransform.h
@@ -1796,14 +1796,16 @@
/// 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 @@
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>
diff --git a/lib/Serialization/ASTReader.cpp b/lib/Serialization/ASTReader.cpp
index 558100d..e0b2b24 100644
--- a/lib/Serialization/ASTReader.cpp
+++ b/lib/Serialization/ASTReader.cpp
@@ -12284,8 +12284,11 @@
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());
diff --git a/lib/Serialization/ASTWriter.cpp b/lib/Serialization/ASTWriter.cpp
index 45de91f..37adcb7 100644
--- a/lib/Serialization/ASTWriter.cpp
+++ b/lib/Serialization/ASTWriter.cpp
@@ -6782,7 +6782,10 @@
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());
diff --git a/test/OpenMP/target_ast_print.cpp b/test/OpenMP/target_ast_print.cpp
index 827d6a4..2734294 100644
--- a/test/OpenMP/target_ast_print.cpp
+++ b/test/OpenMP/target_ast_print.cpp
@@ -14,7 +14,7 @@
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 @@
{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 @@
// 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 @@
// 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 @@
// 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 @@
// 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 @@
// 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();
diff --git a/test/OpenMP/target_data_ast_print.cpp b/test/OpenMP/target_data_ast_print.cpp
index 17dace8..fa67c18 100644
--- a/test/OpenMP/target_data_ast_print.cpp
+++ b/test/OpenMP/target_data_ast_print.cpp
@@ -40,11 +40,16 @@
#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 @@
// 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 @@
// 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 @@
// 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 @@
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 @@
// 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]);
}
diff --git a/test/OpenMP/target_map_messages.cpp b/test/OpenMP/target_map_messages.cpp
index f9547ba..e81e61e 100644
--- a/test/OpenMP/target_map_messages.cpp
+++ b/test/OpenMP/target_map_messages.cpp
@@ -74,6 +74,8 @@
#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 @@
{}
#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 *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 @@
#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 @@
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 @@
#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}}
diff --git a/test/OpenMP/target_parallel_for_map_messages.cpp b/test/OpenMP/target_parallel_for_map_messages.cpp
index f4f98df..6d82921 100644
--- a/test/OpenMP/target_parallel_for_map_messages.cpp
+++ b/test/OpenMP/target_parallel_for_map_messages.cpp
@@ -163,7 +163,7 @@
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 @@
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();
diff --git a/test/OpenMP/target_parallel_for_simd_map_messages.cpp b/test/OpenMP/target_parallel_for_simd_map_messages.cpp
index 1cab598..a513555 100644
--- a/test/OpenMP/target_parallel_for_simd_map_messages.cpp
+++ b/test/OpenMP/target_parallel_for_simd_map_messages.cpp
@@ -163,7 +163,7 @@
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 @@
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();
diff --git a/test/OpenMP/target_parallel_map_messages.cpp b/test/OpenMP/target_parallel_map_messages.cpp
index 4bb14ea..056fd50 100644
--- a/test/OpenMP/target_parallel_map_messages.cpp
+++ b/test/OpenMP/target_parallel_map_messages.cpp
@@ -163,7 +163,7 @@
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 @@
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();
diff --git a/test/OpenMP/target_simd_map_messages.cpp b/test/OpenMP/target_simd_map_messages.cpp
index 3722ecb..acd6298 100644
--- a/test/OpenMP/target_simd_map_messages.cpp
+++ b/test/OpenMP/target_simd_map_messages.cpp
@@ -159,7 +159,7 @@
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 @@
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();
diff --git a/test/OpenMP/target_teams_distribute_map_messages.cpp b/test/OpenMP/target_teams_distribute_map_messages.cpp
index 826a09d..bbfa7cd 100644
--- a/test/OpenMP/target_teams_distribute_map_messages.cpp
+++ b/test/OpenMP/target_teams_distribute_map_messages.cpp
@@ -163,7 +163,7 @@
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 @@
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();
diff --git a/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp b/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp
index 5cba1e7..f585d0a 100644
--- a/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp
+++ b/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp
@@ -163,7 +163,7 @@
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 @@
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();
diff --git a/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp b/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp
index 63d85eb..5cf0327 100644
--- a/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp
+++ b/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp
@@ -163,7 +163,7 @@
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 @@
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();
diff --git a/test/OpenMP/target_teams_distribute_simd_map_messages.cpp b/test/OpenMP/target_teams_distribute_simd_map_messages.cpp
index 73671d7..99c633a 100644
--- a/test/OpenMP/target_teams_distribute_simd_map_messages.cpp
+++ b/test/OpenMP/target_teams_distribute_simd_map_messages.cpp
@@ -163,7 +163,7 @@
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 @@
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();
diff --git a/test/OpenMP/target_teams_map_messages.cpp b/test/OpenMP/target_teams_map_messages.cpp
index 57d9396..4b2629e 100644
--- a/test/OpenMP/target_teams_map_messages.cpp
+++ b/test/OpenMP/target_teams_map_messages.cpp
@@ -454,7 +454,7 @@
#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 @@
#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();
