From 01e162c2965d114aa4e21c91b6341454210ad8fb Mon Sep 17 00:00:00 2001 From: Michael Kruse Date: Mon, 25 Feb 2019 20:34:15 +0000 Subject: [PATCH] [OpenMP 5.0] Parsing/sema support for from clause with mapper modifier. This patch implements the parsing and sema support for the OpenMP 'from'-clause with potential user-defined mappers attached. User-defined mappers are a new feature in OpenMP 5.0. A 'from'-clause can have an explicit or implicit associated mapper, which instructs the compiler to generate and use customized mapping functions. An example is shown below: struct S { int len; int *d; }; #pragma omp declare mapper(id: struct S s) map(s, s.d[0:s.len]) struct S ss; #pragma omp target update from(mapper(id): ss) // use the mapper with name 'id' to map ss from device Contributed-by: Lingda Li Differential Revision: https://reviews.llvm.org/D58638 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@354817 91177308-0d34-0410-b5e6-96231b3b80d8 --- include/clang/AST/OpenMPClause.h | 24 +++++++-- include/clang/Basic/OpenMPKinds.def | 7 +++ include/clang/Basic/OpenMPKinds.h | 8 +++ include/clang/Sema/Sema.h | 6 ++- lib/AST/OpenMPClause.cpp | 35 +++++++++---- lib/Basic/OpenMPKinds.cpp | 20 +++++++- lib/Parse/ParseOpenMP.cpp | 19 ++++--- lib/Sema/SemaOpenMP.cpp | 63 +++++++++++------------- lib/Sema/TreeTransform.h | 26 +++++----- lib/Serialization/ASTReader.cpp | 10 ++++ lib/Serialization/ASTWriter.cpp | 4 ++ test/OpenMP/declare_mapper_ast_print.c | 4 +- test/OpenMP/declare_mapper_ast_print.cpp | 11 +++++ test/OpenMP/declare_mapper_codegen.cpp | 15 +++++- test/OpenMP/declare_mapper_messages.c | 9 ++++ test/OpenMP/declare_mapper_messages.cpp | 13 +++++ 16 files changed, 203 insertions(+), 71 deletions(-) diff --git a/include/clang/AST/OpenMPClause.h b/include/clang/AST/OpenMPClause.h index 3d6fa0ebc3..1aa1c82b1a 100644 --- a/include/clang/AST/OpenMPClause.h +++ b/include/clang/AST/OpenMPClause.h @@ -5082,6 +5082,9 @@ class OMPFromClause final /// 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 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. @@ -5090,9 +5093,12 @@ class OMPFromClause final /// 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. /// @@ -5107,7 +5113,9 @@ class OMPFromClause final /// 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(); @@ -5126,10 +5134,18 @@ public: /// \param Vars The original expression used in the clause. /// \param Declarations Declarations used in the clause. /// \param ComponentLists Component lists used in the clause. + /// \param 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. /// diff --git a/include/clang/Basic/OpenMPKinds.def b/include/clang/Basic/OpenMPKinds.def index bd3eff45a6..b8de5df02c 100644 --- a/include/clang/Basic/OpenMPKinds.def +++ b/include/clang/Basic/OpenMPKinds.def @@ -125,6 +125,9 @@ #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 @@ -584,6 +587,9 @@ 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) @@ -941,6 +947,7 @@ OPENMP_DECLARE_MAPPER_CLAUSE(map) #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 diff --git a/include/clang/Basic/OpenMPKinds.h b/include/clang/Basic/OpenMPKinds.h index 2e23fffedb..01bd860f19 100644 --- a/include/clang/Basic/OpenMPKinds.h +++ b/include/clang/Basic/OpenMPKinds.h @@ -113,6 +113,14 @@ enum OpenMPToModifierKind { 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, diff --git a/include/clang/Sema/Sema.h b/include/clang/Sema/Sema.h index 276e91ae01..f6eee83819 100644 --- a/include/clang/Sema/Sema.h +++ b/include/clang/Sema/Sema.h @@ -9477,8 +9477,10 @@ public: 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); diff --git a/lib/AST/OpenMPClause.cpp b/lib/AST/OpenMPClause.cpp index 0f1cb0258a..b02a888a8c 100644 --- a/lib/AST/OpenMPClause.cpp +++ b/lib/AST/OpenMPClause.cpp @@ -892,10 +892,11 @@ OMPToClause *OMPToClause::CreateEmpty(const ASTContext &C, 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); @@ -903,8 +904,8 @@ OMPFromClause::Create(const ASTContext &C, const OMPVarListLocTy &Locs, 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 @@ -915,13 +916,15 @@ OMPFromClause::Create(const ASTContext &C, const OMPVarListLocTy &Locs, 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; } @@ -932,7 +935,7 @@ OMPFromClause::CreateEmpty(const ASTContext &C, 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); @@ -1465,7 +1468,19 @@ void OMPClausePrinter::VisitOMPToClause(OMPToClause *Node) { 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 << ")"; } } diff --git a/lib/Basic/OpenMPKinds.cpp b/lib/Basic/OpenMPKinds.cpp index 47830c324d..b10c6d2498 100644 --- a/lib/Basic/OpenMPKinds.cpp +++ b/lib/Basic/OpenMPKinds.cpp @@ -122,6 +122,12 @@ unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, .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) @@ -180,7 +186,6 @@ unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, case OMPC_num_tasks: case OMPC_hint: case OMPC_uniform: - case OMPC_from: case OMPC_use_device_ptr: case OMPC_is_device_ptr: case OMPC_unified_address: @@ -277,6 +282,18 @@ const char *clang::getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind, 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: @@ -350,7 +367,6 @@ const char *clang::getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind, case OMPC_num_tasks: case OMPC_hint: case OMPC_uniform: - case OMPC_from: case OMPC_use_device_ptr: case OMPC_is_device_ptr: case OMPC_unified_address: diff --git a/lib/Parse/ParseOpenMP.cpp b/lib/Parse/ParseOpenMP.cpp index 68843f0c65..782ad12aac 100644 --- a/lib/Parse/ParseOpenMP.cpp +++ b/lib/Parse/ParseOpenMP.cpp @@ -2161,13 +2161,20 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind, if (Tok.is(tok::colon)) Data.ColonLoc = ConsumeToken(); - } else if (Kind == OMPC_to) { + } else if (Kind == OMPC_to || Kind == OMPC_from) { if (Tok.is(tok::identifier)) { bool IsMapperModifier = false; - auto Modifier = static_cast( - getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok))); - if (Modifier == OMPC_TO_MODIFIER_mapper) - IsMapperModifier = true; + 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(); @@ -2282,7 +2289,7 @@ bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind, /// to-clause: /// '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: diff --git a/lib/Sema/SemaOpenMP.cpp b/lib/Sema/SemaOpenMP.cpp index 5a6891d01a..76141afd59 100644 --- a/lib/Sema/SemaOpenMP.cpp +++ b/lib/Sema/SemaOpenMP.cpp @@ -9806,7 +9806,8 @@ OMPClause *Sema::ActOnOpenMPVarListClause( 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); @@ -13190,15 +13191,13 @@ static void checkMappableExpressionList( if (VE->isValueDependent() || VE->isTypeDependent() || VE->isInstantiationDependent() || VE->containsUnexpandedParameterPack()) { - if (CKind != OMPC_from) { - // 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()); - } + // 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); @@ -13230,15 +13229,13 @@ static void checkMappableExpressionList( 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_from) { - // 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()); - } + // 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); @@ -13358,14 +13355,12 @@ static void checkMappableExpressionList( } // Try to find the associated user-defined mapper. - if (CKind != OMPC_from) { - ExprResult ER = buildUserDefinedMapperRef( - SemaRef, DSAS->getCurScope(), MapperIdScopeSpec, MapperId, - Type.getCanonicalType(), UnresolvedMapper); - if (ER.isInvalid()) - continue; - MVLI.UDMapperList.push_back(ER.get()); - } + 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); @@ -14182,18 +14177,20 @@ OMPClause *Sema::ActOnOpenMPToClause(ArrayRef VarList, } OMPClause *Sema::ActOnOpenMPFromClause(ArrayRef VarList, - const OMPVarListLocTy &Locs) { + CXXScopeSpec &MapperIdScopeSpec, + DeclarationNameInfo &MapperId, + const OMPVarListLocTy &Locs, + ArrayRef UnresolvedMappers) { MappableVarListInfo MVLI(VarList); - CXXScopeSpec MapperIdScopeSpec; - DeclarationNameInfo MapperId; - ArrayRef UnresolvedMappers; 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, diff --git a/lib/Sema/TreeTransform.h b/lib/Sema/TreeTransform.h index b8f922a76e..386b23db3a 100644 --- a/lib/Sema/TreeTransform.h +++ b/lib/Sema/TreeTransform.h @@ -1915,8 +1915,12 @@ public: /// 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. @@ -8976,16 +8980,16 @@ OMPClause *TreeTransform::TransformOMPToClause(OMPToClause *C) { 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 diff --git a/lib/Serialization/ASTReader.cpp b/lib/Serialization/ASTReader.cpp index bed58d0e10..9626418f68 100644 --- a/lib/Serialization/ASTReader.cpp +++ b/lib/Serialization/ASTReader.cpp @@ -12448,6 +12448,10 @@ void OMPClauseReader::VisitOMPToClause(OMPToClause *C) { 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(); @@ -12459,6 +12463,12 @@ void OMPClauseReader::VisitOMPFromClause(OMPFromClause *C) { 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) diff --git a/lib/Serialization/ASTWriter.cpp b/lib/Serialization/ASTWriter.cpp index 1575a400ba..3d19c9d823 100644 --- a/lib/Serialization/ASTWriter.cpp +++ b/lib/Serialization/ASTWriter.cpp @@ -6885,8 +6885,12 @@ void OMPClauseWriter::VisitOMPFromClause(OMPFromClause *C) { 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()) diff --git a/test/OpenMP/declare_mapper_ast_print.c b/test/OpenMP/declare_mapper_ast_print.c index f053044a09..e82bc4c2d3 100644 --- a/test/OpenMP/declare_mapper_ast_print.c +++ b/test/OpenMP/declare_mapper_ast_print.c @@ -48,8 +48,8 @@ int main() { #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) -// CHECK: #pragma omp target update to(mapper(id): vv) +#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; } diff --git a/test/OpenMP/declare_mapper_ast_print.cpp b/test/OpenMP/declare_mapper_ast_print.cpp index 1e05fe5469..6462fa38d8 100644 --- a/test/OpenMP/declare_mapper_ast_print.cpp +++ b/test/OpenMP/declare_mapper_ast_print.cpp @@ -83,6 +83,8 @@ T foo(T a) { { 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; } @@ -97,6 +99,8 @@ T foo(T a) { // 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) @@ -109,6 +113,8 @@ T foo(T a) { // 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() { @@ -131,6 +137,11 @@ int main() { #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(N1::id) : vc) +// CHECK: #pragma omp target update from(mapper(N1::id): vc) +#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) { diff --git a/test/OpenMP/declare_mapper_codegen.cpp b/test/OpenMP/declare_mapper_codegen.cpp index 4946ab2cce..6f1d6ec8fd 100644 --- a/test/OpenMP/declare_mapper_codegen.cpp +++ b/test/OpenMP/declare_mapper_codegen.cpp @@ -26,12 +26,14 @@ public: #pragma omp declare mapper(id: C s) map(s.a) -// CHECK-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l52.region_id = weak constant i8 0 +// CHECK-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l54.region_id = weak constant i8 0 // CHECK: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4] // CHECK: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 35] // CHECK: [[TSIZES:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4] // CHECK: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33] +// CHECK: [[FSIZES:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4] +// CHECK: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 34] // CHECK-LABEL: foo{{.*}}( void foo(int a){ @@ -64,6 +66,17 @@ void foo(int a){ // 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) + + // 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) } diff --git a/test/OpenMP/declare_mapper_messages.c b/test/OpenMP/declare_mapper_messages.c index 95014b0d8f..4df14b62d5 100644 --- a/test/OpenMP/declare_mapper_messages.c +++ b/test/OpenMP/declare_mapper_messages.c @@ -58,6 +58,15 @@ int fun(int arg) { #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; diff --git a/test/OpenMP/declare_mapper_messages.cpp b/test/OpenMP/declare_mapper_messages.cpp index d7aa073da9..bcb5ac463a 100644 --- a/test/OpenMP/declare_mapper_messages.cpp +++ b/test/OpenMP/declare_mapper_messages.cpp @@ -99,6 +99,19 @@ int fun(int arg) { #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'}} } -- 2.40.0