Index: include/clang/AST/OpenMPClause.h =================================================================== --- include/clang/AST/OpenMPClause.h +++ include/clang/AST/OpenMPClause.h @@ -4051,7 +4051,7 @@ /// Define the sizes of each trailing object array except the last one. This /// is required for TrailingObjects to work properly. size_t numTrailingObjects(OverloadToken) const { - return varlist_size(); + return 2 * varlist_size(); } size_t numTrailingObjects(OverloadToken) const { return getUniqueDeclarationsNum(); @@ -4068,12 +4068,18 @@ private: /// Map-type-modifiers for the 'map' clause. OpenMPMapModifierKind MapTypeModifiers[NumberOfModifiers] = { - OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown - }; - + OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown, + OMPC_MAP_MODIFIER_unknown}; + /// Location of map-type-modifiers for the 'map' clause. SourceLocation MapTypeModifiersLoc[NumberOfModifiers]; + /// C++ nested name specifier for the assoicated user-defined mapper. + NestedNameSpecifierLoc MapperQualifierLoc; + + /// The associated user-defined mapper identifier information. + DeclarationNameInfo MapperIdInfo; + /// Map type for the 'map' clause. OpenMPMapClauseKind MapType = OMPC_MAP_unknown; @@ -4092,6 +4098,9 @@ /// /// \param MapModifiers Map-type-modifiers. /// \param MapModifiersLoc Locations of map-type-modifiers. + /// \param MapperQualifierLoc C++ nested name specifier for the assoicated + /// user-defined mapper. + /// \param MapperIdInfo The identifier of associated user-defined mapper. /// \param MapType Map type. /// \param MapTypeIsImplicit Map type is inferred implicitly. /// \param MapLoc Location of the map type. @@ -4104,6 +4113,8 @@ /// \param NumComponents Total number of expression components in the clause. explicit OMPMapClause(ArrayRef MapModifiers, ArrayRef MapModifiersLoc, + NestedNameSpecifierLoc MapperQualifierLoc, + DeclarationNameInfo MapperIdInfo, OpenMPMapClauseKind MapType, bool MapTypeIsImplicit, SourceLocation MapLoc, SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc, @@ -4112,16 +4123,16 @@ : OMPMappableExprListClause(OMPC_map, StartLoc, LParenLoc, EndLoc, NumVars, NumUniqueDeclarations, NumComponentLists, NumComponents), - 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)); + MapperQualifierLoc(MapperQualifierLoc), MapperIdInfo(MapperIdInfo), + 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)); + assert(llvm::array_lengthof(MapTypeModifiersLoc) == + MapModifiersLoc.size() && + "Unexpected number of map type modifier locations."); + llvm::copy(MapModifiersLoc, std::begin(MapTypeModifiersLoc)); } /// Build an empty clause. @@ -4170,6 +4181,38 @@ /// Set colon location. void setColonLoc(SourceLocation Loc) { ColonLoc = Loc; } + /// Set the nested name specifier of associated user-defined mapper. + void setMapperQualifierLoc(NestedNameSpecifierLoc NNSL) { + MapperQualifierLoc = NNSL; + } + + /// Set the name of associated user-defined mapper. + void setMapperIdInfo(DeclarationNameInfo MapperId) { + MapperIdInfo = MapperId; + } + + /// Get the user-defined mapper references that are in the trailing objects of + /// the class. + MutableArrayRef getUDMapperRefs() { + return llvm::makeMutableArrayRef( + getTrailingObjects() + varlist_size(), varlist_size()); + } + + /// Get the user-defined mappers references that are in the trailing objects + /// of the class. + ArrayRef getUDMapperRefs() const { + return llvm::makeArrayRef( + getTrailingObjects() + varlist_size(), varlist_size()); + } + + /// Set the user-defined mappers that are in the trailing objects of the + /// class. + void setUDMapperRefs(ArrayRef DMDs) { + assert(DMDs.size() == varlist_size() && + "Unexpected amount of user-defined mappers."); + std::copy(DMDs.begin(), DMDs.end(), getUDMapperRefs().begin()); + } + public: /// Creates clause with a list of variables \a VL. /// @@ -4179,20 +4222,26 @@ /// \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 UDMapperRefs References to user-defined mappers associated with + /// expressions used in the clause. /// \param MapModifiers Map-type-modifiers. /// \param MapModifiersLoc Location of map-type-modifiers. + /// \param UDMQualifierLoc C++ nested name specifier for the assoicated + /// user-defined mapper. + /// \param MapperId The identifier of associated user-defined mapper. /// \param Type Map type. /// \param TypeIsImplicit Map type is inferred implicitly. /// \param TypeLoc Location of the map type. - static OMPMapClause *Create(const ASTContext &C, SourceLocation StartLoc, - SourceLocation LParenLoc, SourceLocation EndLoc, - ArrayRef Vars, - ArrayRef Declarations, - MappableExprComponentListsRef ComponentLists, - ArrayRef MapModifiers, - ArrayRef MapModifiersLoc, - OpenMPMapClauseKind Type, bool TypeIsImplicit, - SourceLocation TypeLoc); + static OMPMapClause * + Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation LParenLoc, + SourceLocation EndLoc, ArrayRef Vars, + ArrayRef Declarations, + MappableExprComponentListsRef ComponentLists, + ArrayRef UDMapperRefs, + ArrayRef MapModifiers, + ArrayRef MapModifiersLoc, + NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId, + OpenMPMapClauseKind Type, bool TypeIsImplicit, SourceLocation TypeLoc); /// Creates an empty clause with the place for \a NumVars original /// expressions, \a NumUniqueDeclarations declarations, \NumComponentLists @@ -4232,7 +4281,7 @@ /// Fetches the map-type-modifier location at 'Cnt' index of array of /// modifiers' locations. /// - /// \param Cnt index for map-type-modifier location. + /// \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."); @@ -4243,18 +4292,47 @@ 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); } + /// Gets the name info for associated user-defined mapper. + const DeclarationNameInfo &getMapperIdInfo() const { return MapperIdInfo; } + + /// Gets the nested name specifier for associated user-defined mapper. + NestedNameSpecifierLoc getMapperQualifierLoc() const { + return MapperQualifierLoc; + } + /// Fetches location of clause mapping kind. SourceLocation getMapLoc() const LLVM_READONLY { return MapLoc; } /// Get colon location. SourceLocation getColonLoc() const { return ColonLoc; } + using mapperlist_iterator = MutableArrayRef::iterator; + using mapperlist_const_iterator = ArrayRef::iterator; + using mapperlist_range = llvm::iterator_range; + using mapperlist_const_range = + llvm::iterator_range; + + mapperlist_iterator mapperlist_begin() { return getUDMapperRefs().begin(); } + mapperlist_iterator mapperlist_end() { return getUDMapperRefs().end(); } + mapperlist_const_iterator mapperlist_begin() const { + return getUDMapperRefs().begin(); + } + mapperlist_const_iterator mapperlist_end() const { + return getUDMapperRefs().end(); + } + mapperlist_range mapperlists() { + return mapperlist_range(mapperlist_begin(), mapperlist_end()); + } + mapperlist_const_range mapperlists() const { + return mapperlist_const_range(mapperlist_begin(), mapperlist_end()); + } + child_range children() { return child_range( reinterpret_cast(varlist_begin()), Index: include/clang/Basic/DiagnosticParseKinds.td =================================================================== --- include/clang/Basic/DiagnosticParseKinds.td +++ include/clang/Basic/DiagnosticParseKinds.td @@ -1163,7 +1163,7 @@ 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' or 'close'">; + "incorrect map type modifier, expected 'always', 'close', or 'mapper'">; def err_omp_map_type_missing : Error< "missing map type">; def err_omp_map_type_modifier_missing : Error< Index: include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- include/clang/Basic/DiagnosticSemaKinds.td +++ include/clang/Basic/DiagnosticSemaKinds.td @@ -8990,6 +8990,8 @@ "only variable %0 is allowed in map clauses of this 'omp declare mapper' directive">; def err_omp_declare_mapper_redefinition : Error< "redefinition of user-defined mapper for type %0 with name %1">; +def err_omp_invalid_mapper: Error< + "cannot find a valid user-defined mapper for type %0 with name %1">; def err_omp_array_section_use : Error<"OpenMP array section is not allowed here">; def err_omp_typecheck_section_value : Error< "subscripted value is not an array or pointer">; Index: include/clang/Basic/OpenMPKinds.def =================================================================== --- include/clang/Basic/OpenMPKinds.def +++ include/clang/Basic/OpenMPKinds.def @@ -576,6 +576,7 @@ // Map-type-modifiers for 'map' clause. OPENMP_MAP_MODIFIER_KIND(always) OPENMP_MAP_MODIFIER_KIND(close) +OPENMP_MAP_MODIFIER_KIND(mapper) // Clauses allowed for OpenMP directive 'taskloop'. OPENMP_TASKLOOP_CLAUSE(if) Index: include/clang/Parse/Parser.h =================================================================== --- include/clang/Parse/Parser.h +++ include/clang/Parse/Parser.h @@ -2899,6 +2899,8 @@ MapTypeModifiers; SmallVector MapTypeModifiersLoc; + CXXScopeSpec MapperIdScopeSpec; + DeclarationNameInfo MapperId; OpenMPMapClauseKind MapType = OMPC_MAP_unknown; bool IsMapTypeImplicit = false; SourceLocation DepLinMapLoc; @@ -2915,6 +2917,10 @@ ParsedType ObjectType, SourceLocation *TemplateKWLoc, UnqualifiedId &Result); + /// Parses map-type-modifiers in map clause. + /// map([ [map-type-modifier[,] [map-type-modifier[,] ...] map-type : ] list) + /// where, map-type-modifier ::= always | close | mapper(mapper-identifier) + bool parseMapTypeModifiers(OpenMPVarListDataTy &Data); private: //===--------------------------------------------------------------------===// Index: include/clang/Sema/Sema.h =================================================================== --- include/clang/Sema/Sema.h +++ include/clang/Sema/Sema.h @@ -9344,6 +9344,7 @@ OpenMPLinearClauseKind LinKind, ArrayRef MapTypeModifiers, ArrayRef MapTypeModifiersLoc, + CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo &MapperId, OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, SourceLocation DepLinMapLoc); /// Called on well-formed 'private' clause. @@ -9426,13 +9427,14 @@ SourceLocation LParenLoc, SourceLocation EndLoc); /// Called on well-formed 'map' clause. - OMPClause * - ActOnOpenMPMapClause(ArrayRef MapTypeModifiers, - ArrayRef MapTypeModifiersLoc, - OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, - SourceLocation MapLoc, SourceLocation ColonLoc, - ArrayRef VarList, SourceLocation StartLoc, - SourceLocation LParenLoc, SourceLocation EndLoc); + OMPClause *ActOnOpenMPMapClause( + ArrayRef MapTypeModifiers, + ArrayRef MapTypeModifiersLoc, + CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo &MapperId, + OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, + SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef VarList, + SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc, + ArrayRef UnresolvedMappers = llvm::None); /// Called on well-formed 'num_teams' clause. OMPClause *ActOnOpenMPNumTeamsClause(Expr *NumTeams, SourceLocation StartLoc, SourceLocation LParenLoc, Index: lib/AST/OpenMPClause.cpp =================================================================== --- lib/AST/OpenMPClause.cpp +++ lib/AST/OpenMPClause.cpp @@ -791,15 +791,15 @@ return TotalNum; } -OMPMapClause * -OMPMapClause::Create(const ASTContext &C, SourceLocation StartLoc, - SourceLocation LParenLoc, SourceLocation EndLoc, - ArrayRef Vars, ArrayRef Declarations, - MappableExprComponentListsRef ComponentLists, - ArrayRef MapModifiers, - ArrayRef MapModifiersLoc, - OpenMPMapClauseKind Type, bool TypeIsImplicit, - SourceLocation TypeLoc) { +OMPMapClause *OMPMapClause::Create( + const ASTContext &C, SourceLocation StartLoc, SourceLocation LParenLoc, + SourceLocation EndLoc, ArrayRef Vars, + ArrayRef Declarations, + MappableExprComponentListsRef ComponentLists, ArrayRef UDMapperRefs, + ArrayRef MapModifiers, + ArrayRef MapModifiersLoc, + NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId, + OpenMPMapClauseKind Type, bool TypeIsImplicit, SourceLocation TypeLoc) { unsigned NumVars = Vars.size(); unsigned NumUniqueDeclarations = getUniqueDeclarationsTotalNumber(Declarations); @@ -807,8 +807,8 @@ unsigned NumComponents = getComponentsTotalNumber(ComponentLists); // We need to allocate: - // NumVars x Expr* - we have an original list expression for each clause list - // entry. + // 2 x NumVars x Expr* - we have an original list expression and an associated + // user-defined mapper for each clause list entry. // NumUniqueDeclarations x ValueDecl* - unique base declarations associated // with each component list. // (NumUniqueDeclarations + NumComponentLists) x unsigned - we specify the @@ -819,14 +819,15 @@ void *Mem = C.Allocate( totalSizeToAlloc( - NumVars, NumUniqueDeclarations, + 2 * NumVars, NumUniqueDeclarations, NumUniqueDeclarations + NumComponentLists, NumComponents)); OMPMapClause *Clause = new (Mem) OMPMapClause( - MapModifiers, MapModifiersLoc, Type, TypeIsImplicit, TypeLoc, StartLoc, - LParenLoc, EndLoc, NumVars, NumUniqueDeclarations, NumComponentLists, - NumComponents); + MapModifiers, MapModifiersLoc, UDMQualifierLoc, MapperId, Type, + TypeIsImplicit, TypeLoc, StartLoc, LParenLoc, EndLoc, NumVars, + NumUniqueDeclarations, NumComponentLists, NumComponents); Clause->setVarRefs(Vars); + Clause->setUDMapperRefs(UDMapperRefs); Clause->setClauseInfo(Declarations, ComponentLists); Clause->setMapType(Type); Clause->setMapLoc(TypeLoc); @@ -840,7 +841,7 @@ void *Mem = C.Allocate( totalSizeToAlloc( - NumVars, NumUniqueDeclarations, + 2 * NumVars, NumUniqueDeclarations, NumUniqueDeclarations + NumComponentLists, NumComponents)); return new (Mem) OMPMapClause(NumVars, NumUniqueDeclarations, NumComponentLists, NumComponents); @@ -1432,6 +1433,14 @@ if (Node->getMapTypeModifier(I) != OMPC_MAP_MODIFIER_unknown) { OS << getOpenMPSimpleClauseTypeName(OMPC_map, Node->getMapTypeModifier(I)); + if (Node->getMapTypeModifier(I) == OMPC_MAP_MODIFIER_mapper) { + OS << '('; + NestedNameSpecifier *MapperNNS = + Node->getMapperQualifierLoc().getNestedNameSpecifier(); + if (MapperNNS) + MapperNNS->print(OS, Policy); + OS << Node->getMapperIdInfo() << ')'; + } OS << ','; } } Index: lib/Parse/ParseOpenMP.cpp =================================================================== --- lib/Parse/ParseOpenMP.cpp +++ lib/Parse/ParseOpenMP.cpp @@ -1957,36 +1957,62 @@ /// 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); +/// where, map-type-modifier ::= always | close | mapper(mapper-identifier) +bool Parser::parseMapTypeModifiers(OpenMPVarListDataTy &Data) { + while (getCurToken().isNot(tok::colon)) { + OpenMPMapModifierKind TypeModifier = isMapModifier(*this); if (TypeModifier == OMPC_MAP_MODIFIER_always || TypeModifier == OMPC_MAP_MODIFIER_close) { Data.MapTypeModifiers.push_back(TypeModifier); Data.MapTypeModifiersLoc.push_back(Tok.getLocation()); - P.ConsumeToken(); + ConsumeToken(); + } else if (TypeModifier == OMPC_MAP_MODIFIER_mapper) { + Data.MapTypeModifiers.push_back(TypeModifier); + Data.MapTypeModifiersLoc.push_back(Tok.getLocation()); + ConsumeToken(); + // Parse '('. + BalancedDelimiterTracker T(*this, tok::l_paren, tok::colon); + if (T.expectAndConsume(diag::err_expected_lparen_after, + getOpenMPSimpleClauseTypeName( + OMPC_map, OMPC_MAP_MODIFIER_mapper))) { + SkipUntil(tok::colon, tok::annot_pragma_openmp_end, StopBeforeMatch); + return true; + } + // Parse mapper-identifier + if (getLangOpts().CPlusPlus) + ParseOptionalCXXScopeSpecifier(Data.MapperIdScopeSpec, + /*ObjectType=*/nullptr, + /*EnteringContext=*/false); + if (Tok.isNot(tok::identifier) && Tok.isNot(tok::kw_default)) { + Diag(Tok.getLocation(), diag::err_omp_mapper_illegal_identifier); + SkipUntil(tok::colon, tok::annot_pragma_openmp_end, StopBeforeMatch); + return true; + } + auto &DeclNames = Actions.getASTContext().DeclarationNames; + Data.MapperId = DeclarationNameInfo( + DeclNames.getIdentifier(Tok.getIdentifierInfo()), Tok.getLocation()); + ConsumeToken(); + // Parse ')'. + T.consumeClose(); } 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(); + Diag(Tok, diag::err_omp_map_type_modifier_missing); + 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(); + return false; + Diag(Tok, diag::err_omp_unknown_map_type_modifier); + ConsumeToken(); } - if (P.getCurToken().is(tok::comma)) - P.ConsumeToken(); + if (getCurToken().is(tok::comma)) + ConsumeToken(); } + return false; } /// Checks if the token is a valid map-type. @@ -2023,6 +2049,7 @@ OpenMPVarListDataTy &Data) { UnqualifiedId UnqualifiedReductionId; bool InvalidReductionId = false; + bool IsInvalidMapModifier = false; // Parse '('. BalancedDelimiterTracker T(*this, tok::l_paren, tok::annot_pragma_openmp_end); @@ -2111,8 +2138,9 @@ // 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); + IsInvalidMapModifier = parseMapTypeModifiers(Data); + if (!IsInvalidMapModifier) + parseMapType(*this, Data); } if (Data.MapType == OMPC_MAP_unknown) { Data.MapType = OMPC_MAP_tofrom; @@ -2181,7 +2209,8 @@ return (Kind == OMPC_depend && Data.DepKind != OMPC_DEPEND_unknown && Vars.empty()) || (Kind != OMPC_depend && Kind != OMPC_map && Vars.empty()) || - (MustHaveTail && !Data.TailExpr) || InvalidReductionId; + (MustHaveTail && !Data.TailExpr) || InvalidReductionId || + IsInvalidMapModifier; } /// Parsing of OpenMP clause 'private', 'firstprivate', 'lastprivate', @@ -2214,6 +2243,7 @@ /// 'depend' '(' in | out | inout : list | source ')' /// map-clause: /// 'map' '(' [ [ always [,] ] [ close [,] ] +/// [ mapper(mapper-identifier) [,] ] /// to | from | tofrom | alloc | release | delete ':' ] list ')'; /// to-clause: /// 'to' '(' list ')' @@ -2244,7 +2274,7 @@ return Actions.ActOnOpenMPVarListClause( Kind, Vars, Data.TailExpr, Loc, LOpen, Data.ColonLoc, Data.RLoc, Data.ReductionIdScopeSpec, Data.ReductionId, Data.DepKind, Data.LinKind, - Data.MapTypeModifiers, Data.MapTypeModifiersLoc, Data.MapType, - Data.IsMapTypeImplicit, Data.DepLinMapLoc); + Data.MapTypeModifiers, Data.MapTypeModifiersLoc, Data.MapperIdScopeSpec, + Data.MapperId, Data.MapType, Data.IsMapTypeImplicit, Data.DepLinMapLoc); } Index: lib/Sema/SemaLookup.cpp =================================================================== --- lib/Sema/SemaLookup.cpp +++ lib/Sema/SemaLookup.cpp @@ -1071,6 +1071,17 @@ I = IdResolver.begin(Name), IEnd = IdResolver.end(); + if (NameKind == LookupOMPMapperName) { + // Skip declarations in inner scopes. + for (; I != IEnd; ++I) { + Scope *CurS = S; + while (CurS && !CurS->isDeclScope(*I)) + CurS = CurS->getParent(); + if (CurS) + break; + } + } + // First we lookup local scope. // We don't consider using-directives, as per 7.3.4.p1 [namespace.udir] // ...During unqualified name lookup (3.4.1), the names appear as if @@ -1779,6 +1790,14 @@ else if (NameKind == LookupObjCImplicitSelfParam && !isa(*I)) continue; + else if (NameKind == LookupOMPMapperName) { + // Skip declarations in inner scopes. + Scope *CurS = S; + while (CurS && !CurS->isDeclScope(*I)) + CurS = CurS->getParent(); + if (!CurS) + continue; + } R.addDecl(D); Index: lib/Sema/SemaOpenMP.cpp =================================================================== --- lib/Sema/SemaOpenMP.cpp +++ lib/Sema/SemaOpenMP.cpp @@ -3451,11 +3451,13 @@ } } if (!ImplicitMaps.empty()) { + CXXScopeSpec MapperIdScopeSpec; + DeclarationNameInfo MapperId; if (OMPClause *Implicit = ActOnOpenMPMapClause( - llvm::None, llvm::None, OMPC_MAP_tofrom, - /*IsMapTypeImplicit=*/true, SourceLocation(), SourceLocation(), - ImplicitMaps, SourceLocation(), SourceLocation(), - SourceLocation())) { + llvm::None, llvm::None, MapperIdScopeSpec, MapperId, + OMPC_MAP_tofrom, /*IsMapTypeImplicit=*/true, SourceLocation(), + SourceLocation(), ImplicitMaps, SourceLocation(), + SourceLocation(), SourceLocation())) { ClausesWithImplicit.emplace_back(Implicit); ErrorFound |= cast(Implicit)->varlist_size() != ImplicitMaps.size(); @@ -9740,6 +9742,7 @@ OpenMPLinearClauseKind LinKind, ArrayRef MapTypeModifiers, ArrayRef MapTypeModifiersLoc, + CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo &MapperId, OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, SourceLocation DepLinMapLoc) { OMPClause *Res = nullptr; @@ -9792,7 +9795,8 @@ StartLoc, LParenLoc, EndLoc); break; case OMPC_map: - Res = ActOnOpenMPMapClause(MapTypeModifiers, MapTypeModifiersLoc, MapType, + Res = ActOnOpenMPMapClause(MapTypeModifiers, MapTypeModifiersLoc, + MapperIdScopeSpec, MapperId, MapType, IsMapTypeImplicit, DepLinMapLoc, ColonLoc, VarList, StartLoc, LParenLoc, EndLoc); break; @@ -10622,8 +10626,8 @@ } // namespace template -static T filterLookupForUDR(SmallVectorImpl &Lookups, - const llvm::function_ref Gen) { +static T filterLookupForUDReductionAndMapper( + SmallVectorImpl &Lookups, const llvm::function_ref Gen) { for (U &Set : Lookups) { for (auto *D : Set) { if (T Res = Gen(cast(D))) @@ -10743,7 +10747,7 @@ if (SemaRef.CurContext->isDependentContext() || Ty->isDependentType() || Ty->isInstantiationDependentType() || Ty->containsUnexpandedParameterPack() || - filterLookupForUDR(Lookups, [](ValueDecl *D) { + filterLookupForUDReductionAndMapper(Lookups, [](ValueDecl *D) { return !D->isInvalidDecl() && (D->getType()->isDependentType() || D->getType()->isInstantiationDependentType() || @@ -10792,7 +10796,7 @@ } // Perform ADL. argumentDependentLookup(SemaRef, ReductionId, Loc, Ty, Lookups); - if (auto *VD = filterLookupForUDR( + if (auto *VD = filterLookupForUDReductionAndMapper( Lookups, [&SemaRef, Ty](ValueDecl *D) -> ValueDecl * { if (!D->isInvalidDecl() && SemaRef.Context.hasSameType(D->getType(), Ty)) @@ -10801,7 +10805,7 @@ })) return SemaRef.BuildDeclRefExpr(VD, VD->getType().getNonReferenceType(), VK_LValue, Loc); - if (auto *VD = filterLookupForUDR( + if (auto *VD = filterLookupForUDReductionAndMapper( Lookups, [&SemaRef, Ty, Loc](ValueDecl *D) -> ValueDecl * { if (!D->isInvalidDecl() && SemaRef.IsDerivedFrom(Loc, Ty, D->getType()) && @@ -12998,6 +13002,106 @@ return FoundError; } +// Look up the user-defined mapper given the mapper name and mapped type, and +// build a reference to it. +ExprResult buildUserDefinedMapperRef(Sema &SemaRef, Scope *S, + CXXScopeSpec &MapperIdScopeSpec, + DeclarationNameInfo &MapperId, + QualType Type, Expr *UnresolvedMapper) { + if (MapperIdScopeSpec.isInvalid()) + return ExprError(); + // Find all user-defined mappers with the given MapperId. + SmallVector, 4> Lookups; + LookupResult Lookup(SemaRef, MapperId, Sema::LookupOMPMapperName); + Lookup.suppressDiagnostics(); + if (S) { + while (S && SemaRef.LookupParsedName(Lookup, S, &MapperIdScopeSpec)) { + NamedDecl *D = Lookup.getRepresentativeDecl(); + while (S && !S->isDeclScope(D)) + S = S->getParent(); + if (S) + S = S->getParent(); + Lookups.emplace_back(); + Lookups.back().append(Lookup.begin(), Lookup.end()); + Lookup.clear(); + } + } else if (auto *ULE = cast_or_null(UnresolvedMapper)) { + // Extract the user-defined mappers with the given MapperId. + Lookups.push_back(UnresolvedSet<8>()); + for (NamedDecl *D : ULE->decls()) { + auto *DMD = cast(D); + assert(DMD && "Expect valid OMPDeclareMapperDecl during instantiation."); + Lookups.back().addDecl(DMD); + } + } + // Defer the lookup for dependent types. The results will be passed through + // UnresolvedMapper on instantiation. + if (SemaRef.CurContext->isDependentContext() || Type->isDependentType() || + Type->isInstantiationDependentType() || + Type->containsUnexpandedParameterPack() || + filterLookupForUDReductionAndMapper(Lookups, [](ValueDecl *D) { + return !D->isInvalidDecl() && + (D->getType()->isDependentType() || + D->getType()->isInstantiationDependentType() || + D->getType()->containsUnexpandedParameterPack()); + })) { + UnresolvedSet<8> URS; + for (const UnresolvedSet<8> &Set : Lookups) { + if (Set.empty()) + continue; + URS.append(Set.begin(), Set.end()); + } + return UnresolvedLookupExpr::Create( + SemaRef.Context, /*NamingClass=*/nullptr, + MapperIdScopeSpec.getWithLocInContext(SemaRef.Context), MapperId, + /*ADL=*/false, /*Overloaded=*/true, URS.begin(), URS.end()); + } + // [OpenMP 5.0], 2.19.7.3 declare mapper Directive, Restrictions + // The type must be of struct, union or class type in C and C++ + if (!Type->isStructureOrClassType() && !Type->isUnionType()) + return ExprEmpty(); + // Return the first user-defined mapper with the desired type. + SourceLocation Loc = MapperId.getLoc(); + if (auto *VD = filterLookupForUDReductionAndMapper( + Lookups, [&SemaRef, Type](ValueDecl *D) -> ValueDecl * { + if (!D->isInvalidDecl() && + SemaRef.Context.hasSameType(D->getType(), Type)) + return D; + return nullptr; + })) + return SemaRef.BuildDeclRefExpr(VD, Type, VK_LValue, Loc); + // Find the first user-defined mapper with a type derived from the desired + // type. + if (auto *VD = filterLookupForUDReductionAndMapper( + Lookups, [&SemaRef, Type, Loc](ValueDecl *D) -> ValueDecl * { + if (!D->isInvalidDecl() && + SemaRef.IsDerivedFrom(Loc, Type, D->getType()) && + !Type.isMoreQualifiedThan(D->getType())) + return D; + return nullptr; + })) { + CXXBasePaths Paths(/*FindAmbiguities=*/true, /*RecordPaths=*/true, + /*DetectVirtual=*/false); + if (SemaRef.IsDerivedFrom(Loc, Type, VD->getType(), Paths)) { + if (!Paths.isAmbiguous(SemaRef.Context.getCanonicalType( + VD->getType().getUnqualifiedType()))) { + if (SemaRef.CheckBaseClassAccess( + Loc, VD->getType(), Type, Paths.front(), + /*DiagID=*/0) != Sema::AR_inaccessible) { + return SemaRef.BuildDeclRefExpr(VD, Type, VK_LValue, Loc); + } + } + } + } + // Report error if a mapper is specified, but cannot be found. + if (MapperIdScopeSpec.isSet() || MapperId.getAsString() != "default") { + SemaRef.Diag(Loc, diag::err_omp_invalid_mapper) + << Type << MapperId.getName(); + return ExprError(); + } + return ExprEmpty(); +} + namespace { // Utility struct that gathers all the related lists associated with a mappable // expression. @@ -13010,6 +13114,8 @@ OMPClauseMappableExprCommon::MappableExprComponentLists VarComponents; // The base declaration of the variable. SmallVector VarBaseDeclarations; + // The reference to the user-defined mapper associated with every expression. + SmallVector UDMapperList; MappableVarListInfo(ArrayRef VarList) : VarList(VarList) { // We have a list of components and base declarations for each entry in the @@ -13021,19 +13127,30 @@ } // Check the validity of the provided variable list for the provided clause kind -// \a CKind. In the check process the valid expressions, and mappable expression -// components and variables are extracted and used to fill \a Vars, -// \a ClauseComponents, and \a ClauseBaseDeclarations. \a MapType and -// \a IsMapTypeImplicit are expected to be valid if the clause kind is 'map'. -static void -checkMappableExpressionList(Sema &SemaRef, DSAStackTy *DSAS, - OpenMPClauseKind CKind, MappableVarListInfo &MVLI, - SourceLocation StartLoc, - OpenMPMapClauseKind MapType = OMPC_MAP_unknown, - bool IsMapTypeImplicit = false) { +// \a CKind. In the check process the valid expressions, mappable expression +// components, variables, and user-defined mappers are extracted and used to +// fill \a ProcessedVarList, \a VarComponents, \a VarBaseDeclarations, and \a +// UDMapperList in MVLI. \a MapType, \a IsMapTypeImplicit, \a MapperIdScopeSpec, +// and \a MapperId are expected to be valid if the clause kind is 'map'. +static void checkMappableExpressionList( + Sema &SemaRef, DSAStackTy *DSAS, OpenMPClauseKind CKind, + MappableVarListInfo &MVLI, SourceLocation StartLoc, + OpenMPMapClauseKind MapType = OMPC_MAP_unknown, + bool IsMapTypeImplicit = false, CXXScopeSpec *MapperIdScopeSpec = nullptr, + DeclarationNameInfo *MapperId = nullptr, + ArrayRef UnresolvedMappers = llvm::None) { // We only expect mappable expressions in 'to', 'from', and 'map' clauses. assert((CKind == OMPC_map || CKind == OMPC_to || CKind == OMPC_from) && "Unexpected clause kind with mappable expressions!"); + assert( + ((CKind == OMPC_map && MapperIdScopeSpec && MapperId) || + (CKind != OMPC_map && !MapperIdScopeSpec && !MapperId)) && + "Map clauses and only map clauses have user-defined mapper identifiers."); + + // Iterators to find the current unresolved mapper expression. + auto UMIt = UnresolvedMappers.begin(), UMEnd = UnresolvedMappers.end(); + bool UpdateUMIt = false; + Expr *UnresolvedMapper = nullptr; // Keep track of the mappable components and base declarations in this clause. // Each entry in the list is going to have a list of components associated. We @@ -13045,11 +13162,33 @@ assert(RE && "Null expr in omp to/from/map clause"); SourceLocation ELoc = RE->getExprLoc(); + // Find the current unresolved mapper expression. + if (UpdateUMIt && UMIt != UMEnd) { + UMIt++; + assert( + UMIt != UMEnd && + "Expect the size of UnresolvedMappers to match with that of VarList"); + } + UpdateUMIt = true; + if (UMIt != UMEnd) + UnresolvedMapper = *UMIt; + const Expr *VE = RE->IgnoreParenLValueCasts(); if (VE->isValueDependent() || VE->isTypeDependent() || VE->isInstantiationDependent() || VE->containsUnexpandedParameterPack()) { + if (CKind == OMPC_map) { + // Try to find the associated user-defined mapper. + ExprResult ER = buildUserDefinedMapperRef( + SemaRef, DSAS->getCurScope(), *MapperIdScopeSpec, *MapperId, + VE->getType().getCanonicalType(), UnresolvedMapper); + if (ER.isInvalid()) + continue; + MVLI.UDMapperList.push_back(ER.get()); + } else { + MVLI.UDMapperList.push_back(nullptr); + } // We can only analyze this information once the missing information is // resolved. MVLI.ProcessedVarList.push_back(RE); @@ -13081,6 +13220,17 @@ if (const auto *TE = dyn_cast(BE)) { // Add store "this" pointer to class in DSAStackTy for future checking DSAS->addMappedClassesQualTypes(TE->getType()); + if (CKind == OMPC_map) { + // Try to find the associated user-defined mapper. + ExprResult ER = buildUserDefinedMapperRef( + SemaRef, DSAS->getCurScope(), *MapperIdScopeSpec, *MapperId, + VE->getType().getCanonicalType(), UnresolvedMapper); + if (ER.isInvalid()) + continue; + MVLI.UDMapperList.push_back(ER.get()); + } else { + MVLI.UDMapperList.push_back(nullptr); + } // Skip restriction checking for variable or field declarations MVLI.ProcessedVarList.push_back(RE); MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1); @@ -13197,6 +13347,16 @@ continue; } } + + // Try to find the associated user-defined mapper. + ExprResult ER = buildUserDefinedMapperRef( + SemaRef, DSAS->getCurScope(), *MapperIdScopeSpec, *MapperId, + Type.getCanonicalType(), UnresolvedMapper); + if (ER.isInvalid()) + continue; + MVLI.UDMapperList.push_back(ER.get()); + } else { + MVLI.UDMapperList.push_back(nullptr); } // Save the current expression. @@ -13218,19 +13378,17 @@ } } -OMPClause * -Sema::ActOnOpenMPMapClause(ArrayRef MapTypeModifiers, - ArrayRef MapTypeModifiersLoc, - OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, - SourceLocation MapLoc, SourceLocation ColonLoc, - ArrayRef VarList, SourceLocation StartLoc, - SourceLocation LParenLoc, SourceLocation EndLoc) { - MappableVarListInfo MVLI(VarList); - checkMappableExpressionList(*this, DSAStack, OMPC_map, MVLI, StartLoc, - MapType, IsMapTypeImplicit); - - OpenMPMapModifierKind Modifiers[] = { OMPC_MAP_MODIFIER_unknown, - OMPC_MAP_MODIFIER_unknown }; +OMPClause *Sema::ActOnOpenMPMapClause( + ArrayRef MapTypeModifiers, + ArrayRef MapTypeModifiersLoc, + CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo &MapperId, + OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, SourceLocation MapLoc, + SourceLocation ColonLoc, ArrayRef VarList, SourceLocation StartLoc, + SourceLocation LParenLoc, SourceLocation EndLoc, + ArrayRef UnresolvedMappers) { + OpenMPMapModifierKind Modifiers[] = {OMPC_MAP_MODIFIER_unknown, + OMPC_MAP_MODIFIER_unknown, + OMPC_MAP_MODIFIER_unknown}; SourceLocation ModifiersLoc[OMPMapClause::NumberOfModifiers]; // Process map-type-modifiers, flag errors for duplicate modifiers. @@ -13248,12 +13406,25 @@ ++Count; } + // If the identifier of user-defined mapper is not specified, it is "default". + if (MapperId.getName().isEmpty()) { + auto &DeclNames = getASTContext().DeclarationNames; + MapperId.setName( + DeclNames.getIdentifier(&getASTContext().Idents.get("default"))); + } + + MappableVarListInfo MVLI(VarList); + checkMappableExpressionList(*this, DSAStack, OMPC_map, MVLI, StartLoc, + MapType, IsMapTypeImplicit, &MapperIdScopeSpec, + &MapperId, UnresolvedMappers); + // 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, Modifiers, ModifiersLoc, - MapType, IsMapTypeImplicit, MapLoc); + return OMPMapClause::Create( + Context, StartLoc, LParenLoc, EndLoc, MVLI.ProcessedVarList, + MVLI.VarBaseDeclarations, MVLI.VarComponents, MVLI.UDMapperList, + Modifiers, ModifiersLoc, MapperIdScopeSpec.getWithLocInContext(Context), + MapperId, MapType, IsMapTypeImplicit, MapLoc); } QualType Sema::ActOnOpenMPDeclareReductionType(SourceLocation TyLoc, Index: lib/Sema/SemaTemplateInstantiateDecl.cpp =================================================================== --- lib/Sema/SemaTemplateInstantiateDecl.cpp +++ lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -2989,11 +2989,19 @@ } if (!IsCorrect) break; + NestedNameSpecifierLoc NewQualifierLoc = + SemaRef.SubstNestedNameSpecifierLoc(OldC->getMapperQualifierLoc(), + TemplateArgs); + CXXScopeSpec SS; + SS.Adopt(NewQualifierLoc); + DeclarationNameInfo NewNameInfo = SemaRef.SubstDeclarationNameInfo( + OldC->getMapperIdInfo(), TemplateArgs); OMPClause *NewC = SemaRef.ActOnOpenMPMapClause( OldC->getMapTypeModifiers(), OldC->getMapTypeModifiersLoc(), - OldC->getMapType(), OldC->isImplicitMapType(), OldC->getMapLoc(), - OldC->getColonLoc(), NewVars, OldC->getBeginLoc(), - OldC->getLParenLoc(), OldC->getEndLoc()); + SS, NewNameInfo, OldC->getMapType(), + OldC->isImplicitMapType(), OldC->getMapLoc(), OldC->getColonLoc(), + NewVars, OldC->getBeginLoc(), OldC->getLParenLoc(), + OldC->getEndLoc()); Clauses.push_back(NewC); } SemaRef.EndOpenMPDSABlock(nullptr); Index: lib/Sema/TreeTransform.h =================================================================== --- lib/Sema/TreeTransform.h +++ lib/Sema/TreeTransform.h @@ -1795,17 +1795,18 @@ /// /// By default, performs semantic analysis to build the new OpenMP clause. /// Subclasses may override this routine to provide different behavior. - OMPClause * - RebuildOMPMapClause(ArrayRef MapTypeModifiers, - ArrayRef MapTypeModifiersLoc, - OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, - SourceLocation MapLoc, SourceLocation ColonLoc, - ArrayRef VarList, SourceLocation StartLoc, - SourceLocation LParenLoc, SourceLocation EndLoc) { - return getSema().ActOnOpenMPMapClause(MapTypeModifiers, MapTypeModifiersLoc, - MapType, IsMapTypeImplicit, MapLoc, - ColonLoc, VarList, StartLoc, - LParenLoc, EndLoc); + OMPClause *RebuildOMPMapClause( + ArrayRef MapTypeModifiers, + ArrayRef MapTypeModifiersLoc, + CXXScopeSpec MapperIdScopeSpec, DeclarationNameInfo MapperId, + OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, + SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef VarList, + SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc, + ArrayRef UnresolvedMappers) { + return getSema().ActOnOpenMPMapClause( + MapTypeModifiers, MapTypeModifiersLoc, MapperIdScopeSpec, MapperId, + MapType, IsMapTypeImplicit, MapLoc, ColonLoc, VarList, StartLoc, + LParenLoc, EndLoc, UnresolvedMappers); } /// Build a new OpenMP 'num_teams' clause. @@ -8808,10 +8809,47 @@ return nullptr; Vars.push_back(EVar.get()); } + NestedNameSpecifierLoc QualifierLoc; + if (C->getMapperQualifierLoc()) { + QualifierLoc = getDerived().TransformNestedNameSpecifierLoc( + C->getMapperQualifierLoc()); + if (!QualifierLoc) + return nullptr; + } + CXXScopeSpec MapperIdScopeSpec; + MapperIdScopeSpec.Adopt(QualifierLoc); + DeclarationNameInfo MapperIdInfo = C->getMapperIdInfo(); + if (MapperIdInfo.getName()) { + MapperIdInfo = getDerived().TransformDeclarationNameInfo(MapperIdInfo); + if (!MapperIdInfo.getName()) + return nullptr; + } + // Build a list of all candidate OMPDeclareMapperDecls, which is provided by + // the previous user-defined mapper lookup in dependent environment. + llvm::SmallVector UnresolvedMappers; + for (auto *E : C->mapperlists()) { + // Transform all the decls. + if (E) { + auto *ULE = cast(E); + UnresolvedSet<8> Decls; + for (auto *D : ULE->decls()) { + NamedDecl *InstD = + cast(getDerived().TransformDecl(E->getExprLoc(), D)); + Decls.addDecl(InstD, InstD->getAccess()); + } + UnresolvedMappers.push_back(UnresolvedLookupExpr::Create( + SemaRef.Context, /*NamingClass=*/nullptr, + MapperIdScopeSpec.getWithLocInContext(SemaRef.Context), MapperIdInfo, + /*ADL=*/false, ULE->isOverloaded(), Decls.begin(), Decls.end())); + } else { + UnresolvedMappers.push_back(nullptr); + } + } return getDerived().RebuildOMPMapClause( - C->getMapTypeModifiers(), C->getMapTypeModifiersLoc(), C->getMapType(), - C->isImplicitMapType(), C->getMapLoc(), C->getColonLoc(), Vars, - C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc()); + C->getMapTypeModifiers(), C->getMapTypeModifiersLoc(), MapperIdScopeSpec, + MapperIdInfo, C->getMapType(), C->isImplicitMapType(), C->getMapLoc(), + C->getColonLoc(), Vars, C->getBeginLoc(), C->getLParenLoc(), + C->getEndLoc(), UnresolvedMappers); } template Index: lib/Serialization/ASTReader.cpp =================================================================== --- lib/Serialization/ASTReader.cpp +++ lib/Serialization/ASTReader.cpp @@ -12288,6 +12288,10 @@ I, static_cast(Record.readInt())); C->setMapTypeModifierLoc(I, Record.readSourceLocation()); } + C->setMapperQualifierLoc(Record.readNestedNameSpecifierLoc()); + DeclarationNameInfo DNI; + Record.readDeclarationNameInfo(DNI); + C->setMapperIdInfo(DNI); C->setMapType( static_cast(Record.readInt())); C->setMapLoc(Record.readSourceLocation()); @@ -12303,6 +12307,12 @@ Vars.push_back(Record.readExpr()); C->setVarRefs(Vars); + SmallVector UDMappers; + UDMappers.reserve(NumVars); + for (unsigned I = 0; I < NumVars; ++I) + UDMappers.push_back(Record.readExpr()); + C->setUDMapperRefs(UDMappers); + SmallVector Decls; Decls.reserve(UniqueDecls); for (unsigned i = 0; i < UniqueDecls; ++i) Index: lib/Serialization/ASTWriter.cpp =================================================================== --- lib/Serialization/ASTWriter.cpp +++ lib/Serialization/ASTWriter.cpp @@ -6785,11 +6785,15 @@ Record.push_back(C->getMapTypeModifier(I)); Record.AddSourceLocation(C->getMapTypeModifierLoc(I)); } + Record.AddNestedNameSpecifierLoc(C->getMapperQualifierLoc()); + Record.AddDeclarationNameInfo(C->getMapperIdInfo()); Record.push_back(C->getMapType()); Record.AddSourceLocation(C->getMapLoc()); Record.AddSourceLocation(C->getColonLoc()); for (auto *E : C->varlists()) Record.AddStmt(E); + for (auto *E : C->mapperlists()) + Record.AddStmt(E); for (auto *D : C->all_decls()) Record.AddDeclRef(D); for (auto N : C->all_num_lists()) Index: test/OpenMP/declare_mapper_ast_print.c =================================================================== --- test/OpenMP/declare_mapper_ast_print.c +++ test/OpenMP/declare_mapper_ast_print.c @@ -40,6 +40,14 @@ { #pragma omp declare mapper(id: struct vec v) map(v.len) // CHECK: #pragma omp declare mapper (id : struct vec v) map(tofrom: v.len) + struct vec vv; + struct dat dd[10]; +#pragma omp target map(mapper(id) alloc: vv) +// CHECK: #pragma omp target map(mapper(id),alloc: vv) + { vv.len++; } +#pragma omp target map(mapper(default) from: dd[0:10]) +// CHECK: #pragma omp target map(mapper(default),from: dd[0:10]) + { dd[0].i++; } } return 0; } Index: test/OpenMP/declare_mapper_ast_print.cpp =================================================================== --- test/OpenMP/declare_mapper_ast_print.cpp +++ test/OpenMP/declare_mapper_ast_print.cpp @@ -21,6 +21,13 @@ }; // CHECK: }; +// CHECK: class vecchild : public N1::vec { +class vecchild : public vec { +public: + int lenc; +}; +// CHECK: }; + #pragma omp declare mapper(id: vec v) map(v.len) // CHECK: #pragma omp declare mapper (id : N1::vec v) map(tofrom: v.len){{$}} }; @@ -56,34 +63,62 @@ template T foo(T a) { + struct foodatchild { + T k; + }; struct foodat { T a; + struct foodatchild b; }; -#pragma omp declare mapper(struct foodat v) map(v.a) +#pragma omp declare mapper(id: struct foodat v) map(v.a) +#pragma omp declare mapper(idd: struct foodatchild v) map(v.k) #pragma omp declare mapper(id: N1::vec v) map(v.len) { #pragma omp declare mapper(id: N1::vec v) map(v.len) } + struct foodat fd; +#pragma omp target map(mapper(id) alloc: fd) + { fd.a++; } +#pragma omp target map(mapper(idd) alloc: fd.b) + { fd.b.k++; } return 0; } // CHECK: template T foo(T a) { -// CHECK: #pragma omp declare mapper (default : struct foodat v) map(tofrom: v.a) +// CHECK: #pragma omp declare mapper (id : struct foodat v) map(tofrom: v.a) +// CHECK: #pragma omp declare mapper (idd : struct foodatchild v) map(tofrom: v.k) // CHECK: #pragma omp declare mapper (id : N1::vec v) map(tofrom: v.len) // CHECK: { // CHECK: #pragma omp declare mapper (id : N1::vec v) map(tofrom: v.len) // CHECK: } +// CHECK: #pragma omp target map(mapper(id),alloc: fd) +// CHECK: #pragma omp target map(mapper(idd),alloc: fd.b) // CHECK: } // CHECK: template<> int foo(int a) { -// CHECK: #pragma omp declare mapper (default : struct foodat v) map(tofrom: v.a) +// CHECK: #pragma omp declare mapper (id : struct foodat v) map(tofrom: v.a) +// CHECK: #pragma omp declare mapper (idd : struct foodatchild v) map(tofrom: v.k) // CHECK: #pragma omp declare mapper (id : N1::vec v) map(tofrom: v.len) // CHECK: { // CHECK: #pragma omp declare mapper (id : N1::vec v) map(tofrom: v.len) // CHECK: } +// CHECK: #pragma omp target map(mapper(id),alloc: fd) +// CHECK: #pragma omp target map(mapper(idd),alloc: fd.b) // CHECK: } // CHECK: int main() { int main() { + N1::vec vv, vvv; + N1::vecchild vc; + dat dd; +#pragma omp target map(mapper(N1::id) tofrom: vv) map(mapper(dat::id) alloc: vvv) +// CHECK: #pragma omp target map(mapper(N1::id),tofrom: vv) map(mapper(dat::id),alloc: vvv) + { vv.len++; } +#pragma omp target map(mapper(N1::id) tofrom: vc) +// CHECK: #pragma omp target map(mapper(N1::id),tofrom: vc) + { vc.len++; } +#pragma omp target map(mapper(default) tofrom: dd) +// CHECK: #pragma omp target map(mapper(default),tofrom: dd) + { dd.d++; } #pragma omp declare mapper(id: N1::vec v) map(v.len) // CHECK: #pragma omp declare mapper (id : N1::vec v) map(tofrom: v.len) { Index: test/OpenMP/declare_mapper_messages.c =================================================================== --- test/OpenMP/declare_mapper_messages.c +++ test/OpenMP/declare_mapper_messages.c @@ -35,6 +35,21 @@ #pragma omp declare mapper(id: struct vec v) map(v.len) // expected-error {{redefinition of user-defined mapper for type 'struct vec' with name 'id'}} { #pragma omp declare mapper(id: struct vec v) map(v.len) + struct vec vv, v1; +#pragma omp target map(mapper) // expected-error {{use of undeclared identifier 'mapper'}} + {} +#pragma omp target map(mapper:vv) // expected-error {{expected '(' after 'mapper'}} + {} +#pragma omp target map(mapper( :vv) // expected-error {{expected expression}} expected-error {{expected ')'}} expected-warning {{implicit declaration of function 'mapper' is invalid in C99}} expected-note {{to match this '('}} + {} +#pragma omp target map(mapper(aa :vv) // expected-error {{use of undeclared identifier 'aa'}} expected-error {{expected ')'}} expected-warning {{implicit declaration of function 'mapper' is invalid in C99}} expected-note {{to match this '('}} + {} +#pragma omp target map(mapper(ab) :vv) // expected-error {{missing map type}} expected-error {{cannot find a valid user-defined mapper for type 'struct vec' with name 'ab'}} + {} +#pragma omp target map(mapper(aa) :vv) // expected-error {{missing map type}} + {} +#pragma omp target map(mapper(aa) to:vv) map(close mapper(aa) from:v1) + {} } } return arg; Index: test/OpenMP/declare_mapper_messages.cpp =================================================================== --- test/OpenMP/declare_mapper_messages.cpp +++ test/OpenMP/declare_mapper_messages.cpp @@ -63,6 +63,29 @@ #pragma omp declare mapper(id: vec v) map(v.len) // expected-note {{previous definition is here}} { #pragma omp declare mapper(id: vec v) map(v.len) + vec vv, v1; +#pragma omp target map(mapper) // expected-error {{use of undeclared identifier 'mapper'}} + {} +#pragma omp target map(mapper:vv) // expected-error {{expected '(' after 'mapper'}} + {} +#pragma omp target map(mapper( :vv) // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} + {} +#pragma omp target map(mapper(aa :vv) // expected-error {{use of undeclared identifier 'aa'}} expected-error {{expected ')'}} expected-note {{to match this '('}} + {} +#pragma omp target map(mapper(ab) :vv) // expected-error {{missing map type}} expected-error {{cannot find a valid user-defined mapper for type 'vec' with name 'ab'}} + {} +#pragma omp target map(mapper(N2::) :vv) // expected-error {{use of undeclared identifier 'N2'}} expected-error {{illegal identifier on 'omp declare mapper' directive}} + {} +#pragma omp target map(mapper(N1::) :vv) // expected-error {{illegal identifier on 'omp declare mapper' directive}} + {} +#pragma omp target map(mapper(aa) :vv) // expected-error {{missing map type}} + {} +#pragma omp target map(mapper(N1::aa) alloc:vv) // expected-error {{cannot find a valid user-defined mapper for type 'vec' with name 'aa'}} + {} +#pragma omp target map(mapper(aa) to:vv) map(close mapper(aa) from:v1) + {} +#pragma omp target map(mapper(N1::stack::id) to:vv) + {} } #pragma omp declare mapper(id: vec v) map(v.len) // expected-error {{redefinition of user-defined mapper for type 'vec' with name 'id'}} } Index: test/OpenMP/target_map_messages.cpp =================================================================== --- test/OpenMP/target_map_messages.cpp +++ test/OpenMP/target_map_messages.cpp @@ -110,17 +110,17 @@ {} #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( d, f, bf: a) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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( , f, : a) // expected-error {{missing map type modifier}} expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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(always tofrom close: a) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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(tofrom from: a) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} {} #pragma omp target map(close bf: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} {} @@ -516,14 +516,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' or 'close'}} expected-error {{missing map type}} +#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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(tofrom, close: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}} #pragma omp target data map(close, tofrom: close, tofrom, x) foo(); return 0; @@ -613,13 +613,13 @@ #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' or 'close'}} expected-error {{missing map type}} +#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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(tofrom, close: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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}} {} Index: test/OpenMP/target_parallel_for_map_messages.cpp =================================================================== --- test/OpenMP/target_parallel_for_map_messages.cpp +++ 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' or 'close'}} expected-error {{missing map type}} +#pragma omp target parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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' or 'close'}} expected-error {{missing map type}} +#pragma omp target parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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: test/OpenMP/target_parallel_for_simd_map_messages.cpp =================================================================== --- test/OpenMP/target_parallel_for_simd_map_messages.cpp +++ 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' or 'close'}} expected-error {{missing map type}} +#pragma omp target parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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' or 'close'}} expected-error {{missing map type}} +#pragma omp target parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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: test/OpenMP/target_parallel_map_messages.cpp =================================================================== --- test/OpenMP/target_parallel_map_messages.cpp +++ 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' or 'close'}} expected-error {{missing map type}} +#pragma omp target parallel map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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' or 'close'}} expected-error {{missing map type}} +#pragma omp target parallel map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}} foo(); #pragma omp target parallel map(always, tofrom: always, tofrom, x) foo(); Index: test/OpenMP/target_simd_map_messages.cpp =================================================================== --- test/OpenMP/target_simd_map_messages.cpp +++ 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' or 'close'}} expected-error {{missing map type}} +#pragma omp target simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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' or 'close'}} expected-error {{missing map type}} +#pragma omp target simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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: test/OpenMP/target_teams_distribute_map_messages.cpp =================================================================== --- test/OpenMP/target_teams_distribute_map_messages.cpp +++ 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' or 'close'}} expected-error {{missing map type}} +#pragma omp target teams distribute map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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' or 'close'}} expected-error {{missing map type}} +#pragma omp target teams distribute map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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: test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp =================================================================== --- test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp +++ 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' or 'close'}} expected-error {{missing map type}} +#pragma omp target teams distribute parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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' or 'close'}} expected-error {{missing map type}} +#pragma omp target teams distribute parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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: test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp =================================================================== --- test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp +++ 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' or 'close'}} expected-error {{missing map type}} +#pragma omp target teams distribute parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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' or 'close'}} expected-error {{missing map type}} +#pragma omp target teams distribute parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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: test/OpenMP/target_teams_distribute_simd_map_messages.cpp =================================================================== --- test/OpenMP/target_teams_distribute_simd_map_messages.cpp +++ 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' or 'close'}} expected-error {{missing map type}} +#pragma omp target teams distribute simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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' or 'close'}} expected-error {{missing map type}} +#pragma omp target teams distribute simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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: test/OpenMP/target_teams_map_messages.cpp =================================================================== --- test/OpenMP/target_teams_map_messages.cpp +++ 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' or 'close'}} expected-error {{missing map type}} +#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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' or 'close'}} expected-error {{missing map type}} +#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} 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();