[OpenMP 5.0] Parsing/sema support for from clause with mapper modifier

[lildmh]

Summary

This patch implements the parsing and sema support for OpenMP from clause with potential user-defined mappers attached. User defined mapper is 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

[ABataev]

LG

cfe/trunk/include/clang/AST/OpenMPClause.h

: public OMPMappableExprListClause<OMPFromClause>,
           OMPClauseMappableExprCommon::MappableComponent> {
   friend class OMPClauseReader;
   friend OMPMappableExprListClause;
   friend OMPVarListClause;
   friend TrailingObjects;
 
   /// 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.
   /// \param Sizes All required sizes to build a mappable clause. It includes 1)
   /// NumVars: number of expressions listed in this clause; 2)
   /// 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.
   ///
   /// \param Sizes All required sizes to build a mappable clause. It includes 1)
   /// NumVars: number of expressions listed in this clause; 2)
   /// 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 OMPMappableExprListSizeTy &Sizes)
       : OMPMappableExprListClause(OMPC_from, OMPVarListLocTy(), Sizes) {}
 
   /// Define the sizes of each trailing object array except the last one. This
   /// is required for TrailingObjects to work properly.
   size_t numTrailingObjects(OverloadToken<Expr *>) 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<ValueDecl *>) const {
     return getUniqueDeclarationsNum();
   }
   size_t numTrailingObjects(OverloadToken<unsigned>) const {
     return getUniqueDeclarationsNum() + getTotalComponentListNum();
   }
 
 public:
   /// Creates clause with a list of variables \a Vars.
   ///
   /// \param C AST context.
   /// \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.
   /// \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<Expr *> Vars,
                                ArrayRef<ValueDecl *> Declarations,
-                               MappableExprComponentListsRef ComponentLists);
+                               MappableExprComponentListsRef ComponentLists,
+                               ArrayRef<Expr *> UDMapperRefs,
+                               NestedNameSpecifierLoc UDMQualifierLoc,
+                               DeclarationNameInfo MapperId);
 
   /// Creates an empty clause with the place for \a NumVars variables.
   ///
   /// \param C AST context.
   /// \param Sizes All required sizes to build a mappable clause. It includes 1)
   /// NumVars: number of expressions listed in this clause; 2)
   /// NumUniqueDeclarations: number of unique base declarations in this clause;
   /// 3) NumComponentLists: number of component lists in this clause; and 4)

cfe/trunk/include/clang/Basic/OpenMPKinds.h

 /// 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,
 #include "clang/Basic/OpenMPKinds.def"
   OMPC_DIST_SCHEDULE_unknown
 };
 
 /// OpenMP attributes for 'defaultmap' clause.

cfe/trunk/include/clang/Basic/OpenMPKinds.def

 #define OPENMP_MAP_KIND(Name)
 #endif
 #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
 #ifndef OPENMP_DEFAULTMAP_KIND
 #define OPENMP_DEFAULTMAP_KIND(Name)
 #endif
 #ifndef OPENMP_ATOMIC_DEFAULT_MEM_ORDER_KIND
 #define OPENMP_ATOMIC_DEFAULT_MEM_ORDER_KIND(Name)

 // Map-type-modifiers for 'map' clause.
 OPENMP_MAP_MODIFIER_KIND(always)
 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)
 OPENMP_TASKLOOP_CLAUSE(private)
 OPENMP_TASKLOOP_CLAUSE(firstprivate)
 OPENMP_TASKLOOP_CLAUSE(lastprivate)
 OPENMP_TASKLOOP_CLAUSE(default)
 OPENMP_TASKLOOP_CLAUSE(collapse)

 #undef OPENMP_TARGET_PARALLEL_FOR_CLAUSE
 #undef OPENMP_TEAMS_CLAUSE
 #undef OPENMP_SIMD_CLAUSE
 #undef OPENMP_FOR_CLAUSE
 #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
 #undef OPENMP_DEFAULTMAP_MODIFIER
 #undef OPENMP_TARGET_UPDATE_CLAUSE
 #undef OPENMP_DISTRIBUTE_PARALLEL_FOR_CLAUSE
 #undef OPENMP_DISTRIBUTE_PARALLEL_FOR_SIMD_CLAUSE
 #undef OPENMP_DISTRIBUTE_SIMD_CLAUSE

