Index: clang/include/clang/AST/OpenMPClause.h =================================================================== --- clang/include/clang/AST/OpenMPClause.h +++ clang/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 MapModifiers, + ArrayRef 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 Vars, ArrayRef Declarations, MappableExprComponentListsRef ComponentLists, - OpenMPMapClauseKind TypeModifier, + ArrayRef MapModifiers, + ArrayRef 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 getMapTypeModifiers() const LLVM_READONLY { + return llvm::makeArrayRef(MapTypeModifiers); + } + + /// Fetches ArrayRef of location of map-type-modifiers. + ArrayRef getMapTypeModifiersLoc() const LLVM_READONLY { + return llvm::makeArrayRef(MapTypeModifiersLoc); } /// Fetches location of clause mapping kind. Index: clang/include/clang/Basic/DiagnosticParseKinds.td =================================================================== --- clang/include/clang/Basic/DiagnosticParseKinds.td +++ clang/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< Index: clang/include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- clang/include/clang/Basic/DiagnosticSemaKinds.td +++ clang/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< Index: clang/include/clang/Basic/OpenMPKinds.h =================================================================== --- clang/include/clang/Basic/OpenMPKinds.h +++ clang/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, Index: clang/include/clang/Basic/OpenMPKinds.def =================================================================== --- clang/include/clang/Basic/OpenMPKinds.def +++ clang/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 Index: clang/include/clang/Parse/Parser.h =================================================================== --- clang/include/clang/Parse/Parser.h +++ clang/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 + MapTypeModifiers; + SmallVector + MapTypeModifiersLoc; OpenMPMapClauseKind MapType = OMPC_MAP_unknown; bool IsMapTypeImplicit = false; SourceLocation DepLinMapLoc; Index: clang/include/clang/Sema/Sema.h =================================================================== --- clang/include/clang/Sema/Sema.h +++ clang/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 MapTypeModifiers, + ArrayRef 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 MapTypeModifiers, + ArrayRef MapTypeModifiersLoc, OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef VarList, SourceLocation StartLoc, Index: clang/lib/AST/OpenMPClause.cpp =================================================================== --- clang/lib/AST/OpenMPClause.cpp +++ clang/lib/AST/OpenMPClause.cpp @@ -796,8 +796,10 @@ SourceLocation LParenLoc, SourceLocation EndLoc, ArrayRef Vars, ArrayRef Declarations, MappableExprComponentListsRef ComponentLists, - OpenMPMapClauseKind TypeModifier, OpenMPMapClauseKind Type, - bool TypeIsImplicit, SourceLocation TypeLoc) { + ArrayRef MapModifiers, + ArrayRef 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 << ':'; Index: clang/lib/Basic/OpenMPKinds.cpp =================================================================== --- clang/lib/Basic/OpenMPKinds.cpp +++ clang/lib/Basic/OpenMPKinds.cpp @@ -108,8 +108,11 @@ #include "clang/Basic/OpenMPKinds.def" .Default(OMPC_LINEAR_unknown); case OMPC_map: - return llvm::StringSwitch(Str) -#define OPENMP_MAP_KIND(Name) .Case(#Name, OMPC_MAP_##Name) + return llvm::StringSwitch(Str) +#define OPENMP_MAP_KIND(Name) \ + .Case(#Name, static_cast(OMPC_MAP_##Name)) +#define OPENMP_MAP_MODIFIER_KIND(Name) \ + .Case(#Name, static_cast(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; Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp =================================================================== --- clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -6645,17 +6645,17 @@ struct MapInfo { OMPClauseMappableExprCommon::MappableExprComponentListRef Components; OpenMPMapClauseKind MapType = OMPC_MAP_unknown; - OpenMPMapClauseKind MapTypeModifier = OMPC_MAP_unknown; + ArrayRef MapModifiers; bool ReturnDevicePointer = false; bool IsImplicit = false; MapInfo() = default; MapInfo( OMPClauseMappableExprCommon::MappableExprComponentListRef Components, - OpenMPMapClauseKind MapType, OpenMPMapClauseKind MapTypeModifier, + OpenMPMapClauseKind MapType, + ArrayRef 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 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 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 MapModifiers, bool ReturnDevicePointer, bool IsImplicit) { const ValueDecl *VD = D ? cast(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()) 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()) 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()) 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; + OpenMPMapClauseKind, ArrayRef, bool>; SmallVector DeclComponentLists; // FIXME: MSVC 2013 seems to require this-> to find member CurDir. for (const auto *C : this->CurDir.getClausesOfKind()) { @@ -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 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 MapModifiers; bool IsImplicit; - std::tie(Components, MapType, MapTypeModifier, IsImplicit) = L; + std::tie(Components, MapType, MapModifiers, IsImplicit) = L; ArrayRef 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 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() && Index: clang/lib/Parse/ParseOpenMP.cpp =================================================================== --- clang/lib/Parse/ParseOpenMP.cpp +++ clang/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( + 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( + 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( - 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( - 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( - 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); } Index: clang/lib/Sema/SemaOpenMP.cpp =================================================================== --- clang/lib/Sema/SemaOpenMP.cpp +++ clang/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(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 MapTypeModifiers, + ArrayRef 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 MapTypeModifiers, + ArrayRef MapTypeModifiersLoc, OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef 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, Index: clang/lib/Sema/TreeTransform.h =================================================================== --- clang/lib/Sema/TreeTransform.h +++ clang/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 MapTypeModifiers, + ArrayRef MapTypeModifiersLoc, OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef 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 Index: clang/lib/Serialization/ASTReader.cpp =================================================================== --- clang/lib/Serialization/ASTReader.cpp +++ clang/lib/Serialization/ASTReader.cpp @@ -12284,8 +12284,11 @@ void OMPClauseReader::VisitOMPMapClause(OMPMapClause *C) { C->setLParenLoc(Record.readSourceLocation()); - C->setMapTypeModifier( - static_cast(Record.readInt())); + for (unsigned I = 0; I < OMPMapClause::NumberOfModifiers; ++I) { + C->setMapTypeModifier( + I, static_cast(Record.readInt())); + C->setMapTypeModifierLoc(I, Record.readSourceLocation()); + } C->setMapType( static_cast(Record.readInt())); C->setMapLoc(Record.readSourceLocation()); Index: clang/lib/Serialization/ASTWriter.cpp =================================================================== --- clang/lib/Serialization/ASTWriter.cpp +++ clang/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()); Index: clang/test/OpenMP/target_ast_print.cpp =================================================================== --- clang/test/OpenMP/target_ast_print.cpp +++ clang/test/OpenMP/target_ast_print.cpp @@ -14,7 +14,7 @@ template 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(); Index: clang/test/OpenMP/target_data_ast_print.cpp =================================================================== --- clang/test/OpenMP/target_data_ast_print.cpp +++ clang/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 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 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(argc, &argc) + tmain(argv[0][0], argv[0]); } Index: clang/test/OpenMP/target_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_map_messages.cpp +++ clang/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 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}} Index: clang/test/OpenMP/target_parallel_for_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_parallel_for_map_messages.cpp +++ clang/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(); Index: clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp +++ clang/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(); Index: clang/test/OpenMP/target_parallel_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_parallel_map_messages.cpp +++ clang/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(); Index: clang/test/OpenMP/target_simd_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_simd_map_messages.cpp +++ clang/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(); Index: clang/test/OpenMP/target_teams_distribute_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_map_messages.cpp +++ clang/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(); Index: clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp +++ clang/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(); Index: clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp +++ clang/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(); Index: clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp +++ clang/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(); Index: clang/test/OpenMP/target_teams_map_messages.cpp =================================================================== --- clang/test/OpenMP/target_teams_map_messages.cpp +++ clang/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();