Index: cfe/trunk/include/clang/AST/OpenMPClause.h =================================================================== --- cfe/trunk/include/clang/AST/OpenMPClause.h +++ cfe/trunk/include/clang/AST/OpenMPClause.h @@ -5082,6 +5082,9 @@ /// 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 @@ /// 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 @@ /// 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 @@ /// \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: cfe/trunk/include/clang/Basic/OpenMPKinds.h =================================================================== --- cfe/trunk/include/clang/Basic/OpenMPKinds.h +++ cfe/trunk/include/clang/Basic/OpenMPKinds.h @@ -113,6 +113,14 @@ 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: cfe/trunk/include/clang/Basic/OpenMPKinds.def =================================================================== --- cfe/trunk/include/clang/Basic/OpenMPKinds.def +++ cfe/trunk/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 @@ // 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 @@ #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: cfe/trunk/include/clang/Sema/Sema.h =================================================================== --- cfe/trunk/include/clang/Sema/Sema.h +++ cfe/trunk/include/clang/Sema/Sema.h @@ -9477,8 +9477,10 @@ 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: cfe/trunk/lib/AST/OpenMPClause.cpp =================================================================== --- cfe/trunk/lib/AST/OpenMPClause.cpp +++ cfe/trunk/lib/AST/OpenMPClause.cpp @@ -892,10 +892,11 @@ 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 @@ 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 @@ 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 @@ 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::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: cfe/trunk/lib/Basic/OpenMPKinds.cpp =================================================================== --- cfe/trunk/lib/Basic/OpenMPKinds.cpp +++ cfe/trunk/lib/Basic/OpenMPKinds.cpp @@ -122,6 +122,12 @@ .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 @@ 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 @@ 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 @@ 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: Index: cfe/trunk/lib/Parse/ParseOpenMP.cpp =================================================================== --- cfe/trunk/lib/Parse/ParseOpenMP.cpp +++ cfe/trunk/lib/Parse/ParseOpenMP.cpp @@ -2161,13 +2161,20 @@ 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 @@ /// 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: Index: cfe/trunk/lib/Sema/SemaOpenMP.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaOpenMP.cpp +++ cfe/trunk/lib/Sema/SemaOpenMP.cpp @@ -9806,7 +9806,8 @@ 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 @@ 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 @@ 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 @@ } // 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::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, Index: cfe/trunk/lib/Sema/TreeTransform.h =================================================================== --- cfe/trunk/lib/Sema/TreeTransform.h +++ cfe/trunk/lib/Sema/TreeTransform.h @@ -1915,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. @@ -8976,16 +8980,16 @@ 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: cfe/trunk/lib/Serialization/ASTReader.cpp =================================================================== --- cfe/trunk/lib/Serialization/ASTReader.cpp +++ cfe/trunk/lib/Serialization/ASTReader.cpp @@ -12448,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(); @@ -12459,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: cfe/trunk/lib/Serialization/ASTWriter.cpp =================================================================== --- cfe/trunk/lib/Serialization/ASTWriter.cpp +++ cfe/trunk/lib/Serialization/ASTWriter.cpp @@ -6885,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: cfe/trunk/test/OpenMP/declare_mapper_ast_print.c =================================================================== --- cfe/trunk/test/OpenMP/declare_mapper_ast_print.c +++ cfe/trunk/test/OpenMP/declare_mapper_ast_print.c @@ -48,8 +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) -// 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; } Index: cfe/trunk/test/OpenMP/declare_mapper_ast_print.cpp =================================================================== --- cfe/trunk/test/OpenMP/declare_mapper_ast_print.cpp +++ cfe/trunk/test/OpenMP/declare_mapper_ast_print.cpp @@ -83,6 +83,8 @@ { 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 @@ // 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 @@ // 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 @@ #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) { Index: cfe/trunk/test/OpenMP/declare_mapper_codegen.cpp =================================================================== --- cfe/trunk/test/OpenMP/declare_mapper_codegen.cpp +++ cfe/trunk/test/OpenMP/declare_mapper_codegen.cpp @@ -26,12 +26,14 @@ #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 @@ // 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) } Index: cfe/trunk/test/OpenMP/declare_mapper_messages.c =================================================================== --- cfe/trunk/test/OpenMP/declare_mapper_messages.c +++ cfe/trunk/test/OpenMP/declare_mapper_messages.c @@ -58,6 +58,15 @@ #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: cfe/trunk/test/OpenMP/declare_mapper_messages.cpp =================================================================== --- cfe/trunk/test/OpenMP/declare_mapper_messages.cpp +++ cfe/trunk/test/OpenMP/declare_mapper_messages.cpp @@ -99,6 +99,19 @@ #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'}} }