cfe/trunk/include/clang/Sema/Sema.h

OMPClause *ActOnOpenMPDefaultmapClause(
       SourceLocation KindLoc, SourceLocation EndLoc);
   /// Called on well-formed 'to' clause.
   OMPClause *
   ActOnOpenMPToClause(ArrayRef<Expr *> VarList, CXXScopeSpec &MapperIdScopeSpec,
                       DeclarationNameInfo &MapperId,
                       const OMPVarListLocTy &Locs,
                       ArrayRef<Expr *> UnresolvedMappers = llvm::None);
   /// Called on well-formed 'from' clause.
-  OMPClause *ActOnOpenMPFromClause(ArrayRef<Expr *> VarList,
-                                   const OMPVarListLocTy &Locs);
+  OMPClause *ActOnOpenMPFromClause(
+      ArrayRef<Expr *> VarList, CXXScopeSpec &MapperIdScopeSpec,
+      DeclarationNameInfo &MapperId, const OMPVarListLocTy &Locs,
+      ArrayRef<Expr *> UnresolvedMappers = llvm::None);
   /// Called on well-formed 'use_device_ptr' clause.
   OMPClause *ActOnOpenMPUseDevicePtrClause(ArrayRef<Expr *> VarList,
                                            const OMPVarListLocTy &Locs);
   /// Called on well-formed 'is_device_ptr' clause.
   OMPClause *ActOnOpenMPIsDevicePtrClause(ArrayRef<Expr *> VarList,
                                           const OMPVarListLocTy &Locs);
 
   /// The kind of conversion being performed.

cfe/trunk/lib/AST/OpenMPClause.cpp

