Index: include/clang/AST/OpenMPClause.h =================================================================== --- include/clang/AST/OpenMPClause.h +++ include/clang/AST/OpenMPClause.h @@ -158,8 +158,11 @@ /// This structure contains most locations needed for by an OMPVarListClause. struct OMPVarListLocTy { + /// Starting location of the clause (the clause keyword). SourceLocation StartLoc; + /// Location of '('. SourceLocation LParenLoc; + /// Ending location of the clause. SourceLocation EndLoc; OMPVarListLocTy() = default; OMPVarListLocTy(SourceLocation StartLoc, SourceLocation LParenLoc, @@ -3609,9 +3612,13 @@ /// This structure contains all sizes needed for by an /// OMPMappableExprListClause. struct OMPMappableExprListSizeTy { + /// Number of expressions listed. unsigned NumVars; + /// Number of unique base declarations. unsigned NumUniqueDeclarations; + /// Number of component lists. unsigned NumComponentLists; + /// Total number of expression components. unsigned NumComponents; OMPMappableExprListSizeTy() = default; OMPMappableExprListSizeTy(unsigned NumVars, unsigned NumUniqueDeclarations, @@ -4969,6 +4976,10 @@ /// Build clause with number of variables \a NumVars. /// + /// \param MapperQualifierLoc C++ nested name specifier for the associated + /// user-defined mapper. + /// \param MapperIdInfo The identifier of associated user-defined mapper. + /// \param MapType Map type. /// \param Locs Locations needed to build a mappable clause. It includes 1) /// StartLoc: starting location of the clause (the clause keyword); 2) /// LParenLoc: location of '('; 3) EndLoc: ending location of the clause. @@ -4977,9 +4988,12 @@ /// NumUniqueDeclarations: number of unique base declarations in this clause; /// 3) NumComponentLists: number of component lists in this clause; and 4) /// NumComponents: total number of expression components in the clause. - explicit OMPToClause(const OMPVarListLocTy &Locs, + explicit OMPToClause(NestedNameSpecifierLoc MapperQualifierLoc, + DeclarationNameInfo MapperIdInfo, + const OMPVarListLocTy &Locs, const OMPMappableExprListSizeTy &Sizes) - : OMPMappableExprListClause(OMPC_to, Locs, Sizes) {} + : OMPMappableExprListClause(OMPC_to, Locs, Sizes, &MapperQualifierLoc, + &MapperIdInfo) {} /// Build an empty clause. /// @@ -4994,7 +5008,9 @@ /// 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(); + // There are varlist_size() of expressions, and varlist_size() of + // user-defined mappers. + return 2 * varlist_size(); } size_t numTrailingObjects(OverloadToken) const { return getUniqueDeclarationsNum(); @@ -5013,10 +5029,18 @@ /// \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 UDMQualifierLoc C++ nested name specifier for the associated + /// user-defined mapper. + /// \param MapperId The identifier of associated user-defined mapper. static OMPToClause *Create(const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef Vars, ArrayRef Declarations, - MappableExprComponentListsRef ComponentLists); + MappableExprComponentListsRef ComponentLists, + ArrayRef UDMapperRefs, + NestedNameSpecifierLoc UDMQualifierLoc, + DeclarationNameInfo MapperId); /// Creates an empty clause with the place for \a NumVars variables. /// @@ -5059,6 +5083,10 @@ /// Build clause with number of variables \a NumVars. /// + /// \param MapperQualifierLoc C++ nested name specifier for the associated + /// user-defined mapper. + /// \param MapperIdInfo The identifier of associated user-defined mapper. + /// \param MapType Map type. /// \param Locs Locations needed to build a mappable clause. It includes 1) /// StartLoc: starting location of the clause (the clause keyword); 2) /// LParenLoc: location of '('; 3) EndLoc: ending location of the clause. @@ -5067,9 +5095,12 @@ /// NumUniqueDeclarations: number of unique base declarations in this clause; /// 3) NumComponentLists: number of component lists in this clause; and 4) /// NumComponents: total number of expression components in the clause. - explicit OMPFromClause(const OMPVarListLocTy &Locs, + explicit OMPFromClause(NestedNameSpecifierLoc MapperQualifierLoc, + DeclarationNameInfo MapperIdInfo, + const OMPVarListLocTy &Locs, const OMPMappableExprListSizeTy &Sizes) - : OMPMappableExprListClause(OMPC_from, Locs, Sizes) {} + : OMPMappableExprListClause(OMPC_from, Locs, Sizes, &MapperQualifierLoc, + &MapperIdInfo) {} /// Build an empty clause. /// @@ -5084,7 +5115,9 @@ /// 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(); + // There are varlist_size() of expressions, and varlist_size() of + // user-defined mappers. + return 2 * varlist_size(); } size_t numTrailingObjects(OverloadToken) const { return getUniqueDeclarationsNum(); @@ -5103,10 +5136,18 @@ /// \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 UDMQualifierLoc C++ nested name specifier for the associated + /// user-defined mapper. + /// \param MapperId The identifier of associated user-defined mapper. static OMPFromClause *Create(const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef Vars, ArrayRef Declarations, - MappableExprComponentListsRef ComponentLists); + MappableExprComponentListsRef ComponentLists, + ArrayRef UDMapperRefs, + NestedNameSpecifierLoc UDMQualifierLoc, + DeclarationNameInfo MapperId); /// Creates an empty clause with the place for \a NumVars variables. /// Index: include/clang/Basic/DiagnosticParseKinds.td =================================================================== --- include/clang/Basic/DiagnosticParseKinds.td +++ include/clang/Basic/DiagnosticParseKinds.td @@ -1175,7 +1175,7 @@ def err_omp_expected_clause: Error< "expected at least one clause on '#pragma omp %0' directive">; def err_omp_mapper_illegal_identifier : Error< - "illegal identifier on 'omp declare mapper' directive">; + "illegal OpenMP user-defined mapper identifier">; def err_omp_mapper_expected_declarator : Error< "expected declarator on 'omp declare mapper' directive">; Index: include/clang/Basic/OpenMPKinds.h =================================================================== --- include/clang/Basic/OpenMPKinds.h +++ include/clang/Basic/OpenMPKinds.h @@ -105,6 +105,22 @@ OMPC_MAP_MODIFIER_last }; +/// OpenMP modifier kind for 'to' clause. +enum OpenMPToModifierKind { +#define OPENMP_TO_MODIFIER_KIND(Name) \ + OMPC_TO_MODIFIER_##Name, +#include "clang/Basic/OpenMPKinds.def" + OMPC_TO_MODIFIER_unknown +}; + +/// OpenMP modifier kind for 'from' clause. +enum OpenMPFromModifierKind { +#define OPENMP_FROM_MODIFIER_KIND(Name) \ + OMPC_FROM_MODIFIER_##Name, +#include "clang/Basic/OpenMPKinds.def" + OMPC_FROM_MODIFIER_unknown +}; + /// OpenMP attributes for 'dist_schedule' clause. enum OpenMPDistScheduleClauseKind { #define OPENMP_DIST_SCHEDULE_KIND(Name) OMPC_DIST_SCHEDULE_##Name, Index: include/clang/Basic/OpenMPKinds.def =================================================================== --- include/clang/Basic/OpenMPKinds.def +++ include/clang/Basic/OpenMPKinds.def @@ -122,6 +122,12 @@ #ifndef OPENMP_MAP_MODIFIER_KIND #define OPENMP_MAP_MODIFIER_KIND(Name) #endif +#ifndef OPENMP_TO_MODIFIER_KIND +#define OPENMP_TO_MODIFIER_KIND(Name) +#endif +#ifndef OPENMP_FROM_MODIFIER_KIND +#define OPENMP_FROM_MODIFIER_KIND(Name) +#endif #ifndef OPENMP_DIST_SCHEDULE_KIND #define OPENMP_DIST_SCHEDULE_KIND(Name) #endif @@ -578,6 +584,12 @@ OPENMP_MAP_MODIFIER_KIND(close) OPENMP_MAP_MODIFIER_KIND(mapper) +// Modifiers for 'to' clause. +OPENMP_TO_MODIFIER_KIND(mapper) + +// Modifiers for 'from' clause. +OPENMP_FROM_MODIFIER_KIND(mapper) + // Clauses allowed for OpenMP directive 'taskloop'. OPENMP_TASKLOOP_CLAUSE(if) OPENMP_TASKLOOP_CLAUSE(shared) @@ -934,6 +946,8 @@ #undef OPENMP_FOR_SIMD_CLAUSE #undef OPENMP_MAP_KIND #undef OPENMP_MAP_MODIFIER_KIND +#undef OPENMP_TO_MODIFIER_KIND +#undef OPENMP_FROM_MODIFIER_KIND #undef OPENMP_DISTRIBUTE_CLAUSE #undef OPENMP_DIST_SCHEDULE_KIND #undef OPENMP_DEFAULTMAP_KIND Index: include/clang/Parse/Parser.h =================================================================== --- include/clang/Parse/Parser.h +++ include/clang/Parse/Parser.h @@ -2925,6 +2925,8 @@ ParsedType ObjectType, SourceLocation *TemplateKWLoc, UnqualifiedId &Result); + /// Parses the mapper modifier in map, to, and from clauses. + bool parseMapperModifier(OpenMPVarListDataTy &Data); /// 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) Index: include/clang/Sema/Sema.h =================================================================== --- include/clang/Sema/Sema.h +++ include/clang/Sema/Sema.h @@ -9471,11 +9471,16 @@ SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation MLoc, SourceLocation KindLoc, SourceLocation EndLoc); /// Called on well-formed 'to' clause. - OMPClause *ActOnOpenMPToClause(ArrayRef VarList, - const OMPVarListLocTy &Locs); + OMPClause * + ActOnOpenMPToClause(ArrayRef VarList, CXXScopeSpec &MapperIdScopeSpec, + DeclarationNameInfo &MapperId, + const OMPVarListLocTy &Locs, + ArrayRef UnresolvedMappers = llvm::None); /// Called on well-formed 'from' clause. - OMPClause *ActOnOpenMPFromClause(ArrayRef VarList, - const OMPVarListLocTy &Locs); + OMPClause *ActOnOpenMPFromClause( + ArrayRef VarList, CXXScopeSpec &MapperIdScopeSpec, + DeclarationNameInfo &MapperId, const OMPVarListLocTy &Locs, + ArrayRef UnresolvedMappers = llvm::None); /// Called on well-formed 'use_device_ptr' clause. OMPClause *ActOnOpenMPUseDevicePtrClause(ArrayRef VarList, const OMPVarListLocTy &Locs); Index: lib/AST/OpenMPClause.cpp =================================================================== --- lib/AST/OpenMPClause.cpp +++ lib/AST/OpenMPClause.cpp @@ -845,11 +845,11 @@ return new (Mem) OMPMapClause(Sizes); } -OMPToClause *OMPToClause::Create(const ASTContext &C, - const OMPVarListLocTy &Locs, - ArrayRef Vars, - ArrayRef Declarations, - MappableExprComponentListsRef ComponentLists) { +OMPToClause *OMPToClause::Create( + const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef Vars, + ArrayRef Declarations, + MappableExprComponentListsRef ComponentLists, ArrayRef UDMapperRefs, + NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId) { OMPMappableExprListSizeTy Sizes; Sizes.NumVars = Vars.size(); Sizes.NumUniqueDeclarations = getUniqueDeclarationsTotalNumber(Declarations); @@ -857,8 +857,8 @@ Sizes.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 @@ -869,13 +869,14 @@ void *Mem = C.Allocate( totalSizeToAlloc( - Sizes.NumVars, Sizes.NumUniqueDeclarations, + 2 * Sizes.NumVars, Sizes.NumUniqueDeclarations, Sizes.NumUniqueDeclarations + Sizes.NumComponentLists, Sizes.NumComponents)); - OMPToClause *Clause = new (Mem) OMPToClause(Locs, Sizes); + auto *Clause = new (Mem) OMPToClause(UDMQualifierLoc, MapperId, Locs, Sizes); Clause->setVarRefs(Vars); + Clause->setUDMapperRefs(UDMapperRefs); Clause->setClauseInfo(Declarations, ComponentLists); return Clause; } @@ -885,16 +886,17 @@ void *Mem = C.Allocate( totalSizeToAlloc( - Sizes.NumVars, Sizes.NumUniqueDeclarations, + 2 * Sizes.NumVars, Sizes.NumUniqueDeclarations, Sizes.NumUniqueDeclarations + Sizes.NumComponentLists, Sizes.NumComponents)); return new (Mem) OMPToClause(Sizes); } -OMPFromClause * -OMPFromClause::Create(const ASTContext &C, const OMPVarListLocTy &Locs, - ArrayRef Vars, ArrayRef Declarations, - MappableExprComponentListsRef ComponentLists) { +OMPFromClause *OMPFromClause::Create( + const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef Vars, + ArrayRef Declarations, + MappableExprComponentListsRef ComponentLists, ArrayRef UDMapperRefs, + NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId) { OMPMappableExprListSizeTy Sizes; Sizes.NumVars = Vars.size(); Sizes.NumUniqueDeclarations = getUniqueDeclarationsTotalNumber(Declarations); @@ -902,8 +904,8 @@ Sizes.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 @@ -914,13 +916,15 @@ void *Mem = C.Allocate( totalSizeToAlloc( - Sizes.NumVars, Sizes.NumUniqueDeclarations, + 2 * Sizes.NumVars, Sizes.NumUniqueDeclarations, Sizes.NumUniqueDeclarations + Sizes.NumComponentLists, Sizes.NumComponents)); - OMPFromClause *Clause = new (Mem) OMPFromClause(Locs, Sizes); + auto *Clause = + new (Mem) OMPFromClause(UDMQualifierLoc, MapperId, Locs, Sizes); Clause->setVarRefs(Vars); + Clause->setUDMapperRefs(UDMapperRefs); Clause->setClauseInfo(Declarations, ComponentLists); return Clause; } @@ -931,7 +935,7 @@ void *Mem = C.Allocate( totalSizeToAlloc( - Sizes.NumVars, Sizes.NumUniqueDeclarations, + 2 * Sizes.NumVars, Sizes.NumUniqueDeclarations, Sizes.NumUniqueDeclarations + Sizes.NumComponentLists, Sizes.NumComponents)); return new (Mem) OMPFromClause(Sizes); @@ -1444,7 +1448,19 @@ void OMPClausePrinter::VisitOMPToClause(OMPToClause *Node) { if (!Node->varlist_empty()) { OS << "to"; - VisitOMPClauseList(Node, '('); + DeclarationNameInfo MapperId = Node->getMapperIdInfo(); + if (MapperId.getName() && !MapperId.getName().isEmpty()) { + OS << '('; + OS << "mapper("; + NestedNameSpecifier *MapperNNS = + Node->getMapperQualifierLoc().getNestedNameSpecifier(); + if (MapperNNS) + MapperNNS->print(OS, Policy); + OS << MapperId << "):"; + VisitOMPClauseList(Node, ' '); + } else { + VisitOMPClauseList(Node, '('); + } OS << ")"; } } @@ -1452,7 +1468,19 @@ void OMPClausePrinter::VisitOMPFromClause(OMPFromClause *Node) { if (!Node->varlist_empty()) { OS << "from"; - VisitOMPClauseList(Node, '('); + DeclarationNameInfo MapperId = Node->getMapperIdInfo(); + if (MapperId.getName() && !MapperId.getName().isEmpty()) { + OS << '('; + OS << "mapper("; + NestedNameSpecifier *MapperNNS = + Node->getMapperQualifierLoc().getNestedNameSpecifier(); + if (MapperNNS) + MapperNNS->print(OS, Policy); + OS << MapperId << "):"; + VisitOMPClauseList(Node, ' '); + } else { + VisitOMPClauseList(Node, '('); + } OS << ")"; } } Index: lib/Basic/OpenMPKinds.cpp =================================================================== --- lib/Basic/OpenMPKinds.cpp +++ lib/Basic/OpenMPKinds.cpp @@ -116,6 +116,18 @@ .Case(#Name, static_cast(OMPC_MAP_MODIFIER_##Name)) #include "clang/Basic/OpenMPKinds.def" .Default(OMPC_MAP_unknown); + case OMPC_to: + return llvm::StringSwitch(Str) +#define OPENMP_TO_MODIFIER_KIND(Name) \ + .Case(#Name, static_cast(OMPC_TO_MODIFIER_##Name)) +#include "clang/Basic/OpenMPKinds.def" + .Default(OMPC_TO_MODIFIER_unknown); + case OMPC_from: + return llvm::StringSwitch(Str) +#define OPENMP_FROM_MODIFIER_KIND(Name) \ + .Case(#Name, static_cast(OMPC_FROM_MODIFIER_##Name)) +#include "clang/Basic/OpenMPKinds.def" + .Default(OMPC_FROM_MODIFIER_unknown); case OMPC_dist_schedule: return llvm::StringSwitch(Str) #define OPENMP_DIST_SCHEDULE_KIND(Name) .Case(#Name, OMPC_DIST_SCHEDULE_##Name) @@ -174,8 +186,6 @@ case OMPC_num_tasks: case OMPC_hint: case OMPC_uniform: - case OMPC_to: - case OMPC_from: case OMPC_use_device_ptr: case OMPC_is_device_ptr: case OMPC_unified_address: @@ -260,6 +270,30 @@ break; } llvm_unreachable("Invalid OpenMP 'map' clause type"); + case OMPC_to: + switch (Type) { + case OMPC_TO_MODIFIER_unknown: + return "unknown"; +#define OPENMP_TO_MODIFIER_KIND(Name) \ + case OMPC_TO_MODIFIER_##Name: \ + return #Name; +#include "clang/Basic/OpenMPKinds.def" + default: + break; + } + llvm_unreachable("Invalid OpenMP 'to' clause type"); + case OMPC_from: + switch (Type) { + case OMPC_FROM_MODIFIER_unknown: + return "unknown"; +#define OPENMP_FROM_MODIFIER_KIND(Name) \ + case OMPC_FROM_MODIFIER_##Name: \ + return #Name; +#include "clang/Basic/OpenMPKinds.def" + default: + break; + } + llvm_unreachable("Invalid OpenMP 'from' clause type"); case OMPC_dist_schedule: switch (Type) { case OMPC_DIST_SCHEDULE_unknown: @@ -333,8 +367,6 @@ case OMPC_num_tasks: case OMPC_hint: case OMPC_uniform: - case OMPC_to: - case OMPC_from: case OMPC_use_device_ptr: case OMPC_is_device_ptr: case OMPC_unified_address: Index: lib/Parse/ParseOpenMP.cpp =================================================================== --- lib/Parse/ParseOpenMP.cpp +++ lib/Parse/ParseOpenMP.cpp @@ -1959,6 +1959,34 @@ return TypeModifier; } +/// Parse the mapper modifier in map, to, and from clauses. +bool Parser::parseMapperModifier(OpenMPVarListDataTy &Data) { + // Parse '('. + BalancedDelimiterTracker T(*this, tok::l_paren, tok::colon); + if (T.expectAndConsume(diag::err_expected_lparen_after, "mapper")) { + SkipUntil(tok::colon, tok::r_paren, tok::annot_pragma_openmp_end, + StopBeforeMatch); + return true; + } + // Parse mapper-identifier + if (getLangOpts().CPlusPlus) + ParseOptionalCXXScopeSpecifier(Data.ReductionOrMapperIdScopeSpec, + /*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::r_paren, tok::annot_pragma_openmp_end, + StopBeforeMatch); + return true; + } + auto &DeclNames = Actions.getASTContext().DeclarationNames; + Data.ReductionOrMapperId = DeclarationNameInfo( + DeclNames.getIdentifier(Tok.getIdentifierInfo()), Tok.getLocation()); + ConsumeToken(); + // Parse ')'. + return T.consumeClose(); +} + /// Parse 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) @@ -1974,30 +2002,8 @@ 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); + if (parseMapperModifier(Data)) return true; - } - // Parse mapper-identifier - if (getLangOpts().CPlusPlus) - ParseOptionalCXXScopeSpecifier(Data.ReductionOrMapperIdScopeSpec, - /*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.ReductionOrMapperId = 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 @@ -2053,7 +2059,7 @@ OpenMPVarListDataTy &Data) { UnqualifiedId UnqualifiedReductionId; bool InvalidReductionId = false; - bool IsInvalidMapModifier = false; + bool IsInvalidMapperModifier = false; // Parse '('. BalancedDelimiterTracker T(*this, tok::l_paren, tok::annot_pragma_openmp_end); @@ -2142,9 +2148,11 @@ // Only parse map-type-modifier[s] and map-type if a colon is present in // the map clause. if (ColonPresent) { - IsInvalidMapModifier = parseMapTypeModifiers(Data); - if (!IsInvalidMapModifier) + IsInvalidMapperModifier = parseMapTypeModifiers(Data); + if (!IsInvalidMapperModifier) parseMapType(*this, Data); + else + SkipUntil(tok::colon, tok::annot_pragma_openmp_end, StopBeforeMatch); } if (Data.MapType == OMPC_MAP_unknown) { Data.MapType = OMPC_MAP_tofrom; @@ -2153,6 +2161,35 @@ if (Tok.is(tok::colon)) Data.ColonLoc = ConsumeToken(); + } else if (Kind == OMPC_to || Kind == OMPC_from) { + if (Tok.is(tok::identifier)) { + bool IsMapperModifier = false; + if (Kind == OMPC_to) { + auto Modifier = static_cast( + getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok))); + if (Modifier == OMPC_TO_MODIFIER_mapper) + IsMapperModifier = true; + } else { + auto Modifier = static_cast( + getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok))); + if (Modifier == OMPC_FROM_MODIFIER_mapper) + IsMapperModifier = true; + } + if (IsMapperModifier) { + // Parse the mapper modifier. + ConsumeToken(); + IsInvalidMapperModifier = parseMapperModifier(Data); + if (Tok.isNot(tok::colon)) { + if (!IsInvalidMapperModifier) + Diag(Tok, diag::warn_pragma_expected_colon) << ")"; + SkipUntil(tok::colon, tok::r_paren, tok::annot_pragma_openmp_end, + StopBeforeMatch); + } + // Consume ':'. + if (Tok.is(tok::colon)) + ConsumeToken(); + } + } } bool IsComma = @@ -2214,7 +2251,7 @@ Vars.empty()) || (Kind != OMPC_depend && Kind != OMPC_map && Vars.empty()) || (MustHaveTail && !Data.TailExpr) || InvalidReductionId || - IsInvalidMapModifier; + IsInvalidMapperModifier; } /// Parsing of OpenMP clause 'private', 'firstprivate', 'lastprivate', @@ -2247,12 +2284,12 @@ /// 'depend' '(' in | out | inout : list | source ')' /// map-clause: /// 'map' '(' [ [ always [,] ] [ close [,] ] -/// [ mapper(mapper-identifier) [,] ] +/// [ mapper '(' mapper-identifier ')' [,] ] /// to | from | tofrom | alloc | release | delete ':' ] list ')'; /// to-clause: -/// 'to' '(' list ')' +/// 'to' '(' [ mapper '(' mapper-identifier ')' ':' ] list ')' /// from-clause: -/// 'from' '(' list ')' +/// 'from' '(' [ mapper '(' mapper-identifier ')' ':' ] list ')' /// use_device_ptr-clause: /// 'use_device_ptr' '(' list ')' /// is_device_ptr-clause: Index: lib/Sema/SemaOpenMP.cpp =================================================================== --- lib/Sema/SemaOpenMP.cpp +++ lib/Sema/SemaOpenMP.cpp @@ -9802,10 +9802,12 @@ DepLinMapLoc, ColonLoc, VarList, Locs); break; case OMPC_to: - Res = ActOnOpenMPToClause(VarList, Locs); + Res = ActOnOpenMPToClause(VarList, ReductionOrMapperIdScopeSpec, + ReductionOrMapperId, Locs); break; case OMPC_from: - Res = ActOnOpenMPFromClause(VarList, Locs); + Res = ActOnOpenMPFromClause(VarList, ReductionOrMapperIdScopeSpec, + ReductionOrMapperId, Locs); break; case OMPC_use_device_ptr: Res = ActOnOpenMPUseDevicePtrClause(VarList, Locs); @@ -13056,7 +13058,7 @@ return UnresolvedLookupExpr::Create( SemaRef.Context, /*NamingClass=*/nullptr, MapperIdScopeSpec.getWithLocInContext(SemaRef.Context), MapperId, - /*ADL=*/false, /*Overloaded=*/true, URS.begin(), URS.end()); + /*ADL=*/true, /*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++ @@ -13140,17 +13142,23 @@ static void checkMappableExpressionList( Sema &SemaRef, DSAStackTy *DSAS, OpenMPClauseKind CKind, MappableVarListInfo &MVLI, SourceLocation StartLoc, + CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo MapperId, + ArrayRef UnresolvedMappers, OpenMPMapClauseKind MapType = OMPC_MAP_unknown, - bool IsMapTypeImplicit = false, CXXScopeSpec *MapperIdScopeSpec = nullptr, - const DeclarationNameInfo *MapperId = nullptr, - ArrayRef UnresolvedMappers = llvm::None) { + bool IsMapTypeImplicit = false) { // 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."); + + // If the identifier of user-defined mapper is not specified, it is "default". + // We do not change the actual name in this clause to distinguish whether a + // mapper is specified explicitly, i.e., it is not explicitly specified when + // MapperId.getName() is empty. + if (!MapperId.getName() || MapperId.getName().isEmpty()) { + auto &DeclNames = SemaRef.getASTContext().DeclarationNames; + MapperId.setName(DeclNames.getIdentifier( + &SemaRef.getASTContext().Idents.get("default"))); + } // Iterators to find the current unresolved mapper expression. auto UMIt = UnresolvedMappers.begin(), UMEnd = UnresolvedMappers.end(); @@ -13183,17 +13191,13 @@ 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); - } + // 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()); // We can only analyze this information once the missing information is // resolved. MVLI.ProcessedVarList.push_back(RE); @@ -13225,17 +13229,13 @@ 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); - } + // 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()); // Skip restriction checking for variable or field declarations MVLI.ProcessedVarList.push_back(RE); MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1); @@ -13352,18 +13352,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); } + // 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()); + // Save the current expression. MVLI.ProcessedVarList.push_back(RE); @@ -13410,17 +13408,10 @@ ++Count; } - // If the identifier of user-defined mapper is not specified, it is "default". - if (!MapperId.getName() || MapperId.getName().isEmpty()) { - auto &DeclNames = getASTContext().DeclarationNames; - MapperId.setName( - DeclNames.getIdentifier(&getASTContext().Idents.get("default"))); - } - MappableVarListInfo MVLI(VarList); checkMappableExpressionList(*this, DSAStack, OMPC_map, MVLI, Locs.StartLoc, - MapType, IsMapTypeImplicit, &MapperIdScopeSpec, - &MapperId, UnresolvedMappers); + MapperIdScopeSpec, MapperId, UnresolvedMappers, + MapType, IsMapTypeImplicit); // 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. @@ -14169,25 +14160,37 @@ } OMPClause *Sema::ActOnOpenMPToClause(ArrayRef VarList, - const OMPVarListLocTy &Locs) { + CXXScopeSpec &MapperIdScopeSpec, + DeclarationNameInfo &MapperId, + const OMPVarListLocTy &Locs, + ArrayRef UnresolvedMappers) { MappableVarListInfo MVLI(VarList); - checkMappableExpressionList(*this, DSAStack, OMPC_to, MVLI, Locs.StartLoc); + checkMappableExpressionList(*this, DSAStack, OMPC_to, MVLI, Locs.StartLoc, + MapperIdScopeSpec, MapperId, UnresolvedMappers); if (MVLI.ProcessedVarList.empty()) return nullptr; - return OMPToClause::Create(Context, Locs, MVLI.ProcessedVarList, - MVLI.VarBaseDeclarations, MVLI.VarComponents); + return OMPToClause::Create( + Context, Locs, MVLI.ProcessedVarList, MVLI.VarBaseDeclarations, + MVLI.VarComponents, MVLI.UDMapperList, + MapperIdScopeSpec.getWithLocInContext(Context), MapperId); } OMPClause *Sema::ActOnOpenMPFromClause(ArrayRef VarList, - const OMPVarListLocTy &Locs) { + CXXScopeSpec &MapperIdScopeSpec, + DeclarationNameInfo &MapperId, + const OMPVarListLocTy &Locs, + ArrayRef UnresolvedMappers) { MappableVarListInfo MVLI(VarList); - checkMappableExpressionList(*this, DSAStack, OMPC_from, MVLI, Locs.StartLoc); + checkMappableExpressionList(*this, DSAStack, OMPC_from, MVLI, Locs.StartLoc, + MapperIdScopeSpec, MapperId, UnresolvedMappers); if (MVLI.ProcessedVarList.empty()) return nullptr; - return OMPFromClause::Create(Context, Locs, MVLI.ProcessedVarList, - MVLI.VarBaseDeclarations, MVLI.VarComponents); + return OMPFromClause::Create( + Context, Locs, MVLI.ProcessedVarList, MVLI.VarBaseDeclarations, + MVLI.VarComponents, MVLI.UDMapperList, + MapperIdScopeSpec.getWithLocInContext(Context), MapperId); } OMPClause *Sema::ActOnOpenMPUseDevicePtrClause(ArrayRef VarList, Index: lib/Sema/TreeTransform.h =================================================================== --- lib/Sema/TreeTransform.h +++ lib/Sema/TreeTransform.h @@ -1902,8 +1902,12 @@ /// By default, performs semantic analysis to build the new statement. /// Subclasses may override this routine to provide different behavior. OMPClause *RebuildOMPToClause(ArrayRef VarList, - const OMPVarListLocTy &Locs) { - return getSema().ActOnOpenMPToClause(VarList, Locs); + CXXScopeSpec &MapperIdScopeSpec, + DeclarationNameInfo &MapperId, + const OMPVarListLocTy &Locs, + ArrayRef UnresolvedMappers) { + return getSema().ActOnOpenMPToClause(VarList, MapperIdScopeSpec, MapperId, + Locs, UnresolvedMappers); } /// Build a new OpenMP 'from' clause. @@ -1911,8 +1915,12 @@ /// By default, performs semantic analysis to build the new statement. /// Subclasses may override this routine to provide different behavior. OMPClause *RebuildOMPFromClause(ArrayRef VarList, - const OMPVarListLocTy &Locs) { - return getSema().ActOnOpenMPFromClause(VarList, Locs); + CXXScopeSpec &MapperIdScopeSpec, + DeclarationNameInfo &MapperId, + const OMPVarListLocTy &Locs, + ArrayRef UnresolvedMappers) { + return getSema().ActOnOpenMPFromClause(VarList, MapperIdScopeSpec, MapperId, + Locs, UnresolvedMappers); } /// Build a new OpenMP 'use_device_ptr' clause. @@ -8811,34 +8819,37 @@ C->getLParenLoc(), C->getEndLoc()); } -template -OMPClause *TreeTransform::TransformOMPMapClause(OMPMapClause *C) { - llvm::SmallVector Vars; +template +bool transformOMPMappableExprListClause( + TreeTransform &TT, OMPMappableExprListClause *C, + llvm::SmallVector &Vars, CXXScopeSpec &MapperIdScopeSpec, + DeclarationNameInfo &MapperIdInfo, + llvm::SmallVector &UnresolvedMappers) { + // Transform expressions in the list. Vars.reserve(C->varlist_size()); for (auto *VE : C->varlists()) { - ExprResult EVar = getDerived().TransformExpr(cast(VE)); + ExprResult EVar = TT.getDerived().TransformExpr(cast(VE)); if (EVar.isInvalid()) - return nullptr; + return true; Vars.push_back(EVar.get()); } + // Transform mapper scope specifier and identifier. NestedNameSpecifierLoc QualifierLoc; if (C->getMapperQualifierLoc()) { - QualifierLoc = getDerived().TransformNestedNameSpecifierLoc( + QualifierLoc = TT.getDerived().TransformNestedNameSpecifierLoc( C->getMapperQualifierLoc()); if (!QualifierLoc) - return nullptr; + return true; } - CXXScopeSpec MapperIdScopeSpec; MapperIdScopeSpec.Adopt(QualifierLoc); - DeclarationNameInfo MapperIdInfo = C->getMapperIdInfo(); + MapperIdInfo = C->getMapperIdInfo(); if (MapperIdInfo.getName()) { - MapperIdInfo = getDerived().TransformDeclarationNameInfo(MapperIdInfo); + MapperIdInfo = TT.getDerived().TransformDeclarationNameInfo(MapperIdInfo); if (!MapperIdInfo.getName()) - return nullptr; + return true; } // 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) { @@ -8846,18 +8857,31 @@ UnresolvedSet<8> Decls; for (auto *D : ULE->decls()) { NamedDecl *InstD = - cast(getDerived().TransformDecl(E->getExprLoc(), D)); + cast(TT.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())); + TT.getSema().Context, /*NamingClass=*/nullptr, + MapperIdScopeSpec.getWithLocInContext(TT.getSema().Context), + MapperIdInfo, /*ADL=*/true, ULE->isOverloaded(), Decls.begin(), + Decls.end())); } else { UnresolvedMappers.push_back(nullptr); } } + return false; +} + +template +OMPClause *TreeTransform::TransformOMPMapClause(OMPMapClause *C) { OMPVarListLocTy Locs(C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc()); + llvm::SmallVector Vars; + CXXScopeSpec MapperIdScopeSpec; + DeclarationNameInfo MapperIdInfo; + llvm::SmallVector UnresolvedMappers; + if (transformOMPMappableExprListClause( + *this, C, Vars, MapperIdScopeSpec, MapperIdInfo, UnresolvedMappers)) + return nullptr; return getDerived().RebuildOMPMapClause( C->getMapTypeModifiers(), C->getMapTypeModifiersLoc(), MapperIdScopeSpec, MapperIdInfo, C->getMapType(), C->isImplicitMapType(), C->getMapLoc(), @@ -8942,30 +8966,30 @@ template OMPClause *TreeTransform::TransformOMPToClause(OMPToClause *C) { - llvm::SmallVector Vars; - Vars.reserve(C->varlist_size()); - for (auto *VE : C->varlists()) { - ExprResult EVar = getDerived().TransformExpr(cast(VE)); - if (EVar.isInvalid()) - return 0; - Vars.push_back(EVar.get()); - } OMPVarListLocTy Locs(C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc()); - return getDerived().RebuildOMPToClause(Vars, Locs); + llvm::SmallVector Vars; + CXXScopeSpec MapperIdScopeSpec; + DeclarationNameInfo MapperIdInfo; + llvm::SmallVector UnresolvedMappers; + if (transformOMPMappableExprListClause( + *this, C, Vars, MapperIdScopeSpec, MapperIdInfo, UnresolvedMappers)) + return nullptr; + return getDerived().RebuildOMPToClause(Vars, MapperIdScopeSpec, MapperIdInfo, + Locs, UnresolvedMappers); } template OMPClause *TreeTransform::TransformOMPFromClause(OMPFromClause *C) { - llvm::SmallVector Vars; - Vars.reserve(C->varlist_size()); - for (auto *VE : C->varlists()) { - ExprResult EVar = getDerived().TransformExpr(cast(VE)); - if (EVar.isInvalid()) - return 0; - Vars.push_back(EVar.get()); - } OMPVarListLocTy Locs(C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc()); - return getDerived().RebuildOMPFromClause(Vars, Locs); + llvm::SmallVector Vars; + CXXScopeSpec MapperIdScopeSpec; + DeclarationNameInfo MapperIdInfo; + llvm::SmallVector UnresolvedMappers; + if (transformOMPMappableExprListClause( + *this, C, Vars, MapperIdScopeSpec, MapperIdInfo, UnresolvedMappers)) + return nullptr; + return getDerived().RebuildOMPFromClause( + Vars, MapperIdScopeSpec, MapperIdInfo, Locs, UnresolvedMappers); } template Index: lib/Serialization/ASTReader.cpp =================================================================== --- lib/Serialization/ASTReader.cpp +++ lib/Serialization/ASTReader.cpp @@ -12396,6 +12396,10 @@ void OMPClauseReader::VisitOMPToClause(OMPToClause *C) { C->setLParenLoc(Record.readSourceLocation()); + C->setMapperQualifierLoc(Record.readNestedNameSpecifierLoc()); + DeclarationNameInfo DNI; + Record.readDeclarationNameInfo(DNI); + C->setMapperIdInfo(DNI); auto NumVars = C->varlist_size(); auto UniqueDecls = C->getUniqueDeclarationsNum(); auto TotalLists = C->getTotalComponentListNum(); @@ -12407,6 +12411,12 @@ Vars.push_back(Record.readSubExpr()); C->setVarRefs(Vars); + SmallVector UDMappers; + UDMappers.reserve(NumVars); + for (unsigned I = 0; I < NumVars; ++I) + UDMappers.push_back(Record.readSubExpr()); + C->setUDMapperRefs(UDMappers); + SmallVector Decls; Decls.reserve(UniqueDecls); for (unsigned i = 0; i < UniqueDecls; ++i) @@ -12438,6 +12448,10 @@ void OMPClauseReader::VisitOMPFromClause(OMPFromClause *C) { C->setLParenLoc(Record.readSourceLocation()); + C->setMapperQualifierLoc(Record.readNestedNameSpecifierLoc()); + DeclarationNameInfo DNI; + Record.readDeclarationNameInfo(DNI); + C->setMapperIdInfo(DNI); auto NumVars = C->varlist_size(); auto UniqueDecls = C->getUniqueDeclarationsNum(); auto TotalLists = C->getTotalComponentListNum(); @@ -12449,6 +12463,12 @@ Vars.push_back(Record.readSubExpr()); C->setVarRefs(Vars); + SmallVector UDMappers; + UDMappers.reserve(NumVars); + for (unsigned I = 0; I < NumVars; ++I) + UDMappers.push_back(Record.readSubExpr()); + 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 @@ -6861,8 +6861,12 @@ Record.push_back(C->getTotalComponentListNum()); Record.push_back(C->getTotalComponentsNum()); Record.AddSourceLocation(C->getLParenLoc()); + Record.AddNestedNameSpecifierLoc(C->getMapperQualifierLoc()); + Record.AddDeclarationNameInfo(C->getMapperIdInfo()); 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()) @@ -6881,8 +6885,12 @@ Record.push_back(C->getTotalComponentListNum()); Record.push_back(C->getTotalComponentsNum()); Record.AddSourceLocation(C->getLParenLoc()); + Record.AddNestedNameSpecifierLoc(C->getMapperQualifierLoc()); + Record.AddDeclarationNameInfo(C->getMapperIdInfo()); 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 @@ -48,6 +48,8 @@ #pragma omp target map(mapper(default), from: dd[0:10]) // CHECK: #pragma omp target map(mapper(default),from: dd[0:10]) { dd[0].i++; } +#pragma omp target update to(mapper(id): vv) from(mapper(default): dd[0:10]) +// CHECK: #pragma omp target update to(mapper(id): vv) from(mapper(default): dd[0:10]) } return 0; } Index: test/OpenMP/declare_mapper_ast_print.cpp =================================================================== --- test/OpenMP/declare_mapper_ast_print.cpp +++ test/OpenMP/declare_mapper_ast_print.cpp @@ -81,6 +81,10 @@ { fd.a++; } #pragma omp target map(mapper(idd) alloc: fd.b) { fd.b.k++; } +#pragma omp target update to(mapper(id): fd) +#pragma omp target update to(mapper(idd): fd.b) +#pragma omp target update from(mapper(id): fd) +#pragma omp target update from(mapper(idd): fd.b) return 0; } @@ -93,6 +97,10 @@ // CHECK: } // CHECK: #pragma omp target map(mapper(id),alloc: fd) // CHECK: #pragma omp target map(mapper(idd),alloc: fd.b) +// CHECK: #pragma omp target update to(mapper(id): fd) +// CHECK: #pragma omp target update to(mapper(idd): fd.b) +// CHECK: #pragma omp target update from(mapper(id): fd) +// CHECK: #pragma omp target update from(mapper(idd): fd.b) // CHECK: } // CHECK: template<> int foo(int a) { // CHECK: #pragma omp declare mapper (id : struct foodat v) map(tofrom: v.a) @@ -103,6 +111,10 @@ // CHECK: } // CHECK: #pragma omp target map(mapper(id),alloc: fd) // CHECK: #pragma omp target map(mapper(idd),alloc: fd.b) +// CHECK: #pragma omp target update to(mapper(id): fd) +// CHECK: #pragma omp target update to(mapper(idd): fd.b) +// CHECK: #pragma omp target update from(mapper(id): fd) +// CHECK: #pragma omp target update from(mapper(idd): fd.b) // CHECK: } // CHECK: int main() { @@ -119,6 +131,14 @@ #pragma omp target map(mapper(default) tofrom: dd) // CHECK: #pragma omp target map(mapper(default),tofrom: dd) { dd.d++; } +#pragma omp target update to(mapper(N1::id) : vc) +// CHECK: #pragma omp target update to(mapper(N1::id): vc) +#pragma omp target update from(mapper(N1::id) : vc) +// CHECK: #pragma omp target update from(mapper(N1::id): vc) +#pragma omp target update to(mapper(dat::id): vvv) +// CHECK: #pragma omp target update to(mapper(dat::id): vvv) +#pragma omp target update from(mapper(dat::id): vvv) +// CHECK: #pragma omp target update from(mapper(dat::id): vvv) #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_codegen.cpp =================================================================== --- test/OpenMP/declare_mapper_codegen.cpp +++ test/OpenMP/declare_mapper_codegen.cpp @@ -26,10 +26,14 @@ #pragma omp declare mapper(id: C s) map(s.a) -// CHECK-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l49.region_id = weak constant i8 0 +// CHECK-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l54.region_id = weak constant i8 0 -// CHECK-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i{{64|32}}] [i{{64|32}} 4] -// CHECK-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 35] +// CHECK: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4] +// CHECK: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 35] +// CHECK: [[FSIZES:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4] +// CHECK: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34] +// CHECK: [[TSIZES:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4] +// CHECK: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33] // CHECK-LABEL: foo{{.*}}( void foo(int a){ @@ -46,11 +50,33 @@ // CHECK-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C** // CHECK-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]] // CHECK-DAG: store %class.C* [[VAL]], %class.C** [[CP1]] + // CHECK: call void [[KERNEL:@.+]](%class.C* [[VAL]]) #pragma omp target map(mapper(id),tofrom: c) { ++c.a; } - // CHECK: call void [[KERNEL:@.+]](%class.C* [[VAL]]) + + // CHECK-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i[[sz]]* getelementptr {{.+}}[1 x i[[sz]]]* [[FSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FTYPES]]{{.+}}) + // CHECK-DAG: [[FGEPBP]] = getelementptr inbounds {{.+}}[[FBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 + // CHECK-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 + // CHECK-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0 + // CHECK-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0 + // CHECK-DAG: [[FCBP0:%.+]] = bitcast i8** [[FBP0]] to %class.C** + // CHECK-DAG: [[FCP0:%.+]] = bitcast i8** [[FP0]] to %class.C** + // CHECK-DAG: store %class.C* [[VAL]], %class.C** [[FCBP0]] + // CHECK-DAG: store %class.C* [[VAL]], %class.C** [[FCP0]] + #pragma omp target update from(mapper(id): c) + + // CHECK-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i[[sz]]* getelementptr {{.+}}[1 x i[[sz]]]* [[TSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[TTYPES]]{{.+}}) + // CHECK-DAG: [[TGEPBP]] = getelementptr inbounds {{.+}}[[TBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 + // CHECK-DAG: [[TGEPP]] = getelementptr inbounds {{.+}}[[TP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 + // CHECK-DAG: [[TBP0:%.+]] = getelementptr inbounds {{.+}}[[TBP]], i{{.+}} 0, i{{.+}} 0 + // CHECK-DAG: [[TP0:%.+]] = getelementptr inbounds {{.+}}[[TP]], i{{.+}} 0, i{{.+}} 0 + // CHECK-DAG: [[TCBP0:%.+]] = bitcast i8** [[TBP0]] to %class.C** + // CHECK-DAG: [[TCP0:%.+]] = bitcast i8** [[TP0]] to %class.C** + // CHECK-DAG: store %class.C* [[VAL]], %class.C** [[TCBP0]] + // CHECK-DAG: store %class.C* [[VAL]], %class.C** [[TCP0]] + #pragma omp target update to(mapper(id): c) } Index: test/OpenMP/declare_mapper_messages.c =================================================================== --- test/OpenMP/declare_mapper_messages.c +++ test/OpenMP/declare_mapper_messages.c @@ -22,7 +22,7 @@ #pragma omp declare mapper(bb:struct vec v) private(v) // expected-error {{expected at least one clause on '#pragma omp declare mapper' directive}} // expected-error {{unexpected OpenMP clause 'private' in directive '#pragma omp declare mapper'}} #pragma omp declare mapper(cc:struct vec v) map(v) ( // expected-warning {{extra tokens at the end of '#pragma omp declare mapper' are ignored}} -#pragma omp declare mapper(++: struct vec v) map(v.len) // expected-error {{illegal identifier on 'omp declare mapper' directive}} +#pragma omp declare mapper(++: struct vec v) map(v.len) // expected-error {{illegal OpenMP user-defined mapper identifier}} #pragma omp declare mapper(id1: struct vec v) map(v.len, temp) // expected-error {{only variable v is allowed in map clauses of this 'omp declare mapper' directive}} #pragma omp declare mapper(default : struct vec kk) map(kk.data[0:2]) // expected-note {{previous definition is here}} #pragma omp declare mapper(struct vec v) map(v.len) // expected-error {{redefinition of user-defined mapper for type 'struct vec' with name 'default'}} @@ -50,6 +50,23 @@ {} #pragma omp target map(mapper(aa) to:vv) map(close mapper(aa) from:v1) {} + +#pragma omp target update to(mapper) // expected-error {{expected '(' after 'mapper'}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(mapper() // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(mapper:vv) // expected-error {{expected '(' after 'mapper'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(mapper(:vv) // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(mapper(aa :vv) // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(mapper(ab):vv) // expected-error {{cannot find a valid user-defined mapper for type 'struct vec' with name 'ab'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(mapper(aa):vv) + +#pragma omp target update from(mapper) // expected-error {{expected '(' after 'mapper'}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(mapper() // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(mapper:vv) // expected-error {{expected '(' after 'mapper'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(mapper(:vv) // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(mapper(aa :vv) // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(mapper(ab):vv) // expected-error {{cannot find a valid user-defined mapper for type 'struct vec' with name 'ab'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(mapper(aa) a:vv) // expected-warning {{missing ':' after ) - ignoring}} +#pragma omp target update from(mapper(aa):vv) } } return arg; Index: test/OpenMP/declare_mapper_messages.cpp =================================================================== --- test/OpenMP/declare_mapper_messages.cpp +++ test/OpenMP/declare_mapper_messages.cpp @@ -29,7 +29,7 @@ #pragma omp declare mapper(bb: vec v) private(v) // expected-error {{expected at least one clause on '#pragma omp declare mapper' directive}} // expected-error {{unexpected OpenMP clause 'private' in directive '#pragma omp declare mapper'}} #pragma omp declare mapper(cc: vec v) map(v) ( // expected-warning {{extra tokens at the end of '#pragma omp declare mapper' are ignored}} -#pragma omp declare mapper(++: vec v) map(v.len) // expected-error {{illegal identifier on 'omp declare mapper' directive}} +#pragma omp declare mapper(++: vec v) map(v.len) // expected-error {{illegal OpenMP user-defined mapper identifier}} #pragma omp declare mapper(id1: vec v) map(v.len, temp) // expected-error {{only variable v is allowed in map clauses of this 'omp declare mapper' directive}} #pragma omp declare mapper(default : vec kk) map(kk.data[0:2]) // expected-note {{previous definition is here}} #pragma omp declare mapper(vec v) map(v.len) // expected-error {{redefinition of user-defined mapper for type 'vec' with name 'default'}} @@ -74,9 +74,9 @@ {} #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(N2::) :vv) // expected-error {{use of undeclared identifier 'N2'}} expected-error {{illegal OpenMP user-defined mapper identifier}} {} -#pragma omp target map(mapper(N1::) :vv) // expected-error {{illegal identifier on 'omp declare mapper' directive}} +#pragma omp target map(mapper(N1::) :vv) // expected-error {{illegal OpenMP user-defined mapper identifier}} {} #pragma omp target map(mapper(aa) :vv) // expected-error {{missing map type}} {} @@ -86,6 +86,32 @@ {} #pragma omp target map(mapper(N1::stack::id) to:vv) {} + +#pragma omp target update to(mapper) // expected-error {{expected '(' after 'mapper'}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(mapper() // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(mapper:vv) // expected-error {{expected '(' after 'mapper'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(mapper(:vv) // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(mapper(aa :vv) // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(mapper(N2:: :vv) // expected-error {{use of undeclared identifier 'N2'}} expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(mapper(N1:: :vv) // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(mapper(N1::aa) :vv) // expected-error {{cannot find a valid user-defined mapper for type 'vec' with name 'aa'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(mapper(ab):vv) // expected-error {{cannot find a valid user-defined mapper for type 'vec' with name 'ab'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(mapper(aa) a:vv) // expected-warning {{missing ':' after ) - ignoring}} +#pragma omp target update to(mapper(aa):vv) +#pragma omp target update to(mapper(N1::stack::id) :vv) + +#pragma omp target update from(mapper) // expected-error {{expected '(' after 'mapper'}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(mapper() // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(mapper:vv) // expected-error {{expected '(' after 'mapper'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(mapper(:vv) // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(mapper(aa :vv) // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(mapper(N2:: :vv) // expected-error {{use of undeclared identifier 'N2'}} expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(mapper(N1:: :vv) // expected-error {{illegal OpenMP user-defined mapper identifier}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(mapper(N1::aa) :vv) // expected-error {{cannot find a valid user-defined mapper for type 'vec' with name 'aa'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(mapper(ab):vv) // expected-error {{cannot find a valid user-defined mapper for type 'vec' with name 'ab'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(mapper(aa) a:vv) // expected-warning {{missing ':' after ) - ignoring}} +#pragma omp target update from(mapper(aa):vv) +#pragma omp target update from(mapper(N1::stack::id) :vv) } #pragma omp declare mapper(id: vec v) map(v.len) // expected-error {{redefinition of user-defined mapper for type 'vec' with name 'id'}} }