void *Mem = C.Allocate(
       totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
           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<Expr *> Vars, ArrayRef<ValueDecl *> Declarations,
-                      MappableExprComponentListsRef ComponentLists) {
+OMPFromClause *OMPFromClause::Create(
+    const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef<Expr *> Vars,
+    ArrayRef<ValueDecl *> Declarations,
+    MappableExprComponentListsRef ComponentLists, ArrayRef<Expr *> UDMapperRefs,
+    NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId) {
   OMPMappableExprListSizeTy Sizes;
   Sizes.NumVars = Vars.size();
   Sizes.NumUniqueDeclarations = getUniqueDeclarationsTotalNumber(Declarations);
   Sizes.NumComponentLists = ComponentLists.size();
   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
   // number of lists for each unique declaration and the size of each component
   // list.
   // NumComponents x MappableComponent - the total of all the components in all
   // the lists.
   void *Mem = C.Allocate(
       totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
-          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;
 }
 
 OMPFromClause *
 OMPFromClause::CreateEmpty(const ASTContext &C,
                            const OMPMappableExprListSizeTy &Sizes) {
   void *Mem = C.Allocate(
       totalSizeToAlloc<Expr *, ValueDecl *, unsigned,
                        OMPClauseMappableExprCommon::MappableComponent>(
-          Sizes.NumVars, Sizes.NumUniqueDeclarations,
+          2 * Sizes.NumVars, Sizes.NumUniqueDeclarations,
           Sizes.NumUniqueDeclarations + Sizes.NumComponentLists,
           Sizes.NumComponents));
   return new (Mem) OMPFromClause(Sizes);
 }
 
 void OMPUseDevicePtrClause::setPrivateCopies(ArrayRef<Expr *> VL) {
   assert(VL.size() == varlist_size() &&
          "Number of private copies is not the same as the preallocated buffer");
if (!Node->varlist_empty()) {
     }
     OS << ")";
   }
 }
 
 void OMPClausePrinter::VisitOMPFromClause(OMPFromClause *Node) {
   if (!Node->varlist_empty()) {
     OS << "from";
+    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 << ")";
   }
 }
 
 void OMPClausePrinter::VisitOMPDistScheduleClause(OMPDistScheduleClause *Node) {
   OS << "dist_schedule(" << getOpenMPSimpleClauseTypeName(
                            OMPC_dist_schedule, Node->getDistScheduleKind());
   if (auto *E = Node->getChunkSize()) {

cfe/trunk/lib/Basic/OpenMPKinds.cpp

 #include "clang/Basic/OpenMPKinds.def"
         .Default(OMPC_MAP_unknown);
   case OMPC_to:
     return llvm::StringSwitch<unsigned>(Str)
 #define OPENMP_TO_MODIFIER_KIND(Name)                                          \
   .Case(#Name, static_cast<unsigned>(OMPC_TO_MODIFIER_##Name))
 #include "clang/Basic/OpenMPKinds.def"
         .Default(OMPC_TO_MODIFIER_unknown);
+  case OMPC_from:
+    return llvm::StringSwitch<unsigned>(Str)
+#define OPENMP_FROM_MODIFIER_KIND(Name)                                     \
+  .Case(#Name, static_cast<unsigned>(OMPC_FROM_MODIFIER_##Name))
+#include "clang/Basic/OpenMPKinds.def"
+        .Default(OMPC_FROM_MODIFIER_unknown);
   case OMPC_dist_schedule:
     return llvm::StringSwitch<OpenMPDistScheduleClauseKind>(Str)
 #define OPENMP_DIST_SCHEDULE_KIND(Name) .Case(#Name, OMPC_DIST_SCHEDULE_##Name)
 #include "clang/Basic/OpenMPKinds.def"
         .Default(OMPC_DIST_SCHEDULE_unknown);
   case OMPC_defaultmap:
     return llvm::StringSwitch<unsigned>(Str)
 #define OPENMP_DEFAULTMAP_KIND(Name)                                           \
#include "clang/Basic/OpenMPKinds.def"
   case OMPC_num_teams:
   case OMPC_thread_limit:
   case OMPC_priority:
   case OMPC_grainsize:
   case OMPC_nogroup:
   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:
   case OMPC_unified_shared_memory:
   case OMPC_reverse_offload:
   case OMPC_dynamic_allocators:
     break;
   }

 #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:
       return "unknown";
 #define OPENMP_DIST_SCHEDULE_KIND(Name)                                      \
   case OMPC_DIST_SCHEDULE_##Name:                                            \
     return #Name;
 #include "clang/Basic/OpenMPKinds.def"
#include "clang/Basic/OpenMPKinds.def"
   case OMPC_num_teams:
   case OMPC_thread_limit:
   case OMPC_priority:
   case OMPC_grainsize:
   case OMPC_nogroup:
   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:
   case OMPC_unified_shared_memory:
   case OMPC_reverse_offload:
   case OMPC_dynamic_allocators:
     break;
   }

cfe/trunk/lib/Parse/ParseOpenMP.cpp

// Handle dependency type for depend clause.
     }
     if (Data.MapType == OMPC_MAP_unknown) {
       Data.MapType = OMPC_MAP_tofrom;
       Data.IsMapTypeImplicit = true;
     }
 
     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;
+      if (Kind == OMPC_to) {
         auto Modifier = static_cast<OpenMPToModifierKind>(
             getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok)));
         if (Modifier == OMPC_TO_MODIFIER_mapper)
           IsMapperModifier = true;
+      } else {
+        auto Modifier = static_cast<OpenMPFromModifierKind>(
+            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,

 ///       'depend' '(' in | out | inout : list | source ')'
 ///    map-clause:
 ///       'map' '(' [ [ always [,] ] [ close [,] ]
 ///          [ mapper '(' mapper-identifier ')' [,] ]
 ///          to | from | tofrom | alloc | release | delete ':' ] list ')';
 ///    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:
 ///       'is_device_ptr' '(' list ')'
 ///
 /// For 'linear' clause linear-list may have the following forms:
 ///  list
 ///  modifier(list)

cfe/trunk/lib/Sema/SemaOpenMP.cpp

Res = ActOnOpenMPMapClause(MapTypeModifiers, MapTypeModifiersLoc,
                                ReductionOrMapperId, MapType, IsMapTypeImplicit,
                                DepLinMapLoc, ColonLoc, VarList, Locs);
     break;
   case OMPC_to:
     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);
     break;
   case OMPC_is_device_ptr:
     Res = ActOnOpenMPIsDevicePtrClause(VarList, Locs);
     break;
   case OMPC_if:
for (Expr *RE : MVLI.VarList) {
     if (UMIt != UMEnd)
       UnresolvedMapper = *UMIt;
 
     const Expr *VE = RE->IgnoreParenLValueCasts();
 
     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());
-      }
       // We can only analyze this information once the missing information is
       // resolved.
       MVLI.ProcessedVarList.push_back(RE);
       continue;
     }
 
     Expr *SimpleExpr = RE->IgnoreParenCasts();
 
if (!BE)
       continue;
 
     assert(!CurComponents.empty() &&
            "Invalid mappable expression information.");
 
     if (const auto *TE = dyn_cast<CXXThisExpr>(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());
-      }
       // Skip restriction checking for variable or field declarations
       MVLI.ProcessedVarList.push_back(RE);
       MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1);
       MVLI.VarComponents.back().append(CurComponents.begin(),
                                        CurComponents.end());
       MVLI.VarBaseDeclarations.push_back(nullptr);
       continue;
     }
if (CKind == OMPC_map) {
               << getOpenMPDirectiveName(DSAS->getCurrentDirective());
           reportOriginalDsa(SemaRef, DSAS, CurDeclaration, DVar);
           continue;
         }
       }
     }
 
     // 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());
-    }
 
     // Save the current expression.
     MVLI.ProcessedVarList.push_back(RE);
 
     // Store the components in the stack so that they can be used to check
     // against other clauses later on.
     DSAS->addMappableExpressionComponents(CurDeclaration, CurComponents,
                                           /*WhereFoundClauseKind=*/OMPC_map);
OMPClause *Sema::ActOnOpenMPToClause(ArrayRef<Expr *> VarList,
 
   return OMPToClause::Create(
       Context, Locs, MVLI.ProcessedVarList, MVLI.VarBaseDeclarations,
       MVLI.VarComponents, MVLI.UDMapperList,
       MapperIdScopeSpec.getWithLocInContext(Context), MapperId);
 }
 
 OMPClause *Sema::ActOnOpenMPFromClause(ArrayRef<Expr *> VarList,
-                                       const OMPVarListLocTy &Locs) {
+                                       CXXScopeSpec &MapperIdScopeSpec,
+                                       DeclarationNameInfo &MapperId,
+                                       const OMPVarListLocTy &Locs,
+                                       ArrayRef<Expr *> UnresolvedMappers) {
   MappableVarListInfo MVLI(VarList);
-  CXXScopeSpec MapperIdScopeSpec;
-  DeclarationNameInfo MapperId;
-  ArrayRef<Expr *> 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<Expr *> VarList,
                                                const OMPVarListLocTy &Locs) {
   MappableVarListInfo MVLI(VarList);
   SmallVector<Expr *, 8> PrivateCopies;
   SmallVector<Expr *, 8> Inits;
 

cfe/trunk/lib/Sema/TreeTransform.h

return getSema().ActOnOpenMPToClause(VarList, MapperIdScopeSpec, MapperId,
                                          Locs, UnresolvedMappers);
   }
 
   /// Build a new OpenMP 'from' clause.
   ///
   /// By default, performs semantic analysis to build the new statement.
   /// Subclasses may override this routine to provide different behavior.
   OMPClause *RebuildOMPFromClause(ArrayRef<Expr *> VarList,
-                                  const OMPVarListLocTy &Locs) {
-    return getSema().ActOnOpenMPFromClause(VarList, Locs);
+                                  CXXScopeSpec &MapperIdScopeSpec,
+                                  DeclarationNameInfo &MapperId,
+                                  const OMPVarListLocTy &Locs,
+                                  ArrayRef<Expr *> UnresolvedMappers) {
+    return getSema().ActOnOpenMPFromClause(VarList, MapperIdScopeSpec, MapperId,
+                                           Locs, UnresolvedMappers);
   }
 
   /// Build a new OpenMP 'use_device_ptr' clause.
   ///
   /// By default, performs semantic analysis to build the new OpenMP clause.
   /// Subclasses may override this routine to provide different behavior.
   OMPClause *RebuildOMPUseDevicePtrClause(ArrayRef<Expr *> VarList,
                                           const OMPVarListLocTy &Locs) {
if (transformOMPMappableExprListClause<Derived, OMPToClause>(
           *this, C, Vars, MapperIdScopeSpec, MapperIdInfo, UnresolvedMappers))
     return nullptr;
   return getDerived().RebuildOMPToClause(Vars, MapperIdScopeSpec, MapperIdInfo,
                                          Locs, UnresolvedMappers);
 }
 
 template <typename Derived>
 OMPClause *TreeTransform<Derived>::TransformOMPFromClause(OMPFromClause *C) {
-  llvm::SmallVector<Expr *, 16> Vars;
-  Vars.reserve(C->varlist_size());
-  for (auto *VE : C->varlists()) {
-    ExprResult EVar = getDerived().TransformExpr(cast<Expr>(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<Expr *, 16> Vars;
+  CXXScopeSpec MapperIdScopeSpec;
+  DeclarationNameInfo MapperIdInfo;
+  llvm::SmallVector<Expr *, 16> UnresolvedMappers;
+  if (transformOMPMappableExprListClause<Derived, OMPFromClause>(
+          *this, C, Vars, MapperIdScopeSpec, MapperIdInfo, UnresolvedMappers))
+    return nullptr;
+  return getDerived().RebuildOMPFromClause(
+      Vars, MapperIdScopeSpec, MapperIdInfo, Locs, UnresolvedMappers);
 }
 
 template <typename Derived>
 OMPClause *TreeTransform<Derived>::TransformOMPUseDevicePtrClause(
     OMPUseDevicePtrClause *C) {
   llvm::SmallVector<Expr *, 16> Vars;
   Vars.reserve(C->varlist_size());
   for (auto *VE : C->varlists()) {

cfe/trunk/lib/Serialization/ASTReader.cpp

for (unsigned i = 0; i < TotalComponents; ++i) {
     Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
         AssociatedExpr, AssociatedDecl));
   }
   C->setComponents(Components, ListSizes);
 }
 
 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();
   auto TotalComponents = C->getTotalComponentsNum();
 
   SmallVector<Expr *, 16> Vars;
   Vars.reserve(NumVars);
   for (unsigned i = 0; i != NumVars; ++i)
     Vars.push_back(Record.readSubExpr());
   C->setVarRefs(Vars);
 
+  SmallVector<Expr *, 16> UDMappers;
+  UDMappers.reserve(NumVars);
+  for (unsigned I = 0; I < NumVars; ++I)
+    UDMappers.push_back(Record.readSubExpr());
+  C->setUDMapperRefs(UDMappers);
+
   SmallVector<ValueDecl *, 16> Decls;
   Decls.reserve(UniqueDecls);
   for (unsigned i = 0; i < UniqueDecls; ++i)
     Decls.push_back(Record.readDeclAs<ValueDecl>());
   C->setUniqueDecls(Decls);
 
   SmallVector<unsigned, 16> ListsPerDecl;
   ListsPerDecl.reserve(UniqueDecls);

cfe/trunk/lib/Serialization/ASTWriter.cpp

 }
 
 void OMPClauseWriter::VisitOMPFromClause(OMPFromClause *C) {
   Record.push_back(C->varlist_size());
   Record.push_back(C->getUniqueDeclarationsNum());
   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())
     Record.push_back(N);
   for (auto N : C->all_lists_sizes())
     Record.push_back(N);
   for (auto &M : C->all_components()) {
     Record.AddStmt(M.getAssociatedExpression());

cfe/trunk/test/OpenMP/declare_mapper_ast_print.c

// CHECK: #pragma omp declare mapper (id : struct vec v) map(tofrom: v.len)
     struct vec vv;
     struct dat dd[10];
 #pragma omp target map(mapper(id) alloc: vv)
 // CHECK: #pragma omp target map(mapper(id),alloc: vv)
     { vv.len++; }
 #pragma omp target map(mapper(default), from: dd[0:10])
 // CHECK: #pragma omp target map(mapper(default),from: dd[0:10])
     { dd[0].i++; }
-#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;
 }
 // CHECK: }
 
 #endif

cfe/trunk/test/OpenMP/declare_mapper_ast_print.cpp

#pragma omp declare mapper(id: N1::vec v) map(v.len)
   }
   struct foodat fd;
 #pragma omp target map(mapper(id) alloc: fd)
   { fd.a++; }
 #pragma omp target map(mapper(idd) alloc: fd.b)
   { fd.b.k++; }
 #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;
 }
 
 // CHECK: template <typename T> T foo(T a) {
 // CHECK: #pragma omp declare mapper (id : struct foodat v) map(tofrom: v.a)
 // CHECK: #pragma omp declare mapper (idd : struct foodatchild v) map(tofrom: v.k)
 // CHECK: #pragma omp declare mapper (id : N1::vec v) map(tofrom: v.len)
 // CHECK: {
 // CHECK: #pragma omp declare mapper (id : N1::vec v) map(tofrom: v.len)
 // CHECK: }
 // CHECK: #pragma omp target map(mapper(id),alloc: fd)
 // CHECK: #pragma omp target map(mapper(idd),alloc: fd.b)
 // CHECK: #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>(int a) {
 // CHECK: #pragma omp declare mapper (id : struct foodat v) map(tofrom: v.a)
 // CHECK: #pragma omp declare mapper (idd : struct foodatchild v) map(tofrom: v.k)
 // CHECK: #pragma omp declare mapper (id : N1::vec v) map(tofrom: v.len)
 // CHECK: {
 // CHECK: #pragma omp declare mapper (id : N1::vec v) map(tofrom: v.len)
 // CHECK: }
 // CHECK: #pragma omp target map(mapper(id),alloc: fd)
 // CHECK: #pragma omp target map(mapper(idd),alloc: fd.b)
 // CHECK: #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() {
 int main() {
   N1::vec vv, vvv;
   N1::vecchild vc;
   dat<double> dd;
 #pragma omp target map(mapper(N1::id) tofrom: vv) map(mapper(dat<double>::id) alloc: vvv)
 // CHECK: #pragma omp target map(mapper(N1::id),tofrom: vv) map(mapper(dat<double>::id),alloc: vvv)
   { vv.len++; }
 #pragma omp target map(mapper(N1::id) tofrom: vc)
 // CHECK: #pragma omp target map(mapper(N1::id),tofrom: vc)
   { vc.len++; }
 #pragma omp target map(mapper(default) tofrom: dd)
 // CHECK: #pragma omp target map(mapper(default),tofrom: dd)
   { dd.d++; }
 
 #pragma omp target update to(mapper(N1::id) : vc)
 // CHECK: #pragma omp target update to(mapper(N1::id): vc)
 #pragma omp target update to(mapper(dat<double>::id): vvv)
 // CHECK: #pragma omp target update to(mapper(dat<double>::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<double>::id): vvv)
+// CHECK: #pragma omp target update from(mapper(dat<double>::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)
   {
 #pragma omp declare mapper(id: N1::vec v) map(v.len)
 // CHECK: #pragma omp declare mapper (id : N1::vec v) map(tofrom: v.len)
   }
   return foo<int>(0);
 }
 // CHECK: }
 
 #endif

cfe/trunk/test/OpenMP/declare_mapper_codegen.cpp

 
 class C {
 public:
   int a;
 };
 
 #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){
   int i = a;
   C c;
   c.a = a;
 
   // CHECK-DAG: call i32 @__tgt_target(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}})
void foo(int a){
   // 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)
+
+  // 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: define internal void [[KERNEL]](%class.C* {{.+}}[[ARG:%.+]])
 // CHECK: [[ADDR:%.+]] = alloca %class.C*,
 // CHECK: store %class.C* [[ARG]], %class.C** [[ADDR]]
 // CHECK: [[CADDR:%.+]] = load %class.C*, %class.C** [[ADDR]]
 // CHECK: [[CAADDR:%.+]] = getelementptr inbounds %class.C, %class.C* [[CADDR]], i32 0, i32 0
 // CHECK: [[VAL:%[^,]+]] = load i32, i32* [[CAADDR]]
 // CHECK: {{.+}} = add nsw i32 [[VAL]], 1
 // CHECK: }
 
 #endif

cfe/trunk/test/OpenMP/declare_mapper_messages.c

 
 #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;
 }

cfe/trunk/test/OpenMP/declare_mapper_messages.cpp

 #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<int>::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<int>::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'}}
   }
   return arg;
